随着人工智能的发展,科学计算(尤其是矩阵/张量计算)越来越重要;因此,基于CUDA的张量编程也越来越重要。
在上一篇笔记中翻译了《An Even Easier Introduction to CUDA》,但是感觉作者写的不是很好;
这里重新写了一篇。同时,也作为CUDA和并行编程的开篇。
源代码:
<br/>
<!--more-->温馨提示:本文章配合 Colab 一同执行学习效果更佳:
CUDA 是 NVIDIA 开发的并行计算平台和编程模型;
具有以下特点:
<br/>
下面是一个 GPU 硬件单元:
每个核心中包含了多个 SM(Stream Multi-processor),任务在 SM 中处理;
SM 中包含了:
<br/>
CPU 与 GPU 协同工作的流程如下:
首先,习惯上将:
通常,Global Memory 在其范围和生命周期中是全局的!
也就是说,每个在thread block grid 中的 thread 都可以访问Global Memory,并且生命周期与程序的执行时间一样长!
更多内容:
CUDA 程序执行时主要分为以下几个步骤:
CUDA 这种流程实现了 CPU 与 GPU 协同,让 GPU 承担并行计算 heavy - lifting ,提升计算密集型任务效率,广泛用于深度学习训练推理、科学计算等场景!
<br/>
add_cpu.cpp
#include <cmath>
#include <iostream>
#include <vector>
// Step 2: Define add function
void add_cpu(std::vector<float> &c, const std::vector<float> &a, const std::vector<float> &b) {
// CPU use loop to calculate
for (size_t i = 0; i < a.size(); i++) {
c[i] = a[i] + b[i];
}
}
int main() {
// Step 1: Prepare & initialize data
constexpr size_t N = 1 << 20; // ~1M elements
// Initialize data
const std::vector<float> a(N, 1);
const std::vector<float> b(N, 2);
std::vector<float> c(N, 0);
// Step 3: Call the cpu addition function
add_cpu(c, a, b);
// Step 4: Check for errors (all values should be 3.0f)
float maxError = 0.0f;
for (int i = 0; i < N; i++) {
maxError = fmax(maxError, fabs(c[i] - 3.0f));
}
std::cout << "Max error: " << maxError << std::endl;
}
主要分为以下几个步骤:
<br/>
分为以下几个步骤:
*.cu
:例如add_cuda.cu
(表示 CUDA 程序)cudaMalloc
分配显存、使用 cudaMemcpy
复制数据等;下面分别来看;
<br/>
*.cu
CUDA 规定其文件扩展名为 *.cu
,语法和 C++ 类似!
<br/>
这步比较简单:
add_cuda.cu
// Step 1: Prepare & initialize data
constexpr size_t N = 1 << 20; // ~1M elements
constexpr size_t size_bytes = sizeof(float) * N;
// Initialize data
const std::vector<float> h_a(N, 1);
const std::vector<float> h_b(N, 2);
std::vector<float> h_c(N, 0);
此时在 Host 端的 RAM 分配内存;
<br/>
数据传输到 GPU 使用 CUDA 提供的函数:
cudaMalloc
分配显存;cudaMemcpy
复制数据;add_cuda.cu
float *d_a, *d_b, *d_c;
CUDA_CHECK(cudaMalloc(&d_a, size_bytes));
CUDA_CHECK(cudaMalloc(&d_b, size_bytes));
CUDA_CHECK(cudaMalloc(&d_c, size_bytes));
CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), size_bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), size_bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_c, h_c.data(), size_bytes, cudaMemcpyHostToDevice));
这里使用了:
CUDA_CHECK
宏进行校验;cudaMemcpyHostToDevice
指定数据流方向;CUDA_CHECK 宏定义如下:
#define CUDA_CHECK(call) \
{ \
cudaError_t err = call; \
if (err != cudaSuccess) { \
std::cerr << "CUDA Error at " << __FILE__ << ":" << __LINE__ \
<< " - " << cudaGetErrorString(err) << std::endl; \
} \
}
<br/>
在 CPU 中,使用循环进行执行;
而在 GPU 中,使用的是 SIMT
,即:一条命令会同时被多个线程执行!
此时需要指挥每个线程:组织结构和编号!
在 CUDA 中,包含:
如下图:
其中:每一个 Grid 中包含多个已编号的 Block,而每一个 Block 中包含多个已编号的 Thread!
同时,每个 Block 中包含的线程数是一样的!
一共有:0~N-1
个Thread(假设每个 Block 包含 N 个 Thread);
<br/>
在 CUDA 中:
idx = BlockID * BlockSize + ThreadID
;如下图:
<br/>
相对于 CPU 中使用循环的方式执行,在 GPU 中主要使用的是:多线程并行
;
步骤如下:
层级结构定义:
// Set up kernel configuration
dim3 block_dim(256);
dim3 grid_dim((N + block_dim.x - 1) / block_dim.x);
核函数定义:
template<typename T>
__global__ void add_kernel(T *c, const T *a, const T *b, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
只能通过指针的方式传递!
<font color="#f00">**因为像是 vector 等数据结构,都是在 Host 端定义的,并不能在 Global Memory 中分配!**</font>
核函数调用:
// Call cuda add kernel
add_kernel<<<grid_dim, block_dim>>>(d_c, d_a, d_b, N);
其中:
dim3
:CUDA 表示线程层级结构的类型(包括:x、y、z 三个维度);<<<>>>
:传递线程层级信息给核函数;核函数(kernel)
:设备侧的入口函数;__global__
:表示这是个核函数;blockldx
:表示 block 的编号,第几个 block;blockDim
:表示 block 的大小,一个 block 多少个线程;threadldx
:表示 thread 的编号,表示第几个 thread;<br/>
同样,使用 cudaMemcpy
:
CUDA_CHECK(cudaMemcpy(h_c.data(), d_c, size_bytes, cudaMemcpyDeviceToHost));
<br/>
验证结果使用已经复制到 h_c
中的数据;
释放内存使用 cudaFree
:
add_cuda.cu
float maxError = 0.0f;
for (int i = 0; i < N; i++) {
maxError = fmax(maxError, fabs(h_c[i] - 3.0f));
}
std::cout << "Max error: " << maxError << std::endl;
if (d_a) {
CUDA_CHECK(cudaFree(d_a));
}
if (d_b) {
CUDA_CHECK(cudaFree(d_b));
}
if (d_c) {
CUDA_CHECK(cudaFree(d_c));
}
<br/>
需要使用 NVCC(NIVIDEA CUDA Compiler)
的编译器来编译程序;
NVCC 是 CUDA Toolkit 的一部分:
<br/>
如下图所示:
流程如下:
<br/>
NVIDIA 不同年代生产的GPU可能有不同的架构,如下图所示:
以 A100 为例,A100 为 Ampere 架构;同时,各个架构间有区别;
因此提出:Compute Capability (CC)
虽然 A100 举例,但从 CUDA 编程的角度目前各种架构没有本质区别!
<font color="#f00">**正因为如此,所以说CUDA是一个编程平台**</font>
同时,在编译时也可以指定架构编译选项:
-arch
:指定虚拟架构,PTX生成目标。决定代码中可使用的CUDA 功能;-code
:指定实际架构,生成针对特定 GPU 硬件的二进制机器码(cubin);<br/>
通过:
nvcc add_cuda.cu -o add_cuda
即可编译!
运行:
./add_cuda
<br/>
可以通过:
nvidia-smi
观察 GPU 利用率!
分别对比:
<<<1,1>>>
;<<<256,256>>>
;代码如下:
add_cuda_profiling.cu
#include <cmath>
#include <iostream>
#include <vector>
#define CUDA_CHECK(call) \
{ \
cudaError_t err = call; \
if (err != cudaSuccess) { \
std::cerr << "CUDA Error at " << __FILE__ << ":" << __LINE__ \
<< " - " << cudaGetErrorString(err) << std::endl; \
} \
}
// Step 3: Define add kernel
template<typename T>
__global__ void add_kernel(T *c, const T *a, const T *b, const size_t n, const size_t step) {
int idx = blockIdx.x * blockDim.x + threadIdx.x + step;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
template<typename T>
void vector_add(T *c, const T *a, const T *b, size_t n, const dim3& grid_dim, const dim3& block_dim) {
size_t step = grid_dim.x * block_dim.x;
for (size_t i = 0; i < n; i += step) {
add_kernel<<<grid_dim, block_dim>>>(c, a, b, n, i);
}
}
int main() {
// Step 1: Prepare & initialize data
constexpr size_t N = 1 << 20; // ~1M elements
constexpr size_t size_bytes = sizeof(float) * N;
// Initialize data
const std::vector<float> h_a(N, 1);
const std::vector<float> h_b(N, 2);
std::vector<float> h_c(N, 0);
// Step 2: Allocate device memory & transfer to global memory
float *d_a, *d_b, *d_c;
CUDA_CHECK(cudaMalloc(&d_a, size_bytes));
CUDA_CHECK(cudaMalloc(&d_b, size_bytes));
CUDA_CHECK(cudaMalloc(&d_c, size_bytes));
CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), size_bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), size_bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_c, h_c.data(), size_bytes, cudaMemcpyHostToDevice));
// Step 4: Call the cpu addition function
// Set up kernel configuration
dim3 block_dim(1);
dim3 grid_dim(1);
// Call cuda add kernel
vector_add(d_c, d_a, d_b, N, block_dim, grid_dim);
// Step 5: Transfer data from global mem to host mem
CUDA_CHECK(cudaMemcpy(h_c.data(), d_c, size_bytes, cudaMemcpyDeviceToHost));
// Step 6: Check for errors (all values should be 3.0f)
float sumError = 0.0f;
for (int i = 0; i < N; i++) {
sumError += fabs(h_c[i] - 3.0f);
}
std::cout << "Sum error: " << sumError << std::endl;
if (d_a) {
CUDA_CHECK(cudaFree(d_a));
}
if (d_b) {
CUDA_CHECK(cudaFree(d_b));
}
if (d_c) {
CUDA_CHECK(cudaFree(d_c));
}
}
可以修改其中的:
dim3 block_dim(1); dim3 grid_dim(1);
不同情况下的性能如下:
可以使用 Nsight Systems(nsys,NVIDIA系统级性能分析工具)来分析;
启动 profiling:
nsys profile -t cuda,nvtx,osrt -o add_cuda_profiling -f true ./add_cuda_profiling
解析并统计性能信息:
nsys stats add_cuda_profiling.nsys-rep
** OS Runtime Summary (osrt_sum):
Time (%) Total Time (ns) Num Calls Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- ------------- ------------- ---------- ----------- ------------- ----------------------
56.2 7,592,724,284 84 90,389,574.8 100,130,776.0 2,330 370,626,986 45,049,255.4 poll
42.4 5,736,493,727 26 220,634,374.1 189,702,756.5 41,077,614 752,975,386 124,762,585.8 sem_wait
1.2 164,252,099 543 302,490.1 13,509.0 529 111,402,991 4,818,716.4 ioctl
0.1 14,968,499 38 393,907.9 131,267.0 135 5,539,804 890,642.6 pthread_rwlock_wrlock
......
** CUDA API Summary (cuda_api_sum):
Time (%) Total Time (ns) Num Calls Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- ------------ ----------- -------- ----------- ------------- ----------------------
96.9 6,504,565,162 1,048,576 6,203.2 5,159.0 2,928 37,814,020 99,097.6 cudaLaunchKernel
3.0 203,141,797 3 67,713,932.3 103,908.0 73,162 202,964,727 117,130,625.1 cudaMalloc
0.1 4,017,591 4 1,004,397.8 1,012,632.0 941,545 1,050,782 45,652.8 cudaMemcpy
0.0 524,788 3 174,929.3 136,182.0 122,785 265,821 78,999.0 cudaFree
0.0 2,584 1 2,584.0 2,584.0 2,584 2,584 0.0 cuModuleGetLoadingMode
......
各个类型 API Summary 分析结果如下:
可以看到:<<<1,1>>>
cudaLaunchKernel 占比非常高这是由于:
核函数调用有开销,在外面多次循环调用开销巨大!
因此,需要进行优化!
<br/>
由于在循环中频繁的调用核函数具有巨大的性能开销,因此可以将循环放入核函数中:
template<typename T>
__global__ void add_kernel_inner_loop(T *c, const T *a, const T *b, const size_t n, const size_t step) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for (size_t i = idx; i < n; i += step) {
if (i < n) {
c[i] = a[i] + b[i];
}
}
}
template<typename T>
void vector_add(T *c, const T *a, const T *b, size_t n, const dim3& grid_dim, const dim3& block_dim) {
size_t step = grid_dim.x * block_dim.x;
add_kernel_inner_loop<<<grid_dim, block_dim>>>(c, a, b, n, step);
}
分析后结果如下图:
同时使用 nsys 分析:
** CUDA API Summary (cuda_api_sum):
Time (%) Total Time (ns) Num Calls Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- ------------ ----------- -------- ----------- ------------- ----------------------
55.4 204,935,456 3 68,311,818.7 104,741.0 79,097 204,751,618 118,160,333.0 cudaMalloc
44.4 164,057,041 4 41,014,260.3 1,000,521.5 926,775 161,129,223 80,076,651.2 cudaMemcpy
0.2 653,441 3 217,813.7 204,732.0 194,409 254,300 32,016.9 cudaFree
0.1 264,055 1 264,055.0 264,055.0 264,055 264,055 0.0 cudaLaunchKernel
0.0 2,429 1 2,429.0 2,429.0 2,429 2,429 0.0 cuModuleGetLoadingMode
可以看到 cudaLaunchKernel 的确少了很多!
这说明:
<font color="#f00">**核函数的发射数量减少,因此总体执行时间降低!**</font>
<br/>
指标:
加速比 = T<sub>cpu</sub> / T<sub>gpu</sub>
其中:
理想加速比与实际加速比
- 理想加速比:当任务完全并行化且没有任何开销时,加速比等于处理器核心数之比。例如,一个具有 1000 个 CUDA 核心的 GPU 理论上可以实现 1000 倍的加速(相对于单核 CPU)。
- 实际加速比:由于以下因素,实际加速比通常远低于理想值:
- 任务中存在无法并行化的串行部分
- 数据在 CPU 和 GPU 之间的传输开销
- 线程同步和内存访问延迟
- 算法在 GPU 架构上的效率低下
<<<1,1>>>
比 CPU慢?这是由于,单个 GPU 的核心实际上要比 CPU 的能力要弱!
实际上,GPU 是由于干活的人多,所以快!
<br/>
实际上观察 nsys 的输出结果:
** CUDA GPU Kernel Summary (cuda_gpu_kern_sum):
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- ------------- ------------- ----------- ----------- ----------- ---------------------------------------------------------------------------------------------
100.0 160,054,287 1 160,054,287.0 160,054,287.0 160,054,287 160,054,287 0.0 void add_kernel_inner_loop<float>(T1 *, const T1 *, const T1 *, unsigned long, unsigned long)
Processing [add_cuda_profiling2.sqlite] with [/usr/local/cuda-12.1/nsight-systems-2023.1.2/host-linux-x64/reports/cuda_gpu_mem_time_sum.py]...
** CUDA GPU MemOps Summary (by Time) (cuda_gpu_mem_time_sum):
Time (%) Total Time (ns) Count Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Operation
-------- --------------- ----- --------- --------- -------- -------- ----------- ------------------
78.4 2,318,310 3 772,770.0 763,159.0 761,400 793,751 18,191.4 [CUDA memcpy HtoD]
21.6 640,473 1 640,473.0 640,473.0 640,473 640,473 0.0 [CUDA memcpy DtoH]
总体的耗时应当是三个部分:
并且,对于 <<<256,256>>>
来说:HtoD 和 DtoH 的耗时会远大于 kernel 的运行时间!
这就说明,来回的移动和复制数据比计算更消耗时间!
能否对这一部分进行优化呢?
后面的文章中会讲解!
<br/>
对于 GPU 而言:越多的线程 => 越大的并行度 => 越好的性能
GPU 最大可以启动的线程数可以参考:
也可以参考:
<br/>
重点的几个参数:
其中:Blocksize 需同时满足这两组条件:maxBlockSize、maxThreadsPerBlock:
- x、y、z加起来不超过:maxThreadsPerBlock;
- x、y、z各个方向不超过:maxBlockSize;
<br/>
查看 CUDA 版本使用:
# CUDA版本
nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Tue_Feb__7_19:32:13_PST_2023
Cuda compilation tools, release 12.1, V12.1.66
Build cuda_12.1.r12.1/compiler.32415258_0
可以看到 CUDA 为 12.1!
而 nvidia-smi
命令输出的是:驱动支持的的最高版本,而非实际正在使用的版本!
Tue Jul 29 09:30:09 2025
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 550.54.15 Driver Version: 550.54.15 CUDA Version: 12.4 |
|-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 Tesla T4 Off | 00000000:00:04.0 Off | 0 |
| N/A 38C P8 10W / 70W | 0MiB / 15360MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
| No running processes found |
+-----------------------------------------------------------------------------------------+
可以看到,最高支持 12.4!
<br/>
更加详细的内容,可以看:
<br/>
源代码:
参考文章:
<br/>