【CUDA编程部署教程】第一章：GPU架构与CUDA编程核心 1

CUDA编程核心解析与优化技巧

VectorShift

已于 2025-10-04 13:40:44 修改

阅读量49

点赞数

CC 4.0 BY-SA版权

分类专栏：人工智能技术白皮书文章标签：开发语言 cuda 人工智能机器学习

于 2025-10-04 13:33:41 首次发布

本文链接：https://blog.youkuaiyun.com/VectorShift/article/details/152511498

人工智能技术白皮书专栏收录该内容

156 篇文章 ¥19.90 ¥99.00

订阅专栏

超级会员免费看

引言

在追求极致的计算性能时，软件的优化极限取决于其对硬件运行范式的理解深度。本章不罗列冗长的硬件参数，而是直击NVIDIA GPU架构的核心设计哲学，并将其与CUDA编程模型紧密结合。我们将剖析现代GPU的计算心脏——流式多处理器（SM），理解其独特的Warp执行模型，探索Tensor Core带来的革命性矩阵运算加速，并深入探讨决定性能命脉的内存体系。掌握这些底层知识，是编写出能压榨出硬件最后一滴性能的高性能CUDA代码的先决条件。

1.1 现代GPU硬件架构剖析

1.1.1 流式多处理器（SM）与Warp执行模型

原理简介

现代NVIDIA GPU的宏观架构是由一组流式多处理器（Streaming Multiprocessors, SM） 构成的并行计算集群。SM是GPU执行计算的最小独立单元。GPU并行计算的核心调度单位并非单个线程，而是由32个线程组成的线程束（Warp）。SM的调度器以Warp为单位分派指令，在一个Warp内，所有32个线程在同一时刻执行完全相同的指令，但处理的数据不同，这种模型被称为单指令多线程（SIMT）。

核心要点：Warp Divergence (线程束分化)

SIMT模型是GPU高效率的源泉，但也是一个关键的性能陷阱。如果一个Warp内的线程因数据相关的条件判断（如if-else）而需要执行不同的代码路径，就会发生Warp分化。此时，SM必须将不同的执行路径串行化执行，导致部分线程空闲等待，计算资源被严重浪费。

技术手册：可执行代码

以下是一个完整的程序，它通过实验来量化Warp分化带来的性能损失。程序将运行两个版本的核函数：一个存在严重分化，另一个通过重构代码避免了分化。

C++

// =================================================================
// 编译指令:
// nvcc -o warp_divergence warp_divergence.cu
// =================================================================
#include <iostream>
#include <vector>
#include <cuda_runtime.h>

// 辅助宏，用于检查CUDA API调用的返回状态
#define CUDA_CHECK(call) do { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA Error: %s:%d, ", __FILE__, __LINE__); \
        fprintf(stderr, "code: %d, reason: %s\n", err, cudaGetErrorString(err)); \
        exit(1); \
    } \
} while(0)

// --- Device Code: 核函数定义 ---

__device__ float g_a = 1.01f, g_b = 2.02f, g_c = 3.03f, g_d = 4.04f;

// 模拟一种复杂计算
__device__ float compute_A(float val) {
    for(int i = 0; i < 50; ++i) val = val * g_a + g_b;
    return val;
}

// 模拟另一种复杂计算
__device__ float compute_B(float val) {
    for(int i = 0; i < 50; ++i) val = val * g_c - g_d;
    return val;
}

// 内核1: 存在严重Warp分化的版本
__global__ void divergent_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= n) return;

    float val = data[idx];

    // 当Warp内数据有正有负时，此处的if/else会导致分化
    if (val > 0) {
        data[idx] = compute_A(val);
    } else {
        data[idx] = compute_B(val);
    }
}

// 内核2: 通过移除分支避免Warp分化的版本
__global__ void branchless_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= n) return;

    float val = data[idx];
    
    // 计算两种可能的结果
    float result_A = compute_A(val);
    float result_B = compute_B(val);

    // 使用三元运算符（或其他无分支逻辑）进行选择
    // 这种方式通常能被编译器优化为谓词执行，避免分化
    data[idx] = (val > 0) ? result_A : result_B;
}

// --- Host Code: 主函数 ---

