使用 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 时，减和的结果就会出错——这仅仅是由于我在开始时所说的，即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;
    }

然而，令我惊讶的是，所有这些都只会使减和结果始终为零。

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

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

Answer 1

添加error checking，在返回总和为零的情况下发现你的内核根本没有运行。

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

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

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

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

Answer 2

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

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 错误检查，您就会知道这一点。