GPU编程 CUDA C++ 分别使用“线程束内同步函数 __syncwarp( )“、“线程束内洗牌函数 __shfl_down_sync( )”、“线程协作组” 三种方式进行数组归约计算

reduce_3_method.cu

#include "error.cuh"  //CHECK错误检查宏
#include <stdio.h>
#include <cooperative_groups.h>
using namespace cooperative_groups;

#ifdef USE_DP       //编译选项：使用双精度浮点
    typedef double real;
#else               //编译选项：使用单精度浮点
    typedef float real;
#endif

const int NUM_REPEATS = 100;
const int N = 100000000;          //数组中元素个数
const int M = sizeof(real) * N;
const int BLOCK_SIZE = 128;
const unsigned FULL_MASK = 0xffffffff;

void timing(const real *d_x, const int method);

int main(void)
{
    real *h_x = (real *) malloc(M);   //分配主机内存
    for (int n = 0; n < N; ++n)
    {
        h_x[n] = 1.23;       //将数组中所有元素初始化为1.23
    }
    real *d_x;
    CHECK(cudaMalloc(&d_x, M));       //分配GPU显存
    CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));   //将数组从主机内存复制到GPU显存

    printf("\nusing syncwarp:\n");  //使用线程束内同步函数法
    timing(d_x, 0);
    printf("\nusing shfl:\n");      //使用线程束内洗牌函数法
    timing(d_x, 1);
    printf("\nusing cooperative group:\n");    //使用线程协作组法
    timing(d_x, 2);

    free(h_x);     //释放主机内存
    CHECK(cudaFree(d_x));    //释放GPU显存
    return 0;
}

void __global__ reduce_syncwarp(const real *d_x, real *d_y, const int N)
{
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    const int n = bid * blockDim.x + tid;
    extern __shared__ real s_y[];         //使用GPU中共享内存
    s_y[tid] = (n < N) ? d_x[n] : 0.0;
    __syncthreads();

    for (int offset = blockDim.x >> 1; offset >= 32; offset >>= 1)  //block块归约
    {
        if (tid < offset)
        {
            s_y[tid] += s_y[tid + offset];
        }
        __syncthreads();  //线程块同步函数
    }

    for (int offset = 16; offset > 0; offset >>= 1)    //进入束内
    {
        if (tid < offset)
        {
            s_y[tid] += s_y[tid + offset];
        }
        __syncwarp();     //使用束内同步函数（从Nvidia伏特架构开始引入的独立线程调度机制）
    }                     //__syncwarp()比__syncthreads()开销更低

    if (tid == 0)
    {
        atomicAdd(d_y, s_y[0]);  //使用原子函数，避免读写竞争
    }
}

void __global__ reduce_shfl(const real *d_x, real *d_y, const int N)
{
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    const int n = bid * blockDim.x + tid;
    extern __shared__ real s_y[];
    s_y[tid] = (n < N) ? d_x[n] : 0.0;
    __syncthreads();

    for (int offset = blockDim.x >> 1; offset >= 32; offset >>= 1)
    {
        if (tid < offset)
        {
            s_y[tid] += s_y[tid + offset];
        }
        __syncthreads();
    }

    real y = s_y[tid];  //在核函数中定义，不加任何限定符，所以y存放在GPU寄存器中

    for (int offset = 16; offset > 0; offset >>= 1)
    {
        y += __shfl_down_sync(FULL_MASK, y, offset);   //洗牌函数能自动处理“同步”和“读-写竞争”问题
    }

    if (tid == 0)
    {
        atomicAdd(d_y, y);
    }
}

void __global__ reduce_cp(const real *d_x, real *d_y, const int N)
{
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    const int n = bid * blockDim.x + tid;
    extern __shared__ real s_y[];
    s_y[tid] = (n < N) ? d_x[n] : 0.0;
    __syncthreads();

    for (int offset = blockDim.x >> 1; offset >= 32; offset >>= 1)
    {
        if (tid < offset)
        {
            s_y[tid] += s_y[tid + offset];
        }
        __syncthreads();
    }

    real y = s_y[tid];

    thread_block_tile<32> g = tiled_partition<32>(this_thread_block());
                           //模板化
                           //上句使用了线程块分割技术
                           //定义了一个线程“块片”类型的变量g
                           //this_thread_block()：获取当前线程所在的线程块
                           //tiled_partition<32>(this_thread_block())：使用tiled_partition()函数，将当前线程块划分为大小为32的小块。这样做的目的是为了更好地利用线程块中的线程资源。
                           //thread_block_tile<32> g = ...：将划分后的小块存储在g中。thread_block_tile是一个线程块的视图，可以通过它来访问线程块中的元素。

    for (int i = g.size() >> 1; i > 0; i >>= 1)    //使用了“块片”类g中成员函数g.size()获取“块片”的大小
    {
        y += g.shfl_down(y, i);    //使用了“块片”类g中成员函数g.shfl_down()进行归约，换成g.shfl_xor()也可以
                                   //使用线程协作组洗牌函数的核函数和使用线程束的洗牌函数的核函数具有等价的执行效率
    }

    if (tid == 0)
    {
        atomicAdd(d_y, y);
    }
}

