CUDA编程：其五、线程束_cuda 线程束-优快云博客

本文链接：https://blog.youkuaiyun.com/jiexianxiaogege/article/details/139027657

逻辑层面上，一个grid上的所有线程都是并行的，但这仅仅是逻辑层面，物理层面上肯定做不到百万级别的线程并行。就好像CPU线程，物理并行能力要看CPU核数，GPU的并行能力要看线程束（warp）的大小。

一、SP和SM

SP（streaming processor）：最基本的处理单元，也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。这里需要注意一点，一个SP在同一时刻只能处理一个线程，但是可以通过时间片调度，实现多个线程的逻辑并行。这个有点像CPU的核了。

SM（streaming multiprocessor）：多个SP加上其他的一些资源组成一个SM，也叫GPU大核，其他资源如：warp scheduler，register，shared memory等。SM里面SP的数量表示了GPU的并发能力。一个block上的线程只能分配在一个SM上执行，一个SM可以处理多个block。

这是硬件概念！！！

二、线程束（warp）

线程束是软件的概念了。
warp是CUDA调度的基本单位，一个线程束包含 32个线程（我查了很多资料，问了万能的AI，都是告诉我32是定死的）。

线程束中的线程是锁步执行的，即它们会同时执行相同的指令。如果线程束中的线程需要执行不同的指令（例如由于分支条件不同），就会发生分支发散（Warp Divergence），导致性能下降。比如warp里面有32个线程，要根据数据的取值判断进入if还是else分支。16个进入if，16个进入else，那么if和else对应的是不同的指令，就必须先执行其中一个，后执行另一个。

一个SM可以同时执行多个线程束，具体数量取决于SM的资源（如寄存器、共享内存等）。关于这一块，网上资料众说纷纭，主流说法是SM并行执行的warp数量=SM中的SP数量/32。

三、SM、SP、warp、block的关系

**一个block里面的线程，只能在一个SM里面运行。**如果block里面的线程很多，比如192个，那么会被分到6个warp上等待调度。如果SM里面的sp数量不足192，那么6个warp会分批次进入block。

四、线程束分化

先说一下我的显卡环境：
A4000，48个SM，一个SM里面有128个SP。
这是一个非常简单的demo程序，矩阵乘法：

#include <iostream>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdio.h>
#include <stdlib.h>
#include <windows.h>
#include <chrono>

// CUDA 核函数，用于矩阵乘法
__global__ void matrixMultiply(float *A, float *B, float *C, int m, int k, int n) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < m && col < n) {
        float sum = 0.0f;
        for (int i = 0; i < k; ++i) {
            sum += A[row * k + i] * B[i * n + col];
        }
        C[row * n + col] = sum;
    }
}

// 初始化矩阵
void initializeMatrix(float *matrix, int rows, int cols) {
    for (int i = 0; i < rows; ++i) {
        for (int j = 0; j < cols; ++j) {
            matrix[i * cols + j] = static_cast<float>(rand()) / RAND_MAX;
        }
    }
}

int main(int argc, char **argv)
{

    // 调用核函数


    int m = 4096;
    int k = 4096;
    int n = 4096;

    // 分配主机内存
    float *h_A = new float[m * k];
    float *h_B = new float[k * n];
    float *h_C = new float[m * n];

    // 初始化矩阵
    initializeMatrix(h_A, m, k);
    initializeMatrix(h_B, k, n);

    // 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc((void**)&d_A, m * k * sizeof(float));
    cudaMalloc((void**)&d_B, k * n * sizeof(float));
    cudaMalloc((void**)&d_C, m * n * sizeof(float));

    // 将数据从主机复制到设备
    cudaMemcpy(d_A, h_A, m * k * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, k * n * sizeof(float), cudaMemcpyHostToDevice);

    // 定义线程块和网格的维度
    dim3 dimBlock(16, 16);
    dim3 dimGrid((n + dimBlock.x - 1) / dimBlock.x, (m + dimBlock.y - 1) / dimBlock.y);
    
    auto now1 = std::chrono::system_clock::now();
    auto duration1 = now1.time_since_epoch();
    auto sm1 =  std::chrono::duration_cast<std::chrono::milliseconds>(duration1).count();


    matrixMultiply<<<dimGrid, dimBlock>>>(d_A, d_B, d_C, m, k, n);

    cudaDeviceSynchronize();

    auto now2 = std::chrono::system_clock::now();
    auto duration2 = now2.time_since_epoch();
    auto sm2 =  std::chrono::duration_cast<std::chrono::milliseconds>(duration2).count();

    std::cout<<(sm2-sm1)<<std::endl;

    Sleep(100000);

    // 将结果从设备复制到主机
    cudaMemcpy(h_C, d_C, m * n * sizeof(float), cudaMemcpyDeviceToHost);


    // 释放设备内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // 释放主机内存
    delete[] h_A;
    delete[] h_B;
    delete[] h_C;

    return 0;
}

运行时间为120ms。

倘若在核函数里面加上if/else分支？

__global__ void matrixMultiply(float *A, float *B, float *C, int m, int k, int n) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < m && col < n) {
        float sum = 0.0f;
        if(row%2==0)
        {
            for (int i = 0; i < k; ++i) {
                sum += A[row * k + i] * B[i * n + col];
            }
        }
        else
        {
            for (int i = 0; i < k; ++i) {
                sum += A[row * k + i] * B[i * n + col];
            }
        }

        C[row * n + col] = sum;
    }
}

运行时间达到了231ms，时间差不多就是没有分化的两倍。

但是如果分支语句这么写呢？

__global__ void matrixMultiply(float *A, float *B, float *C, int m, int k, int n) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < m && col < n) {
        float sum = 0.0f;

        for (int i = 0; i < k; ++i) {
            if(row%2==0)
            {
                sum += A[row * k + i] * B[i * n + col];
            }
            
            else
            {
                sum += A[row * k + i] * B[i * n + col];
            }
        }

        C[row * n + col] = sum;
    }
}