c++ - CUDA 缩减，大阵列的方法

Question

我有以下“Frankenstein”和减少代码，部分来自 common CUDA reduction slices ，部分来自 CUDA 示例。

    __global__ void  reduce6(float *g_idata, float *g_odata, unsigned int n)
{
    extern __shared__ float 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;
    sdata[tid] = 0;
    float mySum = 0;   

    while (i < n) { 
        sdata[tid] += g_idata[i] + g_idata[i+MAXTREADS]; 
        i += gridSize; 
    }
   __syncthreads();


    // do reduction in shared mem
    if (tid < 256)
        sdata[tid] += sdata[tid + 256];
    __syncthreads();

    if (tid < 128)
        sdata[tid] +=  sdata[tid + 128];
     __syncthreads();

    if (tid <  64)
       sdata[tid] += sdata[tid +  64];
    __syncthreads();


#if (__CUDA_ARCH__ >= 300 )
    if ( tid < 32 )
    {
        // Fetch final intermediate sum from 2nd warp
        mySum = sdata[tid]+ sdata[tid + 32];
        // Reduce final warp using shuffle
        for (int offset = warpSize/2; offset > 0; offset /= 2) 
            mySum += __shfl_down(mySum, offset);
    }
    sdata[0]=mySum;
#else

    // fully unroll reduction within a single warp
    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];
    }
#endif
    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
  }

我将使用它来减少 Tesla k40 GPU 上展开的大尺寸数组(例如 512^3 = 134217728 = n)。

我对 blockSize 变量及其值有一些疑问。

从这里开始，我将尝试解释我对它是如何工作的理解(无论是对还是错):

我选择的 blockSize 越大，这段代码执行得越快，因为它在整个循环中花费的时间更少，但它不会完成整个数组的减少，但会返回一个更小的大小为 dimBlock.x 的数组，对吗？如果我使用 blockSize=1，此代码将在 1 调用缩减值时返回，但它会非常慢，因为它几乎没有利用 CUDA 的功能。因此，我需要多次调用缩减内核，每次都使用较小的 blokSize，并减少之前调用 reduce 的结果，直到我到达最小点。

类似(伪代码)

blocks=number; //where do we start? why?
while(not the min){

    dim3 dimBlock( blocks );
    dim3 dimGrid(n/dimBlock.x);
    int smemSize = dimBlock.x * sizeof(float);
    reduce6<<<dimGrid, dimBlock, smemSize>>>(in, out, n);

    in=out;

    n=dimGrid.x; 
    dimGrid.x=n/dimBlock.x; // is this right? Should I also change dimBlock?
}

我应该从哪个值开始？我想这取决于 GPU。 Tesla k40 应该是哪个值(只是为了让我了解如何选择这些值)？

我的逻辑有问题吗？怎么办？

Answer 1

有一个 CUDA 工具可以为您获得良好的网格和 block 大小:Cuda Occupancy API .

响应 “我选择的 blockSize 越大，这段代码执行得越快” -- 不一定，因为你想要最大 occupancy 的大小(事件扭曲与可能事件扭曲总数的比率)。

有关更多信息，请参阅此答案 How do I choose grid and block dimensions for CUDA kernels? .

最后，对于支持 Kelper 或更高版本的 Nvidia GPU，有 shuffle intrinsics使减少更容易和更快。这是一篇关于如何使用随机播放内在函数的文章:Faster Parallel Reductions on Kepler .

选择线程数的更新:

如果使用最大线程数会降低寄存器的使用效率，您可能不想使用它。来自入住链接:

为了计算占用率，每个线程使用的寄存器数量是关键因素之一。例如，计算能力为 1.1 的设备每个多处理器有 8,192 个 32 位寄存器，最多可以同时驻留 768 个线程(24 个线程 x 每个线程 32 个线程)。这意味着在其中一个设备中，要使多处理器拥有 100% 的占用率，每个线程最多可以使用 10 个寄存器。然而，这种确定寄存器计数如何影响占用率的方法没有考虑寄存器分配粒度。例如，在计算能力为 1.1 的设备上，具有 128 线程 block 的内核每个线程使用 12 个寄存器导致占用率为 83%，每个多处理器有 5 个事件的 128 线程 block ，而具有 256 线程 block 的内核每个线程使用相同的 12 个寄存器会导致占用率为 66%，因为在一个多处理器上只能驻留两个 256 线程 block 。

所以我的理解是，由于寄存器的分配方式，增加线程数可能会限制性能。然而，情况并非总是如此，您需要自己进行计算(如上述语句)以确定每个 block 的最佳线程数。