CUDA-内核应该是动态崩溃取决于块大小
我想做一个稀疏矩阵,密集向量乘法。让我们假设压缩矩阵中唯一的存储格式是压缩行存储CRS。CUDA-内核应该是动态崩溃取决于块大小
我的内核如下所示:
__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使内核也失败。再次,只是最后一个值没有被计算。
您的动态分配的共享内存大小不正确。现在你这样做是:
krnlSpMVmul1<<<grid_dim, block_dim, BLOCK_SIZE>>>(.....)
共享内存的大小应在字节给出。使用每个块的64个线程,这意味着您将为16个浮点大小的单词分配足够的共享内存,并解释为什么每行幻数17个条目导致失败 - 您有共享缓冲区溢出,这将导致GPU并中止内核。
你应该做这样的事情:
krnlSpMVmul1<<<grid_dim, block_dim, BLOCK_SIZE * sizeof(float)>>>(.....)
这会给你正确的动态共享内存的大小,应消除问题。
最后一个问题。我试图用真实数据运行内核。我得到了一个有成千上万行的矩阵。似乎所有行(不是非零非常多的行)都被正确计算。如果内核一旦出现第一次越界访问就会失败,这怎么可能呢? – 2012-02-28 13:02:12
答案可能取决于您使用的是哪种GPU(在较旧的硬件上,结果可能是错误的,在Fermi卡上,如果您正确检查,应该得到未指定的启动失败错误)。我也会推荐用'cuda-memcheck'运行你的代码。它会检测并报告共享和全局内存访问的出界。 – talonmies 2012-02-28 13:07:36
非常感谢您的努力。它真的很感激(其实我真的有一个CC1.1设备运行) – 2012-02-28 13:09:10
你可以发布完整的内核代码吗?这很可能是问题出在代码中,你已经省略了。你还可以显示你用来调用内核的内核参数吗? – talonmies 2012-02-28 12:15:58