并行缩减 #5 展开最后一个扭曲答案

【问题标题】：Parallel reduction #5 unroll the last warp并行缩减 #5 展开最后一个扭曲
【发布时间】：2017-07-21 11:32:05
【问题描述】：

我正在研究 Mark Harris 的著名幻灯片中的缩减。特别是，我实施了优化步骤 #5，但我得到了错误的结果：17 而不是 41。我使用了幻灯片中显示的相同数字序列。我在代码中省略了共享数组的“extern”，因为内核和主机代码在同一个 .cu 文件中。

#include <stdio.h>
#include <cuda_runtime.h>

#define THREAD_PER_BLOCK 16

__global__ void reduce5(int *g_idata, int *g_odata) {
    __shared__ int sdata[THREAD_PER_BLOCK];
    // perform first level of reduction,
    // reading from global memory, writing to shared memory
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;
    sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
    __syncthreads();
    // do reduction in shared mem
    for (unsigned int s=blockDim.x/2; s>32; s>>=1) {
        if (tid < s) sdata[tid] += sdata[tid + s];
        __syncthreads();
    }

    if (tid < 32)
    {
        sdata[tid] += sdata[tid + 32];
        sdata[tid] += sdata[tid + 16];
        sdata[tid] += sdata[tid + 8];
        sdata[tid] += sdata[tid + 4];
        sdata[tid] += sdata[tid + 2];
        sdata[tid] += sdata[tid + 1];
    }
    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

int main()
{
    int inputLength=16;
    int hostInput[16]={10,1,8,-1,0,-2,3,5,-2,-3,2,7,0,11,0,2};
    int hostOutput=0;
    int *deviceInput;
    int *deviceOutput;

    cudaMalloc((void **)&deviceInput, inputLength * sizeof(int));
    cudaMalloc((void **)&deviceOutput, sizeof(int));

    cudaMemcpy(deviceInput, hostInput, inputLength * sizeof(int),cudaMemcpyHostToDevice);

    reduce5<<<1,16>>>(deviceInput, deviceOutput);

    cudaDeviceSynchronize();

    cudaMemcpy(&hostOutput, deviceOutput,sizeof(int), cudaMemcpyDeviceToHost);

    printf("%d\n",hostOutput);

    cudaFree(deviceInput);
    cudaFree(deviceOutput);

    return 0;
}

【问题讨论】：

标签： cuda

【解决方案1】：

在此代码中THREAD_PER_BLOCK 必须是 32 的倍数，且至少为 64。输入数据的长度也必须是块和网格大小乘积的两倍。

您看不到它（因为您没有执行任何类型的错误检查），但由于越界共享内存和全局内存访问，线程和扭曲减少将失败。

还要注意extern __shared__与内核和其他代码是否在同一个文件中无关。这表示该变量的共享内存将在运行时动态分配，而不是在编译时静态分配。分配的大小作为内核启动语法中的第三个参数传递。

【讨论】：

我执行了以下更改 THREAD_PER_BLOCK 64 和 reduce5>>(deviceInput, deviceOutput);但我总是得到 17。

【解决方案2】：

我遇到了同样的问题，发现如果变量没有声明为volatile，线程实际上并没有同步。

在声明sdata时只需添加volatile即可解决问题。请参考我的帖子：cuda Threads in A Warp appear to be not in synchronization

【讨论】：

非常感谢 :) 它解决了我的代码中的类似问题 :)