int main() {
    int n = 1 << 24; // 大约 16 million elements
    size_t bytes = n * sizeof(float);

    // 1. 在主机上初始化数据
    // 我们特意让数据正负交错，以最大化Warp分化的概率
    std::vector<float> h_data(n);
    for (int i = 0; i < n; ++i) {
        h_data[i] = (i % 2 == 0) ? (float)i : (float)-i;
    }

    // 2. 在设备上分配内存
    float *d_data1, *d_data2;
    CUDA_CHECK(cudaMalloc(&d_data1, bytes));
    CUDA_CHECK(cudaMalloc(&d_data2, bytes));

    // 3. 将数据从主机拷贝到设备
    CUDA_CHECK(cudaMemcpy(d_data1, h_data.data(), bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_data2, h_data.data(), bytes, cudaMemcpyHostToDevice));

    // 4. 设置内核启动参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;

    // 5. 创建CUDA事件用于计时
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    float ms;

    // --- 测试Divergent Kernel ---
    CUDA_CHECK(cudaEventRecord(start));
    divergent_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_data1, n);
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
    std::cout << "Divergent Kernel Time:   " << ms << " ms" << std::endl;

    // --- 测试Branchless Kernel ---
    CUDA_CHECK(cudaEventRecord(start));
    branchless_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_data2, n);
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
    std::cout << "Branchless Kernel Time:  " << ms << " ms" << std::endl;

    // 6. 清理资源
    CUDA_CHECK(cudaFree(d_data1));
    CUDA_CHECK(cudaFree(d_data2));
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));

    return 0;
}

分析与预期结果：当您编译并运行此代码时，您会观察到Branchless Kernel的执行时间明显少于Divergent Kernel。尽管Branchless Kernel执行了更多的浮点运算，但它通过保持Warp内线程执行路径的一致性，充分利用了GPU的SIMT并行能力，其性能优势远超额外计算带来的开销。

1.1.2 Tensor Core与混合精度计算

原理简介

Tensor Core 是NVIDIA GPU中的专用计算单元，专门用于硬件加速大规模的矩阵乘加（MMA）运算：D=A×B+C。它通过混合精度计算实现卓越性能：使用低精度（如FP16）输入进行高速乘法，同时使用高精度（FP32）进行累加，以在保持数值稳定性的前提下获得数倍的性能提升和显存节省。

技术手册：可执行代码

直接手写Tensor Core汇编或WMMA intrinsics非常复杂。实践中，我们通过cuBLAS等高性能库来调用它们。下面的完整程序将对比使用Tensor Core (FP16) 与不使用Tensor Core (FP32) 执行大型矩阵乘法的性能。

C++

// =================================================================
// 编译指令:
// nvcc -o tensor_core -lcublas tensor_core.cu
// 注意: 需要Volta (SM 7.0) 或更新架构的GPU才能运行Tensor Core部分
// =================================================================
#include <iostream>
#include <vector>
#include <cuda_runtime.h>
#include <cublas_v2.h>
#include <cuda_fp16.h>

// 辅助宏
#define CUDA_CHECK(call) do { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA Error: %s:%d, ", __FILE__, __LINE__); \
        fprintf(stderr, "code: %d, reason: %s\n", err, cudaGetErrorString(err)); \
        exit(1); \
    } \
} while(0)

#define CUBLAS_CHECK(call) do { \
    cublasStatus_t err = call; \
    if (err != CUBLAS_STATUS_SUCCESS) { \
        fprintf(stderr, "cuBLAS Error: %s:%d, ", __FILE__, __LINE__); \
        fprintf(stderr, "code: %d\n", err); \
        exit(1); \
    } \
} while(0)


