c - cuda共享内存静态分配错误

Question

我使用共享内存编写了一个 CUDA 代码，如下所示:

__global__ void matrix_mul_shared(float *ad,float *bd,float *cd,int N)
{
    float pvalue=0;
    int TILE=blockDim.x;
    int ty=threadIdx.y;
    int tx=threadIdx.x;

    //allocate shared memory per block
    __shared__ float ads[1][1];
    __shared__ float bds[1][1];

 .

。。}

此代码有效，但以下代码失败；

__global__ void matrix_mul_shared(float *ad,float *bd,float *cd,int N)
{
    float pvalue=0;
    int TILE=blockDim.x;
    int ty=threadIdx.y;
    int tx=threadIdx.x;

    //allocate shared memory per block
    __shared__ float ads[TILE][TILE];
    __shared__ float bds[TILE][TILE];

 .
. 
.
}

编译器期望我分配共享内存的行有一些恒定的东西。它说(我忘记了确切的错误，但它是这样的):

The parameters should be a constant

我能够使用 printf 并打印 TILE 的值，结果是 1。那么为什么会出现此错误？

Answer 1

我认为您收到的错误是

error: expression must have a constant value

变量TILE不是编译器意义上的常量。编译器正在询问编译时已知的共享内存数组的维度。

可能的解决方案:

#define TILE 16

__global__ void matrix_mul_shared(float *ad,float *bd,float *cd,int N)
{
    ...

    __shared__ float ads[TILE][TILE];

    ...
}