cuda程序编译流程
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.主机代码路径（左侧）
- C++ 预处理器：处理
.cu文件的主机部分，生成.cpp4.ii文件。 - CUDA 前端编译器
cudafe++：将预处理后的.cpp4.ii文件转换为.cudafe1.cpp文件。这个文件包含主机代码和设备代码的占位符。 - C++ 编译器：处理
.cudafe1.cpp文件，生成主机端的.o或.obj文件。这是CPU可执行的部分。
b.设备代码路径（右侧）
设备代码的编译流程比主机代码复杂，详细流程如下：
- C++ 预处理器：CUDA源代码中的设备代码被提取出来，并进行预处理，生成
.cpp1.ii文件。这个文件是设备代码的初步处理结果。 - CUDA 前端编译器
cudafe：cudafe 将.cpp1.ii转换为.cudafe1.gpu，这是设备端的中间表示文件,它的主要作用是处理CUDA设备代码的结构，提取__global__、__device__和__host__修饰的函数，并将这些代码转换为适合进一步编译的中间形式。 - C 编译器预处理：对
.cudafe1.gpu进行C预处理，生成.cpp2.i文件.在这个步骤中，设备代码再次经过C编译器的预处理器。这是为了进一步处理包含的头文件、宏定义等，并对 cudafe1.gpu 文件中的内容进行进一步的C语言级别的预处理。 - CUDA 前端编译器
cudafe：再次编译，生成.cudafe2.gpu文件，准备进一步转换为GPU代码。在这一步，cudafe 再次处理设备代码，分析C语言的结构，并将 .cpp2.i 文件中的设备代码转换为 cudafe2.gpu 文件。这是准备进行PTX代码生成的关键一步，它将设备代码进一步转化为中间表示，准备进入GPU架构相关的编译步骤。 - C 编译器预处理：将
.cudafe2.gpu再次预处理，生成.cpp3.i文件。 - CUDA 编译器
cicc：将.cpp3.i文件编译为 PTX（并行线程执行）中间代码，输出.ptx文件。cicc 是CUDA编译器的核心组件之一，它负责将设备代码编译为PTX（Parallel Thread Execution）代码。PTX是NVIDIA的中间代码表示，类似于汇编语言，它是高层次的机器码，独立于具体的GPU硬件架构。 - PTX 汇编器
ptxas：将.ptx文件转化为设备可执行的.cubin文件。这是最终的设备二进制文件，GPU可以直接执行这个文件。
4. 生成胖二进制文件 fatbinary
通过 fatbinary 工具，多个 .cubin 文件被打包成 .fatbin 文件。这是“胖二进制”文件，支持不同架构的GPU设备。
5. 生成 .fatbin.c 文件
生成的 fatbin 文件会被转换为 .fatbin.c 文件，这是C语言代码文件，包含了设备代码的二进制部分，最终将和主机代码一起被编译。
6. 最终链接
最后，主机端的 .o 文件和设备端的 .fatbin.c 文件通过标准C/C++编译器（如 gcc 或 g++）进行链接。这个过程会将主机代码和设备代码一起打包成最终的可执行文件。
编译链接全过程
Ref
- https://blog.csdn.net/dark5669/article/details/53869631
- https://www.zhangty15226.com/2023/11/25/NVCC%E7%BC%96%E8%AF%91%E6%B5%81%E7%A8%8B/
- https://blog.csdn.net/fb_help/article/details/80462853
- https://cloud.baidu.com/article/3224884
- https://findhao.net/easycoding/2039
- https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
- https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html
- https://docs.nvidia.com/cuda/cuda-runtime-api/index.html
- https://docs.nvidia.com/cuda/cuda-driver-api/index.html
- https://developer.nvidia.com/cuda-examplehttps://developer.nvidia.com/gpu-computing-sdk