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)>>>(.....) 

这会给你正确的动态共享内存的大小，应消除问题。