cuda程序编译流程

本文以cuda example的matrixMul矩阵乘法为例说明cuda程序的编译流程。

1. 源代码 .cu 文件

matrixMul示例中,源代码文件 matrixMul.cu 是典型的CUDA程序,包含以下部分:

流程图

在这里插入图片描述

  • 主机代码(Host Code):运行在CPU上的代码,用于数据准备、调用CUDA内核(kernel)等。
  • 设备代码(Device Code):运行在GPU上的CUDA内核,负责矩阵乘法计算。
#include <stdio.h>
#include <cuda_runtime.h>

// CUDA内核:用于执行矩阵乘法
__global__ void MatrixMulKernel(float* C, const float* A, const float* B, int width) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;  // 矩阵的列索引
    int y = blockIdx.y * blockDim.y + threadIdx.y;  // 矩阵的行索引

    if (x < width && y < width) {
        float sum = 0;
        for (int i = 0; i < width; ++i) {
            sum += A[y * width + i] * B[i * width + x];  // A的行与B的列对应相乘累加
        }
        C[y * width + x] = sum;  // 结果存储在C矩阵中
    }
}

void randomMatrixInit(float* mat, int size) {
    for (int i = 0; i < size; ++i) {
        mat[i] = rand() % 10;  // 随机初始化矩阵中的每个元素
    }
}

