CUDA 减少查找最大值答案

【问题标题】：CUDA reduce find maxCUDA 减少查找最大值
【发布时间】：2021-12-11 06:42:36
【问题描述】：

我正在尝试编写自己的最大数量的还原搜索实现。这是我的第一个 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];
    
}

__shared__ box sdata[insize / gridDim.x]; - 为什么“这里不能用作常量”？如何明确设置大小？
box (* sh_ptr)[blockDim.x][spt] = (box (*) [blockDim.x][spt]) &sdata; - "错误：表达式必须有一个常量值" , "注意：尝试访问运行时存储", "注意：变量“spt”的值。在我的情况下如何正确书写？
uint32_t spt = (insize / gridDim.x) / blockDim.x; - “这里不能用作常量”
box (* d_out_ptr)[gridDim.x][blockDim.x] = (box (*)[gridDim.x][blockDim.x]) d_out_data; - “错误：表达式必须有一个常量值”，“注意：尝试访问运行时存储”，这里也一样。

请告诉我我在每个段落中哪里出错了。

【问题讨论】：

1 和 3：您需要编译时间常数。 2 和 4：只使用对数组的引用而不是进行复杂的索引。

标签： c++ cuda reduce

【解决方案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）的地址。这是一个一维数组，这样处理。
错误号 3 似乎是诊断消息，而不是错误。
错误编号 4 与错误编号 2 的类型相同。

如需进一步阅读，请参阅Pointer to Multidimensional Array in C?

【讨论】：