Cuda內核與減少 - 2個矩陣點積的邏輯錯誤

Question

我剛剛開始與CUDA，並試圖圍繞CUDA減少算法包裹我的大腦。就我而言，我一直試圖獲得兩個矩陣的點積。但是對於只有大小爲2的矩陣，我得到了正確的答案。對於任何其他大小的矩陣，我錯了。

這只是測試，所以我保持矩陣大小非常小。只有大約100塊，所以只有1塊可以滿足所有需求。 任何幫助將不勝感激。謝謝！

這裏是常規代碼

float* ha = new float[n]; // matrix a 
float* hb = new float[n]; // matrix b 
float* hc = new float[1]; // sum of a.b 

float dx = hc[0]; 
float hx = 0; 
// dot product 
for (int i = 0; i < n; i++) 
    hx += ha[i] * hb[i]; 

這裏是我的CUDA內核

__global__ void sum_reduce(float* da, float* db, float* dc, int n) 
{ 
    int tid = threadIdx.x; 
    dc[tid] = 0; 
    for (int stride = 1; stride < n; stride *= 2) { 
     if (tid % (2 * stride) == 0) 
       dc[tid] += (da[tid] * db[tid]) + (da[tid+stride] * db[tid+stride]); 
     __syncthreads(); 
    } 
} 

我的完整代碼：http://pastebin.com/zS85URX5

Answer 1

希望你能弄清楚爲什麼它的工作原理爲N = 2那麼讓我們跳過這一點，並看看爲什麼它在其他情況下失敗，我們選擇n = 4。當n = 4時，您有4個線程，編號爲0至3。

在您的for循環，跨度= 1，因此通過該if測試的線程是線程0的第一次迭代和2

thread 0: dc[0] += da[0]*db[0] + da[1]*db[1]; 
thread 2: dc[2] += da[2]*db[2] + da[3]*db[3]; 

到目前爲止好。在for循環的第二次迭代中，stride是2，所以通過if測試的線程是線程0（僅）。

thread 0: dc[0] += da[0]*db[0] + da[2]*db[2]; 

但是這沒有意義，也不是我們想要的。我們想要的是這樣的：

dc[0] += dc[2]; 

所以它被打破。我花了一點時間想着如何通過幾個步驟來解決這個問題，但對我來說這只是一種減少。如果你用這段代碼替換你的內核代碼，我想你會有很好的結果。它不像你的代碼，但它是最接近我可以得到的東西，適用於你所設想的所有情況（即最大線程塊大小，使用一個塊）

// CUDA kernel code 
__global__ void sum_reduce(float* da, float* db, float* dc, int n) 
{ 
    int tid = threadIdx.x; 
    // do multiplication in parallel for full width of threads 
    dc[tid] = da[tid] * db[tid]; 
    // wait for all threads to complete multiply step 
    __syncthreads(); 
    int stride = blockDim.x; 
    while (stride > 1){ 
     // handle odd step 
     if ((stride & 1) && (tid == 0)) dc[0] += dc[stride - 1]; 
     // successively divide problem by 2 
     stride >>= 1; 
     // add each upper half element to each lower half element 
     if (tid < stride) dc[tid] += dc[tid + stride]; 
     // wait for all threads to complete add step 
     __syncthreads(); 
     } 
} 

請注意，我並沒有真正使用n參數。由於您使用n線程啓動內核，因此在這種情況下，blockDim.x內置變量等於n。