int main() {
    int width = 16;  // 矩阵的宽度(假设矩阵是正方形,大小为 width * width)
    int size = width * width;  // 矩阵的总元素个数

    // 分配主机内存
    float* h_A = (float*)malloc(size * sizeof(float));
    float* h_B = (float*)malloc(size * sizeof(float));
    float* h_C = (float*)malloc(size * sizeof(float));

    // 初始化矩阵A和B
    randomMatrixInit(h_A, size);
    randomMatrixInit(h_B, size);

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

    // 将主机内存中的数据拷贝到设备内存
    cudaMemcpy(d_A, h_A, size * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size * sizeof(float), cudaMemcpyHostToDevice);

    // 定义CUDA网格和块大小
    int blockSize = 16;  // 每个线程块中的线程数(16 x 16 的线程块)
    dim3 threadsPerBlock(blockSize, blockSize);
    dim3 numBlocks((width + blockSize - 1) / blockSize, (width + blockSize - 1) / blockSize);

    // 调用CUDA内核进行矩阵乘法运算
    MatrixMulKernel<<<numBlocks, threadsPerBlock>>>(d_C, d_A, d_B, width);

    // 将结果从设备内存拷贝回主机内存
    cudaMemcpy(h_C, d_C, size * sizeof(float), cudaMemcpyDeviceToHost);

    // 输出结果(可选)
    printf("Matrix A:\n");
    for (int i = 0; i < size; i++) {
        printf("%f ", h_A[i]);
        if ((i + 1) % width == 0) printf("\n");
    }

    printf("\nMatrix B:\n");
    for (int i = 0; i < size; i++) {
        printf("%f ", h_B[i]);
        if ((i + 1) % width == 0) printf("\n");
    }

    printf("\nMatrix C (Result):\n");
    for (int i = 0; i < size; i++) {
        printf("%f ", h_C[i]);
        if ((i + 1) % width == 0) printf("\n");
    }

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

    // 释放主机内存
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

2. C++ 预处理器

编译的第一步是C++预处理器,它处理宏定义、头文件包含以及代码中的条件编译部分。在此阶段,.cu文件会被预处理,生成中间文件 .cpp.i

3. CUDArtfe 编译步骤(C++ 和 CUDA 编译器)

a.主机代码路径(左侧)
  1. C++ 预处理器:处理.cu文件的主机部分,生成 .cpp4.ii 文件。
  2. CUDA 前端编译器 cudafe++:将预处理后的 .cpp4.ii 文件转换为 .cudafe1.cpp 文件。这个文件包含主机代码和设备代码的占位符。
  3. C++ 编译器:处理 .cudafe1.cpp 文件,生成主机端的 .o.obj 文件。这是CPU可执行的部分。
b.设备代码路径(右侧)

设备代码的编译流程比主机代码复杂,详细流程如下:

  1. C++ 预处理器:CUDA源代码中的设备代码被提取出来,并进行预处理,生成 .cpp1.ii 文件。这个文件是设备代码的初步处理结果。
  2. CUDA 前端编译器 cudafe:cudafe 将 .cpp1.ii 转换为 .cudafe1.gpu,这是设备端的中间表示文件,它的主要作用是处理CUDA设备代码的结构,提取 __global____device____host__ 修饰的函数,并将这些代码转换为适合进一步编译的中间形式。
  3. C 编译器预处理:对 .cudafe1.gpu 进行C预处理,生成 .cpp2.i 文件.在这个步骤中,设备代码再次经过C编译器的预处理器。这是为了进一步处理包含的头文件、宏定义等,并对 cudafe1.gpu 文件中的内容进行进一步的C语言级别的预处理。
  4. CUDA 前端编译器 cudafe:再次编译,生成 .cudafe2.gpu 文件,准备进一步转换为GPU代码。在这一步,cudafe 再次处理设备代码,分析C语言的结构,并将 .cpp2.i 文件中的设备代码转换为 cudafe2.gpu 文件。这是准备进行PTX代码生成的关键一步,它将设备代码进一步转化为中间表示,准备进入GPU架构相关的编译步骤。
  5. C 编译器预处理:将 .cudafe2.gpu 再次预处理,生成 .cpp3.i 文件。
  6. CUDA 编译器 cicc:将 .cpp3.i 文件编译为 PTX(并行线程执行)中间代码,输出 .ptx 文件。cicc 是CUDA编译器的核心组件之一,它负责将设备代码编译为PTX(Parallel Thread Execution)代码。PTX是NVIDIA的中间代码表示,类似于汇编语言,它是高层次的机器码,独立于具体的GPU硬件架构。
  7. PTX 汇编器 ptxas:将 .ptx 文件转化为设备可执行的 .cubin 文件。这是最终的设备二进制文件,GPU可以直接执行这个文件。

4. 生成胖二进制文件 fatbinary

通过 fatbinary 工具,多个 .cubin 文件被打包成 .fatbin 文件。这是“胖二进制”文件,支持不同架构的GPU设备。

5. 生成 .fatbin.c 文件

生成的 fatbin 文件会被转换为 .fatbin.c 文件,这是C语言代码文件,包含了设备代码的二进制部分,最终将和主机代码一起被编译。

6. 最终链接

最后,主机端的 .o 文件和设备端的 .fatbin.c 文件通过标准C/C++编译器(如 gccg++)进行链接。这个过程会将主机代码和设备代码一起打包成最终的可执行文件。

编译链接全过程

在这里插入图片描述


Ref

  1. https://blog.csdn.net/dark5669/article/details/53869631
  2. https://www.zhangty15226.com/2023/11/25/NVCC%E7%BC%96%E8%AF%91%E6%B5%81%E7%A8%8B/
  3. https://blog.csdn.net/fb_help/article/details/80462853
  4. https://cloud.baidu.com/article/3224884
  5. https://findhao.net/easycoding/2039
  6. https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
  7. https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html
  8. https://docs.nvidia.com/cuda/cuda-runtime-api/index.html
  9. https://docs.nvidia.com/cuda/cuda-driver-api/index.html
  10. https://developer.nvidia.com/cuda-examplehttps://developer.nvidia.com/gpu-computing-sdk
Logo

欢迎来到由智源人工智能研究院发起的Triton中文社区,这里是一个汇聚了AI开发者、数据科学家、机器学习爱好者以及业界专家的活力平台。我们致力于成为业内领先的Triton技术交流与应用分享的殿堂,为推动人工智能技术的普及与深化应用贡献力量。

更多推荐