CUDA 中的块减少
Posted
技术标签:
【中文标题】CUDA 中的块减少【英文标题】:Block reduction in CUDA 【发布时间】:2014-05-21 06:24:41 【问题描述】:我正在尝试减少 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]
数组中的所有项求和。
我想知道:是否有一个内核来进行整个求和?还是我在这里误解了一些事情?如果有人能用这个来教育我,我将不胜感激。非常感谢。
【问题讨论】:
另外请注意上面代码中的__shared__
数据应该是volatile
,否则不能保证正确的最终结果。可以在@Robert 提供的链接中看到。
【参考方案1】:
你的理解是正确的。 here 所展示的缩减最终会在全局内存中存储一系列块和。
要将所有这些块总和相加,需要某种形式的全局同步。您必须等到所有块都完成后才能将它们的总和相加。此时您有多种选择,其中一些是:
-
在主内核之后启动一个新内核以将块和相加在一起
在主机上添加块总和
在主内核末尾使用原子将块总和相加
使用threadfence reduction 之类的方法在主内核中将块总和相加。
使用CUDA cooperative groups 在内核代码中放置一个网格范围的同步。在网格范围同步后(可能在一个块中)对块总和进行求和。
如果您搜索 CUDA 标签,您可以找到所有这些的示例,并讨论它们的优缺点。要查看您发布的主内核如何用于完全缩减,请查看parallel reduction sample code。
【讨论】:
【参考方案2】: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];
....
【讨论】:
【参考方案3】:为了更好地了解这个主题,您可以查看 NVIDIA 的 this pdf,它以图形方式解释了您在代码中使用的所有策略。
【讨论】:
以上是关于CUDA 中的块减少的主要内容,如果未能解决你的问题,请参考以下文章