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，这段代码执行得越快，因为它在整个循环中花费的时间会更少，但它不会完成整个数组的减少，但它会return 一个较小的数组dimBlock.x，对吧？如果我使用 blockSize=1 此代码将 return in 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 工具可以为您获得良好的网格和块大小：Cuda Occupancy API。

响应"The bigger I choose blockSize, the faster this code will execute" -- 不一定，因为你想要最大 occupancy 的大小（活动经纱与总经纱的比率可能的活动扭曲数）。

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

最后，对于支持 Kelper 或更高版本的 Nvidia GPU，有 shuffle intrinsics to make reductions easier and faster. Here is an article on how to use the shuffle intrinsics : Faster Parallel Reductions on Kepler。

选择线程数的更新：

如果使用最大线程数会降低寄存器的使用效率，您可能不想使用它。从link上入住：

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

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

CUDA 缩减，大阵列的方法

CUDA reduction, approach for big arrays

c++

cuda

reduction