int main() {
    // 矩阵维度 (为了最佳性能，通常设为8或16的倍数)
    int m = 4096;
    int n = 4096;
    int k = 4096;

    // --- 1. 初始化主机数据 (FP32) ---
    std::vector<float> h_A(m * k);
    std::vector<float> h_B(k * n);
    for(int i = 0; i < m * k; ++i) h_A[i] = static_cast<float>(rand()) / RAND_MAX;
    for(int i = 0; i < k * n; ++i) h_B[i] = static_cast<float>(rand()) / RAND_MAX;

    // --- 2. 创建cuBLAS句柄 ---
    cublasHandle_t handle;
    CUBLAS_CHECK(cublasCreate(&handle));

    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    float ms;

    // --- 测试1: 标准FP32 GEMM (不使用Tensor Core) ---
    {
        float *d_A, *d_B, *d_C;
        CUDA_CHECK(cudaMalloc(&d_A, m * k * sizeof(float)));
        CUDA_CHECK(cudaMalloc(&d_B, k * n * sizeof(float)));
        CUDA_CHECK(cudaMalloc(&d_C, m * n * sizeof(float)));
        
        CUDA_CHECK(cudaMemcpy(d_A, h_A.data(), m * k * sizeof(float), cudaMemcpyHostToDevice));
        CUDA_CHECK(cudaMemcpy(d_B, h_B.data(), k * n * sizeof(float), cudaMemcpyHostToDevice));
        
        float alpha = 1.0f, beta = 0.0f;
        
        CUDA_CHECK(cudaEventRecord(start));
        // 使用标准SGEMM (Single-precision General Matrix Multiply)
        CUBLAS_CHECK(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &alpha, d_A, m, d_B, k, &beta, d_C, m));
        CUDA_CHECK(cudaEventRecord(stop));
        CUDA_CHECK(cudaEventSynchronize(stop));
        
        CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
        std::cout << "FP32 SGEMM Time (No Tensor Core): " << ms << " ms" << std::endl;

        cudaFree(d_A);
        cudaFree(d_B);
        cudaFree(d_C);
    }

    // --- 测试2: FP16 GEMM (使用Tensor Core) ---
    {
        // 将主机数据从FP32转换为FP16
        std::vector<__half> h_A_half(m * k);
        std::vector<__half> h_B_half(k * n);
        for(int i = 0; i < m * k; ++i) h_A_half[i] = __float2half(h_A[i]);
        for(int i = 0; i < k * n; ++i) h_B_half[i] = __float2half(h_B[i]);
        
        __half *d_A, *d_B, *d_C;
        CUDA_CHECK(cudaMalloc(&d_A, m * k * sizeof(__half)));
        CUDA_CHECK(cudaMalloc(&d_B, k * n * sizeof(__half)));
        CUDA_CHECK(cudaMalloc(&d_C, m * n * sizeof(__half)));

        CUDA_CHECK(cudaMemcpy(d_A, h_A_half.data(), m * k * sizeof(__half), cudaMemcpyHostToDevice));
        CUDA_CHECK(cudaMemcpy(d_B, h_B_half.data(), k * n * sizeof(__half), cudaMemcpyHostToDevice));

        float alpha = 1.0f, beta = 0.0f;

        // 关键: 设置cuBLAS使用Tensor Core路径
        CUBLAS_CHECK(cublasSetMathMode(handle, CUBLAS_TENSOR_OP_MATH));
        
        CUDA_CHECK(cudaEventRecord(start));
        // 使用cublasGemmEx, 指定FP16输入和FP32计算精度
        CUBLAS_CHECK(cublasGemmEx(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &alpha,
                                  d_A, CUDA_R_16F, m,
                                  d_B, CUDA_R_16F, k, &beta,
                                  d_C, CUDA_R_16F, m,
                                  CUDA_R_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP));
        CUDA_CHECK(cudaEventRecord(stop));
        CUDA_CHECK(cudaEventSynchronize(stop));

        CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
        std::cout << "FP16 GEMM Time (Tensor Core):     " << ms << " ms" << std::endl;

        cudaFree(d_A);
        cudaFree(d_B);
        cudaFree(d_C);
    }
    
    // --- 清理 ---
    CUBLAS_CHECK(cublasDestroy(handle));
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));

    return 0;
}

分析与预期结果：在支持Tensor Core的GPU（Volta架构或更新）上运行时，您将看到FP16 GEMM的执行时间远少于FP32 SGEMM，性能提升可能是2倍、4倍甚至更高，具体取决于GPU的型号。这直观地展示了Tensor Core与混合精度计算带来的巨大威力。

1.1.3 GPU内存层级、带宽与延迟

原理简介

GPU的性能不仅取决于计算速度，更受限于其从内存中获取数据的速度。高性能CUDA编程的精髓在于最大化数据复用，并最小化对慢速全局内存的访问。当必须访问全局内存时，则必须遵循**合并访问（Coalesced Access）**的原则，即让一个Warp内的32个线程访问连续的内存地址，从而让硬件将多次访问合并为一次高效的内存事务，以充分利用带宽。

技术手册：可执行代码

以下程序通过矩阵转置任务，对比了非合并访问与利用**共享内存（Shared Memory）**实现合并访问的性能差异，并包含了结果验证。

C++

// =================================================================
// 编译指令:
// nvcc -o transpose transpose.cu
// =================================================================
#include <iostream>
#include <vector>
#include <cmath>
#include <cuda_runtime.h>

#define CUDA_CHECK(call) do { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA Error: %s:%d, ", __FILE__, __LINE__); \
        fprintf(stderr, "code: %d, reason: %s\n", err, cudaGetErrorString(err)); \
        exit(1); \
    } \
} while(0)

