15

我正在尝试减少 CUDA，我真的是一个新手。我目前正在研究来自 NVIDIA 的示例代码。

我想我真的不确定如何设置块大小和网格大小，特别是当我的输入数组大于（512 X 512）大于单个块大小时。

这是代码。

template <unsigned int blockSize>
__global__ void reduce6(int *g_idata, int *g_odata, unsigned int n)
{
    extern __shared__ int sdata[];
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockSize*2) + tid;
    unsigned int gridSize = blockSize*2*gridDim.x;
    sdata[tid] = 0;

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

    __syncthreads();

    if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }
    if (blockSize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); }
    if (blockSize >= 128) { if (tid < 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads(); }

    if (tid < 32) 
    {
        if (blockSize >= 64) sdata[tid] += sdata[tid + 32];
        if (blockSize >= 32) sdata[tid] += sdata[tid + 16];
        if (blockSize >= 16) sdata[tid] += sdata[tid + 8];
        if (blockSize >= 8) sdata[tid] += sdata[tid + 4];
        if (blockSize >= 4) sdata[tid] += sdata[tid + 2];
        if (blockSize >= 2) sdata[tid] += sdata[tid + 1];
    }

    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

但是，在我看来，g_odata[blockIdx.x]保存了所有块的部分总和，并且，如果我想获得最终结果，我需要对g_odata[blockIdx.x]数组中的所有项求和。

我想知道：是否有一个内核来进行整个求和？还是我在这里误解了一些事情？如果有人能用这个来教育我，我将不胜感激。非常感谢。

4

3 回答 3

15

你的理解是正确的。此处展示的缩减以存储在全局内存中的一系列块和而告终。

要将所有这些块总和加在一起，需要某种形式的全局同步。您必须等到所有块都完成后才能将它们的总和相加。此时您有许多选择，其中一些是：

  1. 在主内核之后启动一个新内核以将块和加在一起
  2. 在主机上添加块总和
  3. 在主内核的末尾使用原子将块总和相加
  4. 使用类似threadfence reduction的方法在主内核中将块总和相加。
  5. 使用CUDA 协作组在内核代码中放置网格范围的同步。在网格范围同步（可能在一个块中）之后对块总和进行求和。

如果您搜索 CUDA 标签，您可以找到所有这些示例，并讨论它们的优缺点。要查看您发布的主内核如何用于完全归约，请查看并行归约示例代码。

于 2014-04-08T13:58:18.057 回答
8

Robert Crovella 已经回答了这个问题，主要是关于理解而不是表现。

但是，对于所有遇到这个问题的人，我只想强调CUB提供了块减少功能。下面，我提供了一个关于如何使用 CUB 的简单示例BlockReduce。

#include <cub/cub.cuh>
#include <cuda.h>

#include "Utilities.cuh"

#include <iostream>

#define BLOCKSIZE   32

const int N = 1024;

/**************************/
/* BLOCK REDUCTION KERNEL */
/**************************/
__global__ void sum(const float * __restrict__ indata, float * __restrict__ outdata) {

    unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;

    // --- Specialize BlockReduce for type float. 
    typedef cub::BlockReduce<float, BLOCKSIZE> BlockReduceT; 

    // --- Allocate temporary storage in shared memory 
    __shared__ typename BlockReduceT::TempStorage temp_storage; 

    float result;
    if(tid < N) result = BlockReduceT(temp_storage).Sum(indata[tid]);

    // --- Update block reduction value
    if(threadIdx.x == 0) outdata[blockIdx.x] = result;

    return;  
}

/********/
/* MAIN */
/********/
int main() {

    // --- Allocate host side space for 
    float *h_data       = (float *)malloc(N * sizeof(float));
    float *h_result     = (float *)malloc((N / BLOCKSIZE) * sizeof(float));

    float *d_data;      gpuErrchk(cudaMalloc(&d_data, N * sizeof(float)));
    float *d_result;    gpuErrchk(cudaMalloc(&d_result, (N / BLOCKSIZE) * sizeof(float)));

    for (int i = 0; i < N; i++) h_data[i] = (float)i;

    gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice));

    sum<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_data, d_result);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaMemcpy(h_result, d_result, (N / BLOCKSIZE) * sizeof(float), cudaMemcpyDeviceToHost));

    std::cout << "output: ";
    for(int i = 0; i < (N / BLOCKSIZE); i++) std::cout << h_result[i] << " ";
    std::cout << std::endl;

    gpuErrchk(cudaFree(d_data));
    gpuErrchk(cudaFree(d_result));

    return 0;
}

在此示例中，创建了一个长度数组，N结果是32连续元素的总和。所以

result[0] = data[0] + ... + data[31];
result[1] = data[32] + ... + data[63];
....
于 2015-07-30T17:32:42.537 回答
7

为了更好地了解这个主题，您可以查看 NVIDIA 的这个 pdf，它以图形方式解释了您在代码中使用的所有策略。

于 2015-11-18T13:23:56.657 回答