规约算法是用于将大量数据聚合为一个值,例如计算数组中所有元素的总和、最大值、最小值等,可以利用到GPU的并行特性。
数组求和算法优化过程:
比如一个求数组和的方法,其方法为:
//reduce 测试
__global__ void reduce1(int* g_idata, int* g_odata)
{
extern __shared__ int sdata[1024];
// each thread loads one element from global to shared mem
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
sdata[tid] = g_idata[i];
__syncthreads();
// do reduction in shared mem
for (unsigned int s = 1; s < blockDim.x; s *= 2) {
if (tid % (2 * s) == 0) {
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}
// write result for this block to global mem
if (tid == 0)
{
g_odata[blockIdx.x] = sdata[0];
}
}
但该方法中的取余操作比较慢,因此改为:
__global__ void reduce2(int* g_idata, int* g_odata)
{
extern __shared__ int sdata[1024];
// each thread loads one element from global to shared mem
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
sdata[tid] = g_idata[i];
__syncthreads();
// do reduction in shared mem
for (unsigned int s = 1; s < blockDim.x; s *= 2) {
int index = 2 * s * tid;
if (index < blockDim.x)
{
sdata[index] += sdata[index + s];
}
__syncthreads();
}
// write result for this block to global mem
if (tid == 0)
{
g_odata[blockIdx.x] = sdata[0];
}
}
在RTX 4080 GPU上,性能提升十分恐怖,效率直接翻倍。
但这个实现仍然有其问题,即bank冲突,这里需要将for循环逆着来,改为:
__global__ void reduce3(int* g_idata, int* g_odata)
{
extern __shared__ int sdata[1024];
// each thread loads one element from global to shared mem
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
sdata[tid] = g_idata[i];
__syncthreads();
// do reduction in shared mem</