CUDA 并行加速基础之 Reduce 和 Scan 的实现

GPU加速求和算法

最新推荐文章于 2025-10-30 15:05:08 发布

原创

最新推荐文章于 2025-10-30 15:05:08 发布 · 3.1k 阅读

10 ·

CC 4.0 BY-SA版权

文章标签：

#cuda #并行加速 #并行算法

本文介绍了如何利用GPU的并行性来加速求和运算，详细解释了Reduce和Scan两种并行算法的原理，并提供了CUDA实现代码。

前言

我们知道硬件拥有其独特的并行性，为了发挥这一特色。我们要将平时串行执行的程序用并行性算法重新改写才能充分发挥 GPU 的优势。

实例：做求和：1+2+3+4+···

为了做这样一个累加和的加速，有两种简单的实现方法，分别是 Redece 进行归约（二分），或者是用 Scan 通过控制步长进行扫描求和。

Reduce

在这里插入图片描述
如上图所示为了并行执行累加，我们要构造出一些线程，每个线程并行工作，从而达到加速的目的。
Reduce 的原理如上图所示。我们构建1024个线程块，每个线程块中含有1024个线程，每个block中的线程（每次迭代只有上次迭代一半的线程数）并行的去计算求和。

求和规则是：每次迭代中左半边的线程分别加上右半边的线程分配得到的数，即只有一半的线程在工作，如此循环往复，最终0号线程得到的结果就是该线程块所有数的和。最后再运行最后一个线程块，将之前得到的线程块和当作输入数据，进行最后一次并行计算，从而得到最终和。

实现代码如下，通过两种访存方式进行实现：

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

__global__ void global_reduce_kernel(float * d_out, float * d_in)
{
   
   
    int myId = threadIdx.x + blockDim.x * blockIdx.x;
    int tid  = threadIdx.x;

    // do reduction in global mem
    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1)
    {
   
   
        if (tid < s)
        {
   
   
            d_in[myId] += d_in[myId + s];
        }
        __syncthreads();        // make sure all adds at one stage are done!
    }

    // only thread 0 writes result for this block back to global mem
    if (tid == 0)
    {
   
   
        d_out[blockIdx.x] = d_in[myId];
    }
}

__global__ void shmem_reduce_kernel(float * d_out, const float * d_in)
{
   
   
    // sdata is allocated in the kernel call: 3rd arg to <<<b, t, shmem>>>
    extern __shared__ float sdata[];

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

    // load shared mem from global mem
    sdata[tid] = d_in[myId];
    __syncthreads();            // make sure entire block is loaded!

    // do reduction in shared mem
    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1)
    {
   
   
        if (tid < s)
        {
   
   
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();        // make sure all adds at one stage are done!
    }

    // only thread 0 writes result for this block back to global mem
    if (tid == 0)
    {
   
   
        d_out[blockIdx.x] = sdata[0];
    }
}

void reduce(float * d_out, float * d_intermediate