CUDA 内核的线程顺序执行

Question

我有两个内核按顺序处理一些数据（只用一个线程启动）。我想将两者结合起来，这样我就可以用一个内核来启动两个线程。这样做之后，我期望获得 max(kernel1, kernel2) 的执行时间，但我得到的是两个执行时间的总和。我将问题缩小到类似下面的代码。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include<iostream>
#include<string>
#include<vector>
#include<random>
#include<functional>
#include<algorithm>
#include<iterator>

__global__ void dummyKernel(const float *d_data_Re, const float *d_data_Im,
    float *d_out_Re, float *d_out_Im, const int dataLen) {
    int i{ threadIdx.x };
    if (i == 0) {
        printf("Thread zero started \n");
        for (int j{}; j < 1000000; j++)
            d_out_Re[j%dataLen] = sqrtf(2) + d_data_Re[j%dataLen] * (j % 4 == 1);
        printf("Thread zero finished \n");
    }
    else if (i == 1) {
        printf("Thread one started \n");
        for (int j{}; j < 1000000; j++)
            d_out_Im[j%dataLen] = sqrtf(2) + d_data_Im[j%dataLen] * (j % 4 == 1);
        printf("Thread one finished \n");
    }
}

__global__ void dummyKernel2(const float *d_data_Re, const float *d_data_Im,
    float *d_out_Re, float *d_out_Im, const int dataLen) {
    int i{ threadIdx.x };
    //if (i == 0) {
        printf("Thread zero started \n");
        for (int j{}; j < 1000000; j++)
            d_out_Re[j%dataLen] = sqrtf(2) + d_data_Re[j%dataLen] * (j % 4 == 1);
        printf("Thread zero finished \n");
    //}
    //else if (i == 1) {
    //  printf("Thread one started \n");
    //  for (int j{}; j < 1000000; j++)
    //      d_out_Im[j%dataLen] = sqrtf(2) + d_data_Im[j%dataLen] * (j % 4 == 1);
    //  printf("Thread one finished \n");
    //}
}

int main()
{
    cudaError_t cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        return 1;
    }

    const int sizeOfFrame = 2 * 1024 * 1024;
    std::vector<float> data_re(sizeOfFrame), data_im;
    //random number generator
    std::uniform_real_distribution<float> distribution(0.0f, 2.0f); //Values between 0 and 2
    std::mt19937 engine; // Mersenne twister MT19937
    auto generator = std::bind(distribution, engine);
    std::generate_n(data_re.begin(), sizeOfFrame, generator);
    std::copy(data_re.begin(), data_re.end(), std::back_inserter(data_im));
    //

    float *d_data_re, *d_data_im;
    cudaMalloc(&d_data_re, sizeOfFrame * sizeof(float));
    cudaMalloc(&d_data_im, sizeOfFrame * sizeof(float));
    cudaMemcpy(d_data_re, data_re.data(), sizeOfFrame * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_data_im, data_im.data(), sizeOfFrame * sizeof(float), cudaMemcpyHostToDevice);

    float *d_pll_out_re, *d_pll_out_im;
    cudaMalloc(&d_pll_out_re, sizeOfFrame * sizeof(float));
    cudaMalloc(&d_pll_out_im, sizeOfFrame * sizeof(float));

    dummyKernel << <1, 2 >> >(d_data_re, d_data_im,
        d_pll_out_re, d_pll_out_im, sizeOfFrame);
    cudaDeviceSynchronize();

    // cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}

顺便说一下，我从 this 问题的答案中得到了随机数生成器的代码。所以，dummyKernel 没有做任何有用的事情，我只是想要一个需要相对较长时间才能完成的内核。如果启动 dummyKernel，输出的顺序将是“线程 0 已启动”、“线程 0 已完成”、“线程 1 已启动”、“线程 1 已完成”。顺序的。但是如果启动 dummyKernel2，输出的顺序将是“线程零启动”、“线程零启动”、“线程零完成”、“线程零完成”，执行时间几乎是 dummyKernel 的一半。我不明白这种行为以及我使用的 if-else 的影响。 OS: Windows 10, GTX 1050 Ti, CUDA Driver/Runtime 版本: 11.1/10.1.

Answer 1

每个 Cuda 多处理器都有执行单元（每个用于 int、float、特殊函数等）。那些像管道一样工作，需要几个周期来完成一个计算，但是在每个周期中可以插入一个新的计算（=计划）并且在管道的不同阶段同时处理多个计算。

一个块中的 32 个线程组（warps）被同时调度到相同的指令（相同的周期或通常是两个周期，具体取决于架构上有多少执行和数据路径资源可用以及该指令需要），连同一个位域，说明应该为哪些线程主动执行该指令。如果 warp 的某些线程将 if 子句评估为 false，则它们会暂时停用。或者某些线程可能已经退出内核。

效果是，如果 32 个线程束发散（分支不同），则 32 个线程中的每一个的每个执行路径都必须运行通过（每个路径都停用了一些线程）。出于性能原因，应该避免这种情况，因为计算资源仍然保留。来自不同 warp 的线程没有这种相互依赖性。算法的结构应该考虑到这一点。

Volta 引入了独立线程调度。每个线程都有自己的指令计数器（并管理一个单独的函数调用栈）。但是调度程序仍然会为活动线程安排 32 个线程组（warps）和位域。改变的是调度程序可以交错发散路径。如果可用的执行单元或内存延迟更合适，它可以执行 CCCIIIEEEICCC pre-Volta（指令：C=common，I=if branch，e=else branch）而不是执行 CCCIIIEEIIECCC。作为程序员，必须小心，因为不能再假设线程没有发散，即使在执行相同的指令时也是如此。这就是为什么 __syncwarp 被引入并且所有类型的合作函数（例如洗牌指令）都有一个同步变体。尽管如此（尽管我们不能确定线程是否分离）仍然必须以一种方式进行编程，以便所有 32 个线程可以一起工作，如果同步执行，尤其是对于联合内存访问。在每个可能发散的指令之后放置 __syncwarp 有助于确保收敛。（但是做性能分析）。

独立线程调度也是为什么必须在 RTX 3080 上正确调用 __syncthreads 的原因 - 每个线程都参与。您在评论中提到的死锁情况的典型纠正解决方案是关闭 if 子句，同步所有线程并打开一个与前一个条件相同的新 if 子句。

CUDA 内核的线程顺序执行

Threads of a CUDA kernel execute sequentially

cuda