real reduce(const real *d_x, const int method)    //将3种方法封装起来
{
    const int grid_size = (N + BLOCK_SIZE - 1) / BLOCK_SIZE;
    const int smem = sizeof(real) * BLOCK_SIZE;

    real h_y[1] = {0};
    real *d_y;
    CHECK(cudaMalloc(&d_y, sizeof(real)));
    CHECK(cudaMemcpy(d_y, h_y, sizeof(real), cudaMemcpyHostToDevice));

    switch (method)
    {
        case 0:
            reduce_syncwarp<<<grid_size, BLOCK_SIZE, smem>>>(d_x, d_y, N);
            break;
        case 1:
            reduce_shfl<<<grid_size, BLOCK_SIZE, smem>>>(d_x, d_y, N);
            break;
        case 2:
            reduce_cp<<<grid_size, BLOCK_SIZE, smem>>>(d_x, d_y, N);
            break;
        default:
            printf("Wrong method.\n");
            exit(1);
    }

    CHECK(cudaMemcpy(h_y, d_y, sizeof(real), cudaMemcpyDeviceToHost));
    CHECK(cudaFree(d_y));

    return h_y[0];
}

void timing(const real *d_x, const int method)   //计时框架，性能测试
{
    real sum = 0;
    
    for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)   //每种方法重复跑100遍
    {
        cudaEvent_t start, stop;
        CHECK(cudaEventCreate(&start));
        CHECK(cudaEventCreate(&stop));
        CHECK(cudaEventRecord(start));
        cudaEventQuery(start);

        sum = reduce(d_x, method);    //被计时代码块

        CHECK(cudaEventRecord(stop));
        CHECK(cudaEventSynchronize(stop));
        float elapsed_time;
        CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
        printf("Time = %g ms.\n", elapsed_time);

        CHECK(cudaEventDestroy(start));
        CHECK(cudaEventDestroy(stop));
    }

    printf("sum = %f.\n", sum);
}

其中头文件error.cuh为CHECK错误检查宏函数：

#pragma once
#include <stdio.h>

#define CHECK(call)                                   \
do                                                    \
{                                                     \
    const cudaError_t error_code = call;              \
    if (error_code != cudaSuccess)                    \
    {                                                 \
        printf("CUDA Error:\n");                      \
        printf("    File:       %s\n", __FILE__);     \
        printf("    Line:       %d\n", __LINE__);     \
        printf("    Error code: %d\n", error_code);   \
        printf("    Error text: %s\n",                \
            cudaGetErrorString(error_code));          \
        exit(1);                                      \
    }                                                 \
} while (0)

编译：

$ nvcc reduce_3_method.cu -o reduce_3_method

运行：

$ ./reduce_3_method

性能测试结果输出：

