CUDA histogram2d 不工作答案

【问题标题】：CUDA histogram2d not workingCUDA histogram2d 不工作
【发布时间】：2016-12-26 02:59:32
【问题描述】：

由于 CUDA 似乎缺乏像样的 2D 直方图（我可以找到...欢迎指点），我正在尝试使用 pyCUDA 自己实现它。

这是直方图的样子（使用 Numpy）：

这是我目前所得到的：

code = '''
__global__ void histogram2d(const float *in_x, const float *in_y, const float *in_w, float *out) {{
    int start = blockIdx.x * blockDim.x + threadIdx.x;

    float *block_out = &out[{xres} * {yres} * {num_chans} * blockIdx.x];

    for(int i = 0; i < {length}; i++) {{
        float x = in_x[start + i];
        float y = in_y[start + i];
        int w_idx = (start + i) * {num_chans};

        int xbin = (int) (((x - {xmin}) / {xptp}) * {xres});
        int ybin = (int) (((y - {ymin}) / {yptp}) * {yres});

        if (0 <= xbin && xbin < {xres} && 0 <= ybin && ybin < {yres}) {{
            for(int c = 0; c < {num_chans}; c++) {{
                atomicAdd(&block_out[(ybin * {xres} + xbin) * {num_chans} + c], in_w[w_idx + c]);
            }}
        }}
    }}
}}
'''.format(**args)

------

__global__ void histogram2d(const float *in_x, const float *in_y, const float *in_w, float *out) {
    int start = blockIdx.x * blockDim.x + threadIdx.x;

    float *block_out = &out[50 * 50 * 4 * blockIdx.x];

    for(int i = 0; i < 100; i++) {
        float x = in_x[start + i];
        float y = in_y[start + i];
        int w_idx = (start + i) * 4;

        int xbin = (int) (((x - -10.0) / 20.0) * 50);
        int ybin = (int) (((y - -10.0) / 20.0) * 50);

        if (0 <= xbin && xbin < 50 && 0 <= ybin && ybin < 50) {
            for(int c = 0; c < 4; c++) {
                atomicAdd(&block_out[(ybin * 50 + xbin) * 4 + c], in_w[w_idx + c]);
            }
        }
    }
}

索引好像有问题，不过我之前没做过太多纯CUDA，所以说不清是什么。这就是我认为等效的python：

def slow_hist(in_x, in_y, in_w, out, blockx, blockdimx, threadx):
    start = blockx * blockdimx + threadx

    block_out_addr = args['xres'] * args['yres'], args['num_chans'] * blockx

    for i in range(args['length']):
        x = in_x[start + i]
        y = in_y[start + i]
        w_idx = (start + i) * args['num_chans']

        xbin = int(((x - args['xmin']) / args['xptp']) * args['xres'])
        ybin = int(((y - args['ymin']) / args['yptp']) * args['yres'])

        if 0 <= xbin < args['xres'] and 0 <= ybin < args['yres']:
            for c in range(args['num_chans']):
                out[(ybin * args['xres'] + xbin) * args['num_chans'] + c] += in_w[w_idx + c]

所有代码都是可见的，包括这些图片，at the Github page of this notebook（这个单元格在底部）。

我在这个 CUDA 代码中做错了什么？我已经尝试了很多小调整（将 atomicAdd 地址跨越 1、4、8、16，转置输出等），但似乎我遗漏了一些微妙的东西，可能是指针算法的工作方式。任何帮助将不胜感激。

【问题讨论】：

对于 n 个样本，in_x 和 in_y 的形状为 (n,)，而 in_w 的形状为 (n,4)。 out 具有形状（num_blocks，yres，xres，4）。我的目标是为每个 CUDA 块分配足够的空间，以便有自己的 (yres, xres, 4) 区域来写入（因此原子添加不会相互阻塞），然后我在轴 0 上求和以获得最终的直方图

标签： python numpy cuda pycuda histogram2d

【解决方案1】：

为 CUDA 部分的输出数组分配的数组使用 Numpy 的默认 float64 而不是 float32，因此内存是预期的两倍。这是新的直方图输出：

我仍然非常感谢 cmets 或帮助解释为什么这些直方图彼此如此不同的答案。

【讨论】：