pycuda 共享内存错误“pycuda._driver.LogicError: cuLaunchKernel failed: invalid value”答案

【问题标题】：pycuda shared memory error "pycuda._driver.LogicError: cuLaunchKernel failed: invalid value"pycuda 共享内存错误“pycuda._driver.LogicError: cuLaunchKernel failed: invalid value”
【发布时间】：2013-08-27 13:34:43
【问题描述】：

我有一个奇怪的问题，我无法确定其来源：

我有一个用于一些特殊矩阵向量乘法的工作内核，我想加快它的速度。基本上，大矩阵（10^6 乘以 10^6）是由几个小矩阵构成的。所以我想把这些数据放在共享内存中。但是，当我尝试添加共享内存时，我只收到错误消息：

pycuda._driver.LogicError: cuLaunchKernel failed: invalid value

所以我的工作内核是：

#define FIELD_SIZE {field}
#define BLOCK_SIZE {block}

__global__ void MatrixMulKernel(double *gpu_matrix, double *gpu_b, double *gpu_y)
{
    int tx = ... + threadIdx.x;

    if(tx < FIELD_SIZE*FIELD_SIZE*BLOCK_SIZE) 
    { ... multiplication ... }
}

如果我尝试添加共享内存部分，它看起来像

#define FIELD_SIZE {field}
#define BLOCK_SIZE {block}

__global__ void MatrixMulKernel(double *gpu_matrix_ptr, double *gpu_b, double *gpu_y)
{
    __shared__ double gpu_matrix[BLOCK_SIZE*BLOCK_SIZE*13];

    int tx = ... + threadIdx.x;
    if(tx < BLOCK_SIZE*BLOCK_SIZE*13) {  gpu_matrix[tx] = gpu_matrix_ptr[tx];  }
    __syncthreads();

    if(tx < FIELD_SIZE*FIELD_SIZE*BLOCK_SIZE) 
    { ... multiplication ... }
}

这是我唯一改变的部分，所以基本上它必须是 gpu_matrix[tx] = gpu_matrix_ptr[tx] 语句，不是吗？但我看不出那应该是怎样的。我基本上试图从 pycuda 示例中复制平铺矩阵乘法示例。 http://wiki.tiker.net/PyCuda/Examples/MatrixmulTiled

调用是：

self.kernel.prepare([np.intp, np.intp, np.intp])
self.kernel.prepared_call(grid_shape,
              block_shape,
              self.matrix_gpu.gpudata,
              b_gpu.gpudata,
              y_gpu.gpudata)

其中 matrix_gpu、b_gpu 和 y_gpu 是 pycuda.gpuarray 实例。

希望你能解开我的一些困惑……

【问题讨论】：

无效值。您的块大小和字段大小的值是多少？也许他们太大了
block_size 是 64，field_size 是 1024。为什么会有问题？编辑：但是你是正确的，将块大小减少到 16 是有效的。无论如何我可以增加共享内存的大小吗？

标签： cuda shared-memory pycuda

【解决方案1】：

根据你的描述，你分配的shared mem太大了。

__shared__ double gpu_matrix[BLOCK_SIZE*BLOCK_SIZE*13];

shared mem是cuda gpu的硬件资源之一。总大小约为 48KBytes，不能增加。

CUDA其实在下面的目录中提供了一个工具来帮助你计算你可以使用的硬件资源。

$CUDA_ROOT/tools/CUDA_Occupancy_Calculator.xls

另一方面，mat-vec-mul-like 内核所需的共享内存大小应该能够从 O(BLOCK_SIZE^2) 减少到 O(BLOCK_SIZE)。在实现自己的之前，您可能需要阅读一些成功的 mat-vec-mul 内核的代码，例如 MAGMA。

【讨论】：