1.題面
題目要求
向量加法
實現一個程序,在GPU上對兩個包含32位浮點數的向量執行逐元素加法。該程序應接受兩個長度相等的輸入向量,并生成一個包含它們和的輸出向量。
實現要求
禁止使用外部庫
solve函數簽名必須保持不變
最終結果必須存儲在向量C中
示例1:
輸入:A = [1.0, 2.0, 3.0, 4.0]
B = [5.0, 6.0, 7.0, 8.0]
輸出:C = [6.0, 8.0, 10.0, 12.0]
示例2:
輸入:A = [1.5, 1.5, 1.5]
B = [2.3, 2.3, 2.3]
輸出:C = [3.8, 3.8, 3.8]
約束條件
輸入向量A和B長度相同
1 ≤ N ≤ 100,000,000
2.已有代碼解析
函數參數與目的
void solve(const float* A, const float* B, float* C, int N)
- 該函數接收三個指向GPU內存的指針:
A
和B
是輸入向量,C
是輸出向量。 N
表示向量的長度,即元素個數。
線程與線程塊配置
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
- 線程塊大小(
threadsPerBlock
):每個線程塊包含256個線程。這是CUDA編程中常用的配置,適合大多數GPU架構。 - 網格大小(
blocksPerGrid
):根據向量長度N
計算所需的線程塊數量。使用向上取整公式確保所有元素都被處理。例如:- 當
N=1000
時,blocksPerGrid = (1000 + 256 - 1) / 256 = 4
。 - 每256個線程就打包成一個block,blocksPerGrid就表示打包了多少個
- 當
核函數調用
vector_add<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, N);
- 核函數(
vector_add
):這是一個在GPU上執行的函數,負責逐元素地將A
和B
相加,結果存入C
。 - 執行配置(
<<<...>>>
):指定網格和線程塊的維度。這里使用一維網格和一維線程塊。 - 核函數會被并行調用,每個線程負責計算一個元素:
C[i] = A[i] + B[i]
。
同步操作
cudaDeviceSynchronize();
- 確保核函數執行完成后才返回。這對于需要立即使用結果的場景很重要。
3.QA環節
Q1: Grid是什么,為什么bocksPerGrid就是塊的數目?
Grid是CUDA編程模型中的一個概念,表示線程塊(Block)的集合。在GPU并行計算中,一個Grid是核函數(Kernel)的一次執行實例,它包含多個線程塊,這些線程塊可以并行或串行執行,具體取決于硬件資源。
blocksPerGrid是根據向量長度N
和每個線程塊的線程數量threadsPerBlock
計算得出的。在CUDA中,網格(Grid)和線程塊(Block)都可以是一維、二維或三維的。在這個向量加法的例子中,我們使用了一維網格,因此blocksPerGrid
直接表示網格中一維方向上的線程塊數量。計算公式為:
blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
這個公式實現了向上取整的效果,確保所有N
個元素都有對應的線程來處理。例如:
- 當
N=1000
,threadsPerBlock=256
時,blocksPerGrid = (1000 + 256 - 1) / 256 = 4
- 每個線程塊處理256個元素,4個線程塊總共可以處理1024個元素,覆蓋了
N=1000
的需求。
Q2: 這里使用“一維網格”和“一維線程塊”是什么含義,還有其他選擇嗎,比如二維的?
一維網格和一維線程塊
在CUDA中,線程組織是分級的:
- 線程塊(Block):由多個線程組成,可以是一維、二維或三維
- 網格(Grid):由多個線程塊組成,同樣可以是一維、二維或三維
在向量加法的例子中:
- 一維線程塊:每個線程塊中的線程按一維數組排列,線程索引由
threadIdx.x
表示 - 一維網格:線程塊按一維數組排列,線程塊索引由
blockIdx.x
表示
每個線程的全局索引計算方式為:
int i = blockIdx.x * blockDim.x + threadIdx.x;
其中:
blockDim.x
是每個線程塊的線程數量(256)blockIdx.x
是當前線程塊的索引threadIdx.x
是當前線程在線程塊內的索引
二維網格和二維線程塊
是的,CUDA支持二維甚至三維的網格和線程塊,適用于處理二維或三維數據結構(如圖像、矩陣)。例如:
二維線程塊
對于圖像處理,可以使用二維線程塊:
dim3 threadsPerBlock(16, 16); // 每個線程塊有16×16=256個線程
線程索引由threadIdx.x
和threadIdx.y
表示,線程塊內的線程按二維矩陣排列。
二維網格
對應的網格也可以是二維的:
dim3 blocksPerGrid((width + threadsPerBlock.x - 1) / threadsPerBlock.x,(height + threadsPerBlock.y - 1) / threadsPerBlock.y
);
線程塊索引由blockIdx.x
和blockIdx.y
表示,網格中的線程塊按二維矩陣排列。
全局索引計算
對于二維網格和二維線程塊,全局索引計算方式為:
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row * width + col; // 轉換為一維索引
為什么向量加法使用一維組織?
向量加法處理的是一維數組,每個元素的計算只依賴于自身索引,使用一維組織更直觀且高效:
- 線程索引計算簡單
- 內存訪問模式連續,提高緩存利用率
對于二維數據(如圖像、矩陣),使用二維組織更自然,能夠更好地映射數據結構和計算邏輯。
解題思路
__global__ void vector_add(const float* A, const float* B, float* C, int N) {int i = blockIdx.x * blockDim.x + threadIdx.x;if(i<N){C[i]=A[i]+B[i];}
}
根據第幾個block和在block內的局部地址,得到總的id。注意i<N的邊界約束,這是因為我們無法保證N是256的倍數,可能會產生數組越界。