CUDA-內核應該是動態崩潰取決於塊大小

Question

我想做一個稀疏矩陣，密集向量乘法。讓我們假設壓縮矩陣中唯一的存儲格式是壓縮行存儲CRS。

我的內核如下所示：

__global__ void 
krnlSpMVmul1(
     float *data_mat, 
     int num_nonzeroes, 
     unsigned int *row_ptr, 
     float *data_vec, 
     float *data_result) 
{ 
    extern __shared__ float local_result[]; 
    local_result[threadIdx.x] = 0; 

    float vector_elem = data_vec[blockIdx.x]; 

    unsigned int start_index = row_ptr[blockIdx.x]; 
    unsigned int end_index = row_ptr[blockIdx.x + 1]; 

    for (int index = (start_index + threadIdx.x); (index < end_index) && (index < num_nonzeroes); index += blockDim.x) 
     local_result[threadIdx.x] += (data_mat[index] * vector_elem); 

    __syncthreads(); 

    // Reduction 

    // Writing accumulated sum into result vector 
} 

正如你所看到的內核應該是儘可能的天真，它甚至做了幾件事情錯誤（如vector_elem是不總是正確的值）。我知道這些事情。

現在我的問題： 假設我使用的32個或64個線程塊大小。一旦我的矩陣中有一行有16個以上的非零值（例如17），只有前16個乘法完成並保存到共享內存中。我知道第17次乘法的結果local_result[16]的值僅爲零。使用16或128個線程塊可修復解釋的問題。

因爲我是相當新的CUDA，我可能都忽略了最簡單的事情，但我不能彌補任何更多的情況來看待。

幫助非常感謝！

---

朝向talonmies編輯評論：

我打印哪些是在local_result[16]計算後直接的值。這是0。儘管如此，這裏是遺漏碼：

的減少部分：

int k = blockDim.x/2; 
while (k != 0) 
{ 
    if (threadIdx.x < k) 
     local_result[threadIdx.x] += local_result[threadIdx.x + k]; 
    else 
     return; 

    __syncthreads(); 

    k /= 2; 
} 

，以及如何我寫的結果返回給全局存儲器：

data_result[blockIdx.x] = local_result[0]; 

這就是我的一切。

現在我測試方案，其中包含由單排的有17件，所有都是非零矩陣。該緩衝區是這樣的僞代碼：

float data_mat[17] = { val0, .., val16 } 
unsigned int row_ptr[2] = { 0, 17 } 
float data_vec[17] = { val0 } // all values are the same 
float data_result[1] = { 0 } 

並且那我的包裝功能的摘錄：

float *dev_data_mat; 
unsigned int *dev_row_ptr; 
float *dev_data_vec; 
float *dev_data_result; 

// Allocate memory on the device 
HANDLE_ERROR(cudaMalloc((void**) &dev_data_mat, num_nonzeroes * sizeof(float))); 
HANDLE_ERROR(cudaMalloc((void**) &dev_row_ptr, num_row_ptr * sizeof(unsigned int))); 
HANDLE_ERROR(cudaMalloc((void**) &dev_data_vec, dim_x * sizeof(float))); 
HANDLE_ERROR(cudaMalloc((void**) &dev_data_result, dim_y * sizeof(float))); 

// Copy each buffer into the allocated memory 
HANDLE_ERROR(cudaMemcpy(
     dev_data_mat, 
     data_mat, 
     num_nonzeroes * sizeof(float), 
     cudaMemcpyHostToDevice)); 
HANDLE_ERROR(cudaMemcpy(
     dev_row_ptr, 
     row_ptr, 
     num_row_ptr * sizeof(unsigned int), 
     cudaMemcpyHostToDevice)); 
HANDLE_ERROR(cudaMemcpy(
     dev_data_vec, 
     data_vec, 
     dim_x * sizeof(float), 
     cudaMemcpyHostToDevice)); 
HANDLE_ERROR(cudaMemcpy(
     dev_data_result, 
     data_result, 
     dim_y * sizeof(float), 
     cudaMemcpyHostToDevice)); 

// Calc grid dimension and block dimension 
dim3 grid_dim(dim_y); 
dim3 block_dim(BLOCK_SIZE); 

// Start kernel 
krnlSpMVmul1<<<grid_dim, block_dim, BLOCK_SIZE>>>(
     dev_data_mat, 
     num_nonzeroes, 
     dev_row_ptr, 
     dev_data_vec, 
     dev_data_result); 

我希望這是簡單，但將解釋的事情，如果它是任何權益。

一兩件事：我剛剛意識到使用128 BLOCK_SIZE，並具有33個nonzeroes使內核也失敗。再次，只是最後一個值沒有被計算。

Answer 1

您的動態分配的共享內存大小不正確。現在你這樣做是：

krnlSpMVmul1<<<grid_dim, block_dim, BLOCK_SIZE>>>(.....) 

共享內存的大小應在字節給出。使用每個塊的64個線程，這意味着您將爲16個浮點大小的單詞分配足夠的共享內存，並解釋爲什麼每行幻數17個條目導致失敗 - 您有共享緩衝區溢出，這將導致GPU並中止內核。

你應該做這樣的事情：

krnlSpMVmul1<<<grid_dim, block_dim, BLOCK_SIZE * sizeof(float)>>>(.....) 

這會給你正確的動態共享內存的大小，應消除問題。