CUDA 内核的线程顺序执行答案

【问题标题】：Threads of a CUDA kernel execute sequentiallyCUDA 内核的线程顺序执行
【发布时间】：2022-01-01 10:39:36
【问题描述】：

我有两个按顺序处理一些数据的内核（仅使用一个线程启动）。我想将两者结合起来，这样我就可以使用两个线程来启动一个内核。这样做之后，我期望得到一个 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，输出的顺序将是“线程零开始”、“线程零完成”、“线程一开始”、“线程一完成”。顺序的。但是如果你启动 dummyKernel2，输出的顺序将是“线程零启动”、“线程零启动”、“线程零完成”、“线程零完成”，执行时间几乎是 dummyKernel 的一半。我不理解这种行为以及我使用的 if-else 的效果。操作系统：Windows 10、GTX 1050 Ti、CUDA 驱动程序/运行时版本：11.1/10.1。

【问题讨论】：

所有条件分支都在warp中串行执行。这在 GPU（更具体地说是 SIMT 模型）上是预期的。请阅读：stackoverflow.com/questions/11687500/…（可能还有 CUDA 手册）。
因此尝试启动 33 个线程 (0..32)，让线程 1-31 什么都不做，并使用线程 0 和 32 来查看它是否进行了更改（仅用于测试/教育目的 -这不是一个好的并行设置）
@JérômeRichard 谢谢，现在开始有意义了
@Sebastian 谢谢，经过测试，它解决了我的问题。
在比你的 (Pascal) 更新的 Nvidia GPU (>= Volta) 中，warp 的每个线程可能执行不同的程序代码位置（这称为独立线程调度），但让warp 的线程不会发散，尤其是对于合并的内存访问。并且不能保证 Nvidia 不会再次为未来的架构恢复行为以节省芯片面积或例如将有 8 个子组，每个子组 4 个线程，它们不能发散。

标签： cuda

【解决方案1】：

每个 Cuda 多处理器都有执行单元（每个都有几个用于 int、float、特殊函数...）。它们作为管道工作，需要几个周期才能完成计算，但在每个周期中都可以插入（=已调度）新计算，并且在管道的不同阶段同时处理多个计算。

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

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

在 Volta 中，引入了独立线程调度。每个线程都有自己的指令计数器（并管理一个单独的函数调用堆栈）。但是调度程序仍然会为活动线程调度具有位域的 32 个线程（warp）组。改变的是调度程序可以交错不同的路径。如果可用的执行单元或内存延迟更适合，它可以执行 CCCIEEIIECCC，而不是执行 CCCIIIEEECCC pre-Volta（指令：C=common，I=if 分支，e=else 分支）。作为程序员，必须小心，因为不能再假设线程没有分歧，即使在执行相同的指令时也是如此。这就是引入 __syncwarp 并且所有类型的协作功能（例如 shuffle 指令）都有同步变体的原因。尽管如此（尽管我们无法确定线程是否分叉），如果同步执行，特别是对于合并的内存访问，仍然必须以所有 32 个线程可以一起工作的方式进行编程。在每条可能发散的指令之后放置 __syncwarp 有助于确保收敛。（但要进行性能分析）。

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

【讨论】：