GPU编程 CUDA C++ 用非默认CUDA流重叠多个核函数在GPU中并发执行

核函数外部并行（可分为5类）：

1. 核函数计算与数据传输之间的并行

2. 主机计算与数据传输之间的并行

3. 不同的数据传输之间的并行

4. 核函数计算与主机计算之间的并行

5. 不同核函数之间的并行

以下代码示例为第5类情况分别对RTX3080Ti在同时并行1-128个核函数的情况进行测试：

kernel-kernel_1024.cu

#include "error.cuh"
#include <math.h>
#include <stdio.h>

#ifdef USE_DP
    typedef double real;
#else
    typedef float real;
#endif

const int NUM_REPEATS = 10;
const int N1 = 1024;   //每个核函数使用1024个线程，后改为2048进行第二次测试
const int MAX_NUM_STREAMS = 128;  //分别递增测试到最大同时并行128个kernel函数
const int N = N1 * MAX_NUM_STREAMS;
const int M = sizeof(real) * N;
const int block_size = 128;
const int grid_size = (N1 - 1) / block_size + 1;
cudaStream_t streams[MAX_NUM_STREAMS];

void timing(const real *d_x, const real *d_y, real *d_z, const int num);

int main(void)
{
    real *h_x = (real*) malloc(M);
    real *h_y = (real*) malloc(M);
    for (int n = 0; n < N; ++n)
    {
        h_x[n] = 1.23;
        h_y[n] = 2.34;
    }

    real *d_x, *d_y, *d_z;
    CHECK(cudaMalloc(&d_x, M));
    CHECK(cudaMalloc(&d_y, M));
    CHECK(cudaMalloc(&d_z, M));
    CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice));

    for (int n = 0 ; n < MAX_NUM_STREAMS; ++n)
    {
        CHECK(cudaStreamCreate(&(streams[n])));
    }

    for (int num = 1; num <= MAX_NUM_STREAMS; ++num)
    {
        timing(d_x, d_y, d_z, num);   //num是流编号，最大128
    }

    for (int n = 0 ; n < MAX_NUM_STREAMS; ++n)
    {
        CHECK(cudaStreamDestroy(streams[n]));
    }

    free(h_x);
    free(h_y);
    CHECK(cudaFree(d_x));
    CHECK(cudaFree(d_y));
    CHECK(cudaFree(d_z));
    return 0;
}

void __global__ add(const real *d_x, const real *d_y, real *d_z)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    if (n < N1)       //每个核函数使用N1个线程
    {
        for (int i = 0; i < 100000; ++i)   //注意，这个for循环是故意让每个线程重复100000加法，拖延时间
        {
            d_z[n] = d_x[n] + d_y[n];
        }
    }
}

void timing(const real *d_x, const real *d_y, real *d_z, const int num)   //num是流编号
{
    float t_sum = 0;
    float t2_sum = 0;

    for (int repeat = 0; repeat <= NUM_REPEATS; ++repeat)
    {
        cudaEvent_t start, stop;
        CHECK(cudaEventCreate(&start));
        CHECK(cudaEventCreate(&stop));
        CHECK(cudaEventRecord(start));
        cudaEventQuery(start);

        for (int n = 0; n < num; ++n)   //num是流编号
        {
            int offset = n * N1;        //每个核函数使用N1个线程
            add<<<grid_size, block_size, 0, streams[n]>>>
            (d_x + offset, d_y + offset, d_z + offset);
        }
 
        CHECK(cudaEventRecord(stop));
        CHECK(cudaEventSynchronize(stop));
        float elapsed_time;
        CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));

        if (repeat > 0)
        {
            t_sum += elapsed_time;
            t2_sum += elapsed_time * elapsed_time;
        }

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

    const float t_ave = t_sum / NUM_REPEATS;
    const float t_err = sqrt(t2_sum / NUM_REPEATS - t_ave * t_ave);
    printf("%g\n", t_ave);
}

头文件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 kernel-kernel_1024.cu -o kernel-kernel_1024
./kernel-kernel_1024 > k-k_1024.dat

nvcc kernel-kernel_2048.cu -o kernel-kernel_2048
./kernel-kernel_2048 > k-k_2048.dat

分别对 k-k_1024.dat 和 k-k_2048.dat 两个结果文件进行绘图，结果如下：

 

可以看出，我使用的 RTX3080Ti 在使用32个SM（流多处理器）以下，每个SM常驻1024个线程以下时，计算时间不变。（因为在横坐标CUDA stream低于32时，每个SM驻1024个线程时间均为4ms；在横坐标CUDA stream低于16时，每个SM驻2048个线程时间均为4ms） 

然而我的 RTX3080Ti 理论结果应该是有80个SM（流多处理器），每个SM理论上可以常驻1536个线程。

$ ./query 
Device id:                                 0
Device name:                               NVIDIA GeForce RTX 3080 Ti
Compute capability:                        8.6
Amount of global memory:                   11.7689 GB
Amount of constant memory:                 64 KB
Maximum grid size:                         2147483647 65535 65535
Maximum block size:                        1024 1024 64
Number of SMs:                             80
Maximum amount of shared memory per block: 48 KB
Maximum amount of shared memory per SM:    100 KB
Maximum number of registers per block:     64 K
Maximum number of registers per SM:        64 K
Maximum number of threads per block:       1024
Maximum number of threads per SM:          1536
(base)