pycuda共享内存错误“pycuda._driver.LogicError：cuLaunchKernel失败：无效值”

Question

我有一个奇怪的问题，我不能确定哪个产地：

我有一个工作的内核，并获得一些特殊的矩阵向量乘法，这我想加快速度。基本上，大矩阵（10^6×10^6）由少量小矩阵构成。所以我想把这些数据放在共享内存中。然而，当我尝试添加共享内存，我只得到了错误：

pycuda._driver.LogicError：cuLaunchKernel失败：无效值

所以我的工作核心是：

#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实例。

希望你能澄清一些我的困惑......

Answer 1

根据你的描述，共享MEM您分配太大。

__shared__ double gpu_matrix[BLOCK_SIZE*BLOCK_SIZE*13]; 

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

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

$CUDA_ROOT/tools/CUDA_Occupancy_Calculator.xls 

在另一方面，通过垫-VEC-MUL状内核需要共享MEM的尺寸应该能够从O（BLOCK_SIZE^2）降低到O（BLOCK_SIZE）。你可能想在实现你自己的代码之前阅读一些成功的mat-vec-mul内核的代码，如MAGMA。