CUDA编程【7】 线程束

线程束和线程块

线程束

• 线程束是SM中基本的执行单元
• 当一个网格被启动（即一个核函数被启动），它内部的线程块会被分配到一个SM上。 SM 会把线程块中的线程分成多个线程束。
• 在每个线程束中，所有线程按照单指令多线程SIMT的方式执行，每一步执行相同的指令，但是数据不同。

线程块

• 在cuda中，线程块是可以二维、三维的组织的，这只是为了方便写程序，处理图像或者三维的数据
• 事实上在计算机中，内存总是一维存在的，不论线程块的组织方式是几维，只能通过一维的方式来访问内存

对于一个三维组织的线程块，其
tid=threadIdx.x+threadIdx.y×blockDim.x+threadIdx.z×blockDim.x×blockDim.y

线程束的分化问题

你提到的内容是关于 线程束分化（Warp Divergence） 的概念，这是CUDA编程中一个非常重要的性能优化点。以下是对你描述内容的总结和补充：

线程束分化（Warp Divergence）

• 定义：当一个线程束中的线程执行不同的控制流路径时（例如，部分线程执行if块，另一部分执行else块），就会发生线程束分化。
• 原因：CUDA支持C语言的控制流（如if-else、for、while等），但如果线程束中的线程根据条件选择不同的路径，就会导致分化。
• 影响：
  • GPU的硬件调度器只能让线程束中的所有线程执行相同的指令。
  • 如果线程束中的线程需要执行不同的路径，GPU会串行化执行这些路径（先执行if块，再执行else块），导致性能下降。
  • 未执行当前路径的线程会处于等待状态，浪费计算资源。

线程束分化的执行机制

• 执行所有路径：当线程束分化发生时，GPU会让线程束中的所有线程执行所有可能的路径（例如，先执行if块，再执行else块）。
• 掩码机制：GPU会使用掩码（mask）来禁用不需要执行当前路径的线程。例如：
  • 执行if块时，只有条件为true的线程会实际执行，其他线程被禁用。
  • 执行else块时，只有条件为false的线程会实际执行，其他线程被禁用。
• 性能损失：由于所有路径都需要执行，线程束分化的性能损失与分支的数量成正比。

如何避免线程束的分化

• 在if else的情况下，让所有执行if的语句尽量放在一个线程束中，所有执行else的语句放在另一个线程束中。如果顺序是乱的，我们可以事后再来调整顺序。

#include <cuda_runtime.h>
#include <iostream>

__global__ void kernel_with_divergence(float *c)
{
    float a = 0.0;
    float b = 0.0;

    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx % 2)
    {
        a = 100.0;
    }
    else
    {
        b = 200.0;
    }

    c[idx] = a + b;
}

__global__ void kernel_without_divergence(float *c)
{
    float a = 0.0;
    float b = 0.0;

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    idx = idx / warpSize;
    if (idx % 2 == 0)
    {
        a = 100.0;
    }
    else
    {
        b = 200.0;
    }

    c[idx] = a + b;
}

int main()
{

    int gridSize = 256, blockSize = 256;
    int N = gridSize * blockSize;

    float *c_h, *c_d;
    c_h = new float[N];

    cudaMalloc((void **)&c_d, N * sizeof(float));

    cudaEvent_t start, end;
    cudaEventCreate(&start);
    cudaEventCreate(&end);

    cudaEventRecord(start);
    kernel_with_divergence<<<gridSize, blockSize>>>(c_d);
    cudaEventRecord(end);

    cudaEventSynchronize(end);

    float time_with_divergence = 0;
    cudaEventElapsedTime(&time_with_divergence, start, end);

    cudaMemcpy(c_h, c_d, N * sizeof(float), cudaMemcpyDeviceToHost);

    for (int i = 0; i < N; i++)
    {
        printf("%.0f ", c_h[i]);
    }

    printf("\n");
    printf("\n");
    printf("\n");

    cudaEventRecord(start);
    kernel_without_divergence<<<gridSize, blockSize>>>(c_d);
    cudaEventRecord(end);

    cudaEventSynchronize(end);

    float time_without_divergence = 0;
    cudaEventElapsedTime(&time_without_divergence, start, end);

    cudaMemcpy(c_h, c_d, N * sizeof(float), cudaMemcpyDeviceToHost);

    for (int i = 0; i < N; i++)
    {
        printf("%.0f ", c_h[i]);
    }

    printf("\n");

    printf("time_with_divergence: %f , time without divergence: %f", time_with_divergence, time_without_divergence);
}

output

time_with_divergence: 0.014336 , time without divergence: 0.011872

可以看到有分支的要慢于无分支的代码