使用 CUDA 对 unsigned char 数组求和：如何使用 uchars 正确累积线程块结果？

Question

依靠 NVIDIA 的示例和在 SO 上找到的一些好的建议，我一直在设法实现我的项目所需的一些数组缩减内核。但是，一个特定的问题仍然给我带来麻烦。即，如何正确地对无符号字符数组进行求和缩减 (uchar)。

因为uchar可以保存从0到255的值，当然线程块不能累积每个线程块大于255的值。我的直觉是，尽管输入是 uchar，但这只是在 int 中收集减和函数内的总和的情况。但是，它不起作用。

让我详细展示一下我有什么。下面是我的内核，用于对 uchar 的数组求和 - 它是 NVIDIA 示例中著名的 reduce6 函数的略微修改版本：

template <class T, unsigned int blockSize>
__global__ void reduce6(int n, T *g_idata, int *g_odata)
{
    extern __shared__ T sdata[];

    // perform first level of reduction,
    // reading from global memory, writing to shared memory
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockSize * 2 + threadIdx.x;
    unsigned int gridSize = blockSize * 2 * gridDim.x;

    int mySum = 0;

    // we reduce multiple elements per thread.  The number is determined by the
    // number of active thread blocks (via gridDim).  More blocks will result
    // in a larger gridSize and therefore fewer elements per thread
    while (i < n)
    {
        mySum += g_idata[i];
        // ensure we don't read out of bounds
        if (i + blockSize < n) mySum += g_idata[i + blockSize];
        i += gridSize;
    }

    // each thread puts its local sum into shared memory
    sdata[tid] = mySum;
    __syncthreads();


    // do reduction in shared mem
    if ((blockSize >= 512) && (tid < 256))
        sdata[tid] = mySum = mySum + sdata[tid + 256];
    __syncthreads();

    if ((blockSize >= 256) && (tid < 128))
        sdata[tid] = mySum = mySum + sdata[tid + 128];
     __syncthreads();

    if ((blockSize >= 128) && (tid <  64))
        sdata[tid] = mySum = mySum + sdata[tid + 64];
    __syncthreads();

    // fully unroll reduction within a single warp
    if ((blockSize >= 64) && (tid < 32))
        sdata[tid] = mySum = mySum + sdata[tid + 32];
    __syncthreads();

    if ((blockSize >= 32) && (tid < 16))
        sdata[tid] = mySum = mySum + sdata[tid + 16];
    __syncthreads();

    if ((blockSize >= 16) && (tid <  8))
        sdata[tid] = mySum = mySum + sdata[tid + 8];
    __syncthreads();

    if ((blockSize >= 8) && (tid <  4))
        sdata[tid] = mySum = mySum + sdata[tid + 4];
    __syncthreads();

    if ((blockSize >= 4) && (tid <  2))
        sdata[tid] = mySum = mySum + sdata[tid + 2];
    __syncthreads();

    if ((blockSize >= 2) && (tid <  1))
        mySum += sdata[tid + 1];
    __syncthreads();

    // write result for this block to global mem
    if (tid == 0)  atomicAdd(g_odata, mySum);
}

当使用 reduce6<uchar, Blocksize> 调用此类内核时 Blocksize*num.threads = 256，一切正常，求和得到正确的结果。每当这样的比率不是 256 时，总和减少的结果就会出错 - 这仅仅是由于我在 bebinning 中所说的，即 uchar 不能累加大于 255 的值。

对我来说，直观的解决方案是简单地更改行：

extern __shared__ T sdata[];

收件人：

extern __shared__ int sdata[];

由于 sdata 是在求和内核中创建的共享数组，我认为它可以是任何类型，因此可以正确地累加线程块求和产生的任何值。也许，为了确保这一点，我什至编写了 while 循环，将收入数据显式转换为 int:

    while (i < n)
    {
        mySum += (int)g_idata[i];
        // ensure we don't read out of bounds
        if (i + blockSize < n) mySum += (int)g_idata[i + blockSize];
        i += gridSize;
    }

然而，令我惊讶的是，所有tha仅使和减少结果始终为零。

我错过了什么？我怎样才能改变这样的内核，以便传递的 uchar 数组可以用任意数量的线程块和线程正确地减少总和？

如果需要，可以在以下位置找到完整的示例代码：http://pastebin.com/nq1VRJCs

Answer 1

问题很可能出现在您未在此处显示的代码中：

int smemSize = (threads <= 256) ?
    2 * threads * sizeof(uchar) :
    threads * sizeof(uchar);
reduce6<uchar, 256> <<< dimGrid, dimBlock, smemSize>>>
    (DATA_LENGTH, d_data1, d_reduced);

如果您在内核中更改了共享内存缓冲区的类型，则还必须在内核调用中更改其大小。

在这种情况下，结果为零的原因是内核永远不会运行完成。如果您运行使用 cuda-memcheck 的代码，或者添加了适当的运行时 API 错误检查，您就已经知道了。

Answer 2

添加error checking发现在返回和为零的情况下你的内核根本不是运行

运行您在 cuda-memcheck 下的程序发现当您更改共享内存数组的类型或将块大小增加到超出范围时，您的内核正在产生越界共享内存访问256.

然后看到你在 pastebin 上的完整代码中的大小计算对于大于 256 的块大小是不正确的，或者当它显式引用共享内存数组的类型时没有与内核中使用的实际类型一起调整:

int smemSize = (threads <= 256) ?
    2 * threads * sizeof(uchar) :
    threads * sizeof(uchar);

您的内核代码本身没有这种大小写区分。

使用 CUDA 对 unsigned char 数组求和：如何使用 uchars 正确累积线程块结果？

Sum-reducing an array of unsigned char with CUDA: how to properly accumulate thread-block results with uchars?

arrays

cuda

char

reduction

unsigned-char