CUDA reduce find max

Question

我正在尝试编写我自己的简化搜索的最大数量的实现。这是我的第一个CUDA C程序，我遇到了一些我无法理解的错误。

我代码:

__global__
void reduce (box * d_in_data, int insize, box * d_out_data)
{
/*
for example:
array size = 25200 (size variable, keep in d_in_data)
gridDim = 252
blockDim = 10
*/
uint32_t tid = threadIdx.x;
//__shared__ box sdata[insize / gridDim.x]; // 25200 / 252 = 100
extern __shared__ box sdata[]; // 25200 / 252 = 100
/*spt its 'size per thread' for filling sdata*/
/* (25200 / 252) / 10 */ /* each thread should be copy 10 elements from d_in_data to sdata  */
uint32_t spt = (insize / gridDim.x) / blockDim.x;
box (* sh_ptr)[blockDim.x][spt] = (box (*) [blockDim.x][spt]) &sdata;
for (int i = 0;i < spt;i++){
(*sh_ptr)[tid][i] = d_in_data[blockIdx.x * (insize / gridDim.x) + (tid * blockDim.x + i)];
}
__syncthreads();
box temp = (*sh_ptr)[tid][0];
for (int i = 0;i < spt;i++){
box i_box = (*sh_ptr)[tid][i];
temp = temp.s < i_box.s ? i_box : temp;
}
(*sh_ptr)[tid][0] = temp;
__syncthreads();
box (* d_out_ptr)[gridDim.x][blockDim.x] =  (box (*)[gridDim.x][blockDim.x]) d_out_data;
(*d_out_ptr)[blockIdx.x][tid]  = (*sh_ptr)[tid][0];

}

1. __shared__ box sdata[insize / gridDim.x];-为什么这里的"不能用作常量">?如何明确设置大小?

2. box (* sh_ptr)[blockDim.x][spt] = (box (*) [blockDim.x][spt]) &sdata;-"error:表达式必须有一个常量值">，"注意:试图访问运行时存储">，"注意:变量"spt">。在我的情况下如何正确书写?

3. uint32_t spt = (insize / gridDim.x) / blockDim.x;-"这里不能用作常量">

4. box (* d_out_ptr)[gridDim.x][blockDim.x] = (box (*)[gridDim.x][blockDim.x]) d_out_data;-"错误:表达式必须有一个常量">，"注意:试图访问运行时存储">，这里也是一样。

请告诉我每一段我错在哪里。

Answer 1

您的错误与CUDA无关。

• __shared__ box sdata[insize / gridDim.x];-为什么这里不能用作常量？ 如何明确设置大小?
  在C/c++中，数组的大小必须在编译时知道。这里，insize / gridDim.x必须在编译时知道，但insize肯定不是。

  extern __shared__ box sdata[];
  

  在我看来完全没问题。如果你需要这个数组的大小(你当然需要)，你可以在运行时通过函数参数等形式计算它。

• box (* sh_ptr)[blockDim.x][spt] = (box (*) [blockDim.x][spt]) &sdata;- "错误:表达式必须有一个常量"，"备注:尝试访问运行时存储";备注:变量"spt"的值。在我的情况下如何正确书写?
  这里你尝试定义一个指向2D数组的指针。只有在编译时知道数组的最后一个维数，但不知道spt的情况下才有可能这样做。为了规避它，您需要自己计算共享内存(sdata)的地址。这是一个1D数组，请这样处理。

• 3号错误似乎是诊断信息，而不是错误。

• 4号错误与2号错误类型相同。

进一步的阅读参见，例如，指针到多维数组在C?