using syncwarp:
Time = 1.86886 ms.
Time = 1.63546 ms.
Time = 1.63152 ms.
Time = 1.7512 ms.
Time = 1.72413 ms.
Time = 1.71853 ms.
Time = 1.73126 ms.
Time = 1.74672 ms.
Time = 1.71226 ms.
Time = 1.72234 ms.
Time = 1.71088 ms.
Time = 1.72038 ms.
Time = 1.72157 ms.
Time = 1.72541 ms.
Time = 1.72355 ms.
Time = 1.72285 ms.
Time = 1.75555 ms.
Time = 1.71315 ms.
Time = 1.70973 ms.
Time = 1.72528 ms.
Time = 1.7105 ms.
Time = 1.71277 ms.
Time = 2.69322 ms.
Time = 5.66118 ms.
Time = 2.06947 ms.
Time = 2.68797 ms.
Time = 1.98102 ms.
Time = 1.76294 ms.
Time = 1.6241 ms.
Time = 1.78102 ms.
Time = 1.62624 ms.
Time = 1.62438 ms.
Time = 1.88864 ms.
Time = 2.43418 ms.
Time = 1.77187 ms.
Time = 1.6257 ms.
Time = 1.80867 ms.
Time = 1.75219 ms.
Time = 1.6263 ms.
Time = 1.63264 ms.
Time = 1.63219 ms.
Time = 1.69274 ms.
Time = 1.76592 ms.
Time = 1.63005 ms.
Time = 1.63123 ms.
Time = 1.62269 ms.
Time = 1.6233 ms.
Time = 1.63325 ms.
Time = 1.62899 ms.
Time = 2.38752 ms.
Time = 1.82586 ms.
Time = 1.73936 ms.
Time = 1.62227 ms.
Time = 1.6199 ms.
Time = 1.62448 ms.
Time = 1.63142 ms.
Time = 1.69587 ms.
Time = 1.72429 ms.
Time = 1.63709 ms.
Time = 1.6273 ms.
Time = 1.61677 ms.
Time = 1.62512 ms.
Time = 1.63286 ms.
Time = 1.63613 ms.
Time = 1.68253 ms.
Time = 1.63018 ms.
Time = 1.62195 ms.
Time = 2.39059 ms.
Time = 1.88794 ms.
Time = 1.62291 ms.
Time = 1.63046 ms.
Time = 1.63254 ms.
Time = 1.62829 ms.
Time = 1.74877 ms.
Time = 1.62886 ms.
Time = 1.61968 ms.
Time = 1.61984 ms.
Time = 1.62419 ms.
Time = 4.94906 ms.
Time = 1.64096 ms.
Time = 1.62771 ms.
Time = 1.62349 ms.
Time = 1.62621 ms.
Time = 1.62189 ms.
Time = 2.61382 ms.
Time = 1.77274 ms.
Time = 1.75104 ms.
Time = 1.62653 ms.
Time = 1.61824 ms.
Time = 1.62518 ms.
Time = 1.62352 ms.
Time = 1.62461 ms.
Time = 1.62963 ms.
Time = 1.72285 ms.
Time = 1.63264 ms.
Time = 1.63725 ms.
Time = 1.63184 ms.
Time = 1.62883 ms.
Time = 1.62387 ms.
Time = 1.62227 ms.
sum = 123633392.000000.

using shfl:
Time = 1.62742 ms.
Time = 1.62787 ms.
Time = 2.39082 ms.
Time = 1.89862 ms.
Time = 1.6287 ms.
Time = 1.61888 ms.
Time = 1.62483 ms.
Time = 1.62445 ms.
Time = 1.44464 ms.
Time = 1.40582 ms.
Time = 1.49712 ms.
Time = 1.40298 ms.
Time = 1.40282 ms.
Time = 1.39709 ms.
Time = 1.39664 ms.
Time = 1.40288 ms.
Time = 1.40083 ms.
Time = 1.40499 ms.
Time = 1.39882 ms.
Time = 1.39904 ms.
Time = 1.39709 ms.
Time = 1.39862 ms.
Time = 1.9129 ms.
Time = 1.65501 ms.
Time = 1.48547 ms.
Time = 1.40378 ms.
Time = 1.40717 ms.
Time = 1.41021 ms.
Time = 1.49056 ms.
Time = 1.4063 ms.
Time = 1.39818 ms.
Time = 1.39869 ms.
Time = 1.39968 ms.
Time = 1.39782 ms.
Time = 1.4032 ms.
Time = 1.39398 ms.
Time = 1.40483 ms.
Time = 1.40083 ms.
Time = 1.40813 ms.
Time = 1.41552 ms.
Time = 1.3961 ms.
Time = 1.39907 ms.
Time = 1.39475 ms.
Time = 1.39366 ms.
Time = 2.17616 ms.
Time = 1.47789 ms.
Time = 1.40579 ms.
Time = 1.49501 ms.
Time = 1.39933 ms.
Time = 1.40051 ms.
Time = 1.3944 ms.
Time = 1.44355 ms.
Time = 1.3983 ms.
Time = 1.39302 ms.
Time = 1.39635 ms.
Time = 1.39853 ms.
Time = 1.40352 ms.
Time = 1.40877 ms.
Time = 1.40358 ms.
Time = 1.39645 ms.
Time = 1.3953 ms.
Time = 1.40147 ms.
Time = 1.39469 ms.
Time = 1.40301 ms.
Time = 1.40214 ms.
Time = 1.99174 ms.
Time = 1.57549 ms.
Time = 1.47949 ms.
Time = 1.39933 ms.
Time = 1.39827 ms.
Time = 1.39987 ms.
Time = 1.41386 ms.
Time = 1.39405 ms.
Time = 1.4007 ms.
Time = 3.63766 ms.
Time = 1.41165 ms.
Time = 1.3935 ms.
Time = 1.40208 ms.
Time = 1.39322 ms.
Time = 1.39642 ms.
Time = 1.39696 ms.
Time = 1.40122 ms.
Time = 1.39555 ms.
Time = 1.40003 ms.
Time = 1.39277 ms.
Time = 1.93914 ms.
Time = 1.64608 ms.
Time = 1.48074 ms.
Time = 1.40038 ms.
Time = 1.3913 ms.
Time = 1.39072 ms.
Time = 1.53334 ms.
Time = 1.39267 ms.
Time = 1.3967 ms.
Time = 1.39843 ms.
Time = 1.39296 ms.
Time = 1.38973 ms.
Time = 1.39686 ms.
Time = 1.39584 ms.
Time = 1.39206 ms.
sum = 123633392.000000.