#define TILE_DIM 32

// 内核1: 朴素转置，写入时存在非合并访问
__global__ void naive_transpose(float* out, const float* in, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < width && y < height) {
        // 读取是合并的
        float val = in[y * width + x];
        // 写入是非合并的 (strided)
        out[x * height + y] = val;
    }
}

// 内核2: 使用共享内存优化，读写均为合并访问
__global__ void shared_mem_transpose(float* out, const float* in, int width, int height) {
    __shared__ float tile[TILE_DIM][TILE_DIM];

    int x_in = blockIdx.x * TILE_DIM + threadIdx.x;
    int y_in = blockIdx.y * TILE_DIM + threadIdx.y;
    int in_idx = y_in * width + x_in;

    // 从全局内存以合并方式读取到共享内存
    if (x_in < width && y_in < height) {
        tile[threadIdx.y][threadIdx.x] = in[in_idx];
    }
    
    __syncthreads(); // 块内同步，确保tile加载完毕

    int x_out = blockIdx.y * TILE_DIM + threadIdx.x;
    int y_out = blockIdx.x * TILE_DIM + threadIdx.y;
    int out_idx = y_out * height + x_out;

    // 从共享内存读取，以合并方式写入到全局内存
    if (x_out < height && y_out < width) {
        out[out_idx] = tile[threadIdx.x][threadIdx.y];
    }
}

int main() {
    int width = 4096, height = 4096;
    size_t bytes = width * height * sizeof(float);

    // 1. 初始化主机数据
    std::vector<float> h_in(width * height);
    for (int i = 0; i < width * height; ++i) h_in[i] = (float)i;
    std::vector<float> h_out_naive(width * height);
    std::vector<float> h_out_shared(width * height);

    // 2. 分配设备内存
    float *d_in, *d_out_naive, *d_out_shared;
    CUDA_CHECK(cudaMalloc(&d_in, bytes));
    CUDA_CHECK(cudaMalloc(&d_out_naive, bytes));
    CUDA_CHECK(cudaMalloc(&d_out_shared, bytes));
    CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));

    // 3. 计时器
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    float ms;

    // --- 测试 Naive Transpose ---
    dim3 block_naive(TILE_DIM, TILE_DIM);
    dim3 grid_naive( (width + block_naive.x - 1) / block_naive.x, 
                     (height + block_naive.y - 1) / block_naive.y );

    CUDA_CHECK(cudaEventRecord(start));
    naive_transpose<<<grid_naive, block_naive>>>(d_out_naive, d_in, width, height);
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
    std::cout << "Naive Transpose Time:     " << ms << " ms" << std::endl;

    // --- 测试 Shared Memory Transpose ---
    dim3 block_shared(TILE_DIM, TILE_DIM);
    dim3 grid_shared( (width + TILE_DIM - 1) / TILE_DIM, 
                      (height + TILE_DIM - 1) / TILE_DIM );

    CUDA_CHECK(cudaEventRecord(start));
    shared_mem_transpose<<<grid_shared, block_shared>>>(d_out_shared, d_in, width, height);
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
    std::cout << "Shared Memory Transpose Time: " << ms << " ms" << std::endl;

    // 4. 验证结果
    CUDA_CHECK(cudaMemcpy(h_out_naive.data(), d_out_naive, bytes, cudaMemcpyDeviceToHost));
    CUDA_CHECK(cudaMemcpy(h_out_shared.data(), d_out_shared, bytes, cudaMemcpyDeviceToHost));
    
    bool success = true;
    for (int y = 0; y < height; ++y) {
        for (int x = 0; x < width; ++x) {
            float expected = h_in[x * height + y];
            if (fabs(h_out_shared[y * width + x] - expected) > 1e-5) {
                success = false;
                break;
            }
        }
        if(!success) break;
    }
    std::cout << "Verification: " << (success ? "SUCCESS" : "FAILED") << std::endl;

    // 5. 清理
    cudaFree(d_in);
    cudaFree(d_out_naive);
    cudaFree(d_out_shared);
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    return 0;
}

分析与预期结果：运行该程序，您将观察到Shared Memory Transpose的执行时间远小于Naive Transpose，性能提升可能达到一个数量级。这是因为共享内存版本将一次代价高昂的、非合并的全局内存写入，分解成了一次合并的全局内存读取 + 一次高速的共享内存操作 + 一次合并的全局内存写入。最终的验证SUCCESS也证明了优化的正确性。这完美诠释了利用GPU内存层级进行优化的核心思想。