CUDA自定义reduce求最大值核函数编译报错问题排查咨询

问题1：__shared__ box sdata[insize / gridDim.x]; 报错"here cannot be used as a constant" #

错误原因 #

CUDA的静态共享内存（无extern修饰的__shared__变量）的大小必须是编译期就能确定的常量，而insize是核函数传入的参数、gridDim.x是运行时才确定的网格维度，二者都属于运行时变量，不能用来指定静态共享内存的大小。

解决方法 #

你注释中写的extern __shared__ box sdata[];就是正确的动态共享内存声明方式，只需要在核函数启动时显式传入共享内存的大小即可，启动示例如下：

int block_size = 10;
int grid_size = 252;
int shared_mem_size = (insize / grid_size) * sizeof(box);
reduce<<<grid_size, block_size, shared_mem_size>>>(d_in_data, insize, d_out_data);

问题2：box (* sh_ptr)[blockDim.x][spt] = (box (*) [blockDim.x][spt]) &sdata; 报错"expression must have a constant value" #

错误原因 #

C/C++标准要求，数组指针的维度参数必须是编译期常量，而blockDim.x是运行时才能获取的块维度、spt是用运行时参数计算出来的变量，二者都不符合编译期常量的要求，因此无法用来声明多维数组指针。

解决方法 #

放弃使用多维数组指针，直接把共享内存当成一维数组，通过手动计算偏移量访问对应位置即可，原逻辑中的(*sh_ptr)[tid][i]可以直接替换为sdata[tid * spt + i]。

问题3：uint32_t spt = (insize / gridDim.x) / blockDim.x; 报错"here cannot be used as a constant" #

错误原因 #

这行本身的变量定义和赋值语法没有问题，报错是编译器的回溯提示：你后续代码中把运行时才能确定值的spt，用到了需要编译期常量的场景（比如问题2里的数组指针维度声明），错误指向这行是为了告诉你spt是运行时生成的变量，不能作为常量使用。

解决方法 #

删除所有把spt作为编译期常量使用的代码（比如问题2的多维数组指针声明），这行的报错会自动消失。

问题4：box (* d_out_ptr)[gridDim.x][blockDim.x] = (box (*)[gridDim.x][blockDim.x]) d_out_data; 报错"expression must have a constant value" #

错误原因 #

和问题2的原理一致：gridDim.x、blockDim.x都是运行时变量，不能作为数组指针的维度参数使用。

解决方法 #

同样用一维指针偏移访问，你要赋值的逻辑可以直接写成：

d_out_data[blockIdx.x * blockDim.x + tid] = sdata[tid * spt];

修改后可正常编译的核心代码参考 #

__global__
void reduce (box * d_in_data, int insize, box * d_out_data)
{
    uint32_t tid = threadIdx.x;
    extern __shared__ box sdata[];
    uint32_t block_data_len = insize / gridDim.x;
    uint32_t spt = block_data_len / blockDim.x;

    // 拷贝数据到共享内存
    for (int i = 0;i < spt;i++){
        sdata[tid * spt + i] = d_in_data[blockIdx.x * block_data_len + tid * spt + i];
    }
    __syncthreads();

    // 每个线程计算自己负责的spt个元素的最大值
    box temp = sdata[tid * spt];
    for (int i = 0;i < spt;i++){
        box i_box = sdata[tid * spt + i];
        temp = temp.s < i_box.s ? i_box : temp;
    }
    sdata[tid * spt] = temp;
    __syncthreads();

    // 输出到全局内存
    d_out_data[blockIdx.x * blockDim.x + tid] = sdata[tid * spt];
}

内容的提问来源于stack exchange，提问作者alph alph