using cooperative group:
Time = 1.4225 ms.
Time = 1.39811 ms.
Time = 1.39846 ms.
Time = 1.39347 ms.
Time = 1.39386 ms.
Time = 1.40787 ms.
Time = 1.39242 ms.
Time = 2.20211 ms.
Time = 1.48118 ms.
Time = 1.47914 ms.
Time = 1.39907 ms.
Time = 1.4087 ms.
Time = 1.4431 ms.
Time = 1.76768 ms.
Time = 1.39446 ms.
Time = 1.40038 ms.
Time = 1.39888 ms.
Time = 1.39738 ms.
Time = 1.5113 ms.
Time = 1.39965 ms.
Time = 1.39565 ms.
Time = 1.39293 ms.
Time = 1.39594 ms.
Time = 1.39549 ms.
Time = 1.39517 ms.
Time = 1.39283 ms.
Time = 1.39664 ms.
Time = 1.39642 ms.
Time = 2.2871 ms.
Time = 1.64477 ms.
Time = 6.65958 ms.
Time = 1.48026 ms.
Time = 1.48208 ms.
Time = 1.38784 ms.
Time = 1.39475 ms.
Time = 1.38637 ms.
Time = 1.38499 ms.
Time = 1.3865 ms.
Time = 1.38582 ms.
Time = 1.384 ms.
Time = 7.0721 ms.
Time = 1.39411 ms.
Time = 1.38419 ms.
Time = 1.38269 ms.
Time = 2.42253 ms.
Time = 1.50829 ms.
Time = 1.74698 ms.
Time = 1.45808 ms.
Time = 1.3888 ms.
Time = 1.37322 ms.
Time = 1.38157 ms.
Time = 1.38586 ms.
Time = 1.3831 ms.
Time = 1.39453 ms.
Time = 1.39187 ms.
Time = 1.38352 ms.
Time = 1.51811 ms.
Time = 1.39043 ms.
Time = 1.37955 ms.
Time = 1.3735 ms.
Time = 1.38214 ms.
Time = 1.38182 ms.
Time = 1.37891 ms.
Time = 1.99027 ms.
Time = 1.65075 ms.
Time = 1.46733 ms.
Time = 1.38486 ms.
Time = 1.39027 ms.
Time = 1.38307 ms.
Time = 1.37946 ms.
Time = 1.37859 ms.
Time = 1.38509 ms.
Time = 1.38522 ms.
Time = 1.38061 ms.
Time = 1.48733 ms.
Time = 1.38781 ms.
Time = 1.38054 ms.
Time = 1.37194 ms.
Time = 1.38384 ms.
Time = 1.38579 ms.
Time = 1.38237 ms.
Time = 1.37507 ms.
Time = 1.38342 ms.
Time = 1.3879 ms.
Time = 1.37782 ms.
Time = 1.95101 ms.
Time = 1.64682 ms.
Time = 1.45888 ms.
Time = 1.3823 ms.
Time = 1.38982 ms.
Time = 1.38899 ms.
Time = 1.48445 ms.
Time = 1.45632 ms.
Time = 1.38534 ms.
Time = 1.38275 ms.
Time = 1.37773 ms.
Time = 1.37712 ms.
Time = 1.37978 ms.
Time = 1.38755 ms.
Time = 1.38666 ms.
sum = 123633392.000000.
(base)