CUDA编程基础
一、快速理解CUDA编程
1.1 CUDA简介
CUDA(Compute Unified Device Architecture)是由NVIDIA推出的并行计算平台和应用程序接口模型。它允许开发者利用NVIDIA GPU的强大计算能力来加速通用计算任务,而不仅仅是图形渲染。通过CUDA,开发者可以编写C、C++或Fortran代码,并将其扩展以在GPU上运行,从而显著提高性能,特别是在处理大规模数据集和复杂算法时。
1.2 CUDA并行计算的核心思想
1.2.1 并行计算基础与CUDA的独特之处
传统的C/C++编程主要依赖于CPU进行串行执行,即一个指令接一个指令地顺序执行。然而,CUDA编程的核心在于数据并行性和任务并行性。这意味着大量相似的操作可以同时应用于不同的数据元素,或者多个独立的任务可以并行执行。这种并行性特别适合处理大规模的数据集,如图像处理、科学计算等。
CUDA采用的是**单指令多线程(SIMT)**架构,不同于传统的多核CPU上的多线程(SMP)。在SIMT中,一组线程(通常称为warp)会同时执行相同的指令,但作用于不同的数据。这种方式非常适合数据并行的任务,例如矩阵乘法、图像滤波等。
1.2.2 资源的变化:从CPU到GPU
- 计算资源:CPU通常有少量的高性能核心(如4核、8核),每个核心都有复杂的控制逻辑和较大的缓存。相比之下,GPU拥有成百上千个简单的核心,专门设计用于高效处理大量简单任务。
- 内存层次结构:CPU程序主要依赖于少量的高速缓存和大容量的主内存。而在CUDA中,除了全局内存外,还提供了共享内存、寄存器、常量内存和纹理内存等多种内存类型,每种都有其特定用途和访问速度。
- 带宽和延迟:GPU的全局内存带宽远高于CPU,但由于物理距离较长,访问延迟也较高。因此,有效利用共享内存和寄存器是提高CUDA程序性能的关键。
1.2.3 编程思想的变化
- 数据并行化思维:将问题分解为可以并行执行的小任务。例如,向量加法可以通过让每个线程负责一对元素的加法操作来实现。
- 内存管理:需要手动管理设备内存(通过
cudaMalloc
和cudaFree
),并且要考虑到主机与设备之间的数据传输成本。 - 同步机制:由于并行执行的特点,线程间的同步变得至关重要。例如,块内的线程可能需要使用
__syncthreads()
确保它们在继续执行之前完成某些关键步骤。
1.3 CUDA编程的核心流程
1.3.1 初始化与环境设置
- 选择设备:通过调用
cudaSetDevice
选择要使用的GPU设备。 - 分配内存:使用
cudaMalloc
为设备端变量分配内存,使用cudaMemcpy
将数据从主机复制到设备。
1.3.2 编写核函数
- 定义核函数:使用
__global__
关键字声明核函数,指定输入输出参数和执行逻辑。 - 启动核函数:通过
<<<...>>>
语法配置网格和块尺寸,并启动核函数。例如,kernel<<<gridDim, blockDim>>>(args)
。
1.3.3 执行与同步
- 异步执行:可选地使用流(Stream)来并发执行多个内核或拷贝操作,提升效率。
- 同步操作:使用
cudaDeviceSynchronize
等待所有先前启动的内核执行完毕,确保结果可用。
1.3.4 结果回收与清理
- 获取结果:使用
cudaMemcpy
将计算结果从设备复制回主机。 - 释放资源:调用
cudaFree
释放设备端分配的内存,调用cudaDeviceReset
重置设备状态。
1.4 线程层次结构
CUDA编程模型基于一个分层的线程组织结构:
- Grid(网格):由多个线程块组成。
- Block(块):每个块包含多个线程。块是调度的基本单位,同一块内的线程可以协作,例如共享内存。
- Thread(线程):执行实际计算工作的最小单位。
这种层次化的结构允许程序员灵活地根据问题规模调整并行度。
1.4.1 内存层次
CUDA提供多种类型的内存,每种都有其特定用途和访问速度:
- 全局内存(Global Memory):容量大但访问延迟高,所有线程都可以访问。
- 共享内存(Shared Memory):位于每个线程块内部,用于块内线程间的数据交换,访问速度快但容量有限。
- 寄存器(Registers):每个线程私有的高速存储,用于临时变量。
- 常量内存(Constant Memory):只读且缓存优化,适合存储不变的数据。
- 纹理内存(Texture Memory):支持二维空间局部性访问模式,适用于图像处理等应用。
1.4.2 核函数(Kernel)
核函数是在GPU上执行的函数,用__global__
关键字声明。它们不能直接调用,而是需要从主机代码中启动,语法为kernel<<<grid, block>>>(args)
。核函数没有返回值,参数列表可以包括输入输出指针以及尺寸信息等。
1.5 CUDA编程模型
CUDA采用的是**单指令多线程(SIMT)**架构,在这种架构下,一组线程(通常称为warp)会同时执行相同的指令,但作用于不同的数据。这种方式非常适合数据并行的任务,如矩阵运算、图像处理等。
1.5.1 线程索引
为了使每个线程知道它应该处理的数据位置,CUDA提供了几个内置变量:
blockIdx
:当前线程块在整个网格中的索引。threadIdx
:当前线程在其所属块内的索引。blockDim
:当前块的维度大小。gridDim
:整个网格的维度大小。
通过组合这些变量,我们可以计算出每个线程的唯一ID,进而确定该线程应处理的数据位置。
1.5.2 同步机制
在同一块内的线程可以通过调用__syncthreads()
函数实现同步,确保所有线程到达这一点后继续执行。这在需要保证块内线程之间的协调时非常有用,比如在共享内存中读写数据之前。
第二章 详解CUDA架构
2.1 CUDA架构的核心组件
CUDA架构是NVIDIA GPU的并行计算基础,其核心组件包括硬件和软件两部分,共同支撑高性能并行计算。
2.1.1 流式多处理器(Streaming Multiprocessor, SM)
SM是GPU的核心计算单元,每个SM包含多个CUDA核心(CUDA Cores),负责执行线程。SM的架构设计直接影响CUDA程序的性能。以下是SM的关键特性:
- CUDA Core(CUDA核心):执行浮点运算和整数运算的基本单元。例如,最新的Blackwell架构的B200 GPU拥有2080亿个晶体管,每个SM包含数百个CUDA Core。
- 线程调度:SM通过warp(线程组)管理线程。一个warp通常由32个线程组成(在NVIDIA架构中),这些线程并行执行相同的指令(SIMT模型),但作用于不同的数据。
- 资源分配:每个SM包含寄存器、共享内存(Shared Memory)、L1/L2缓存等资源,这些资源的容量和分配策略直接影响线程的执行效率。
2.1.2 内存层次结构
CUDA的内存层次结构分为多层,每层的访问速度和容量不同,开发者需根据需求合理使用:
内存类型 | 访问速度 | 容量 | 用途 |
---|---|---|---|
寄存器(Registers) | 极快 | 线程私有 | 临时变量,访问延迟最低,需合理分配以避免溢出。 |
共享内存(Shared Memory) | 快 | 块内共享 | 块内线程协作,减少全局内存访问(如矩阵乘法中的Tile方法)。 |
L1/L2缓存 | 快 | 小型缓存 | 加速对全局内存的访问,L1缓存位于SM内,L2缓存为全局共享。 |
全局内存(Global Memory) | 较慢 | 大容量 | 存储所有线程可访问的数据,需通过优化访问模式(如内存合并)提升带宽。 |
常量内存(Constant Memory) | 快 | 有限 | 存储只读数据,有独立缓存机制,适合共享不变的数据(如算法参数)。 |
纹理内存(Texture Memory) | 快 | 有限 | 优化空间局部性访问(如图像处理),支持硬件缓存和过滤。 |
关键优化策略:
- 内存合并(Memory Coalescing):线程块内线程按顺序访问连续的全局内存地址时,GPU会将这些访问合并为一个请求,显著提升带宽。例如,线程
i
访问A[i]
时,若线程ID按顺序排列,访问会被合并。 - Bank Conflict:共享内存访问时,若多个线程访问同一bank(共享内存的存储单元),会导致冲突。需通过数据对齐或访问模式调整避免。
2.2 CUDA线程模型与执行流程
CUDA的线程模型基于分层结构,开发者需明确线程、块、网格的组织方式,以最大化并行性。
2.2.1 线程层次结构
CUDA程序的线程分为三个层次:
- Grid(网格):由多个线程块(Blocks)组成,代表整个任务的范围。例如,计算一个矩阵的乘积时,每个线程块负责计算矩阵的一部分。
- Block(线程块):由多个线程(Threads)组成,是GPU调度的最小单位。块内的线程可以协作(如共享内存、同步)。
- Thread(线程):执行核函数的最小单位,每个线程有唯一的ID。
线程ID的计算:
int tid = blockIdx.x * blockDim.x + threadIdx.x; // 一维线程ID
int tid = (blockIdx.y * gridDim.x + blockIdx.x) * (blockDim.x * blockDim.y) +
threadIdx.y * blockDim.x + threadIdx.x; // 二维线程ID
2.2.2 SIMT执行模型
CUDA采用**单指令多线程(SIMT)**架构,其核心思想是:
- Warp级并行:每个warp(32线程)同时执行同一指令,但作用于不同数据(数据并行)。
- 指令分发:SM的调度器将warp的指令分发到CUDA Core,最大化硬件利用率。
- 分支处理:若warp内线程执行不同分支(如
if-else
),会串行化执行所有分支,导致性能下降(称为Warp Divergence)。需尽量减少分支或确保线程路径一致。
2.2.3 同步与通信
- 块内同步:通过
__syncthreads()
确保所有线程到达该点后再继续执行。例如:
__global__ void kernel() {
sharedMem[threadIdx.x] = computeValue();
__syncthreads(); // 确保所有线程完成写入后再读取
result[threadIdx.x] = computeWithSharedMem();
}
块间通信:通过全局内存或原子操作(如atomicAdd
)实现。例如:
__global__ void reduce(int* input, int* output) {
int sum = 0;
for (int i = threadIdx.x; i < N; i += blockDim.x)
sum += input[i];
atomicAdd(output, sum); // 块间累加结果
}
2.3 CUDA编程模型与流程
CUDA编程的核心流程分为以下步骤,需结合硬件架构优化:
2.3.1 核函数(Kernel)设计
核函数是CUDA程序的核心,需遵循以下原则:
- 并行粒度:确保每个线程执行足够多的计算,避免“细粒度”任务(如每个线程只执行简单加法)。
- 资源限制:每个SM的寄存器、共享内存容量有限,需避免因资源不足导致线程块被阻塞。
- 线程块配置:选择合适的块尺寸(如256或512线程)以最大化SM利用率。
2.3.2 内存管理
主机与设备内存分配:
// 主机内存(Host)
float* h_data = (float*)malloc(N * sizeof(float));
// 设备内存(Device)
float* d_data;
cudaMalloc(&d_data, N * sizeof(float));
- 数据传输优化:
- 使用异步传输(
cudaMemcpyAsync
)与**流(Stream)**并行执行计算和传输。 - 通过** pinned memory**(
cudaHostAlloc
)减少CPU-GPU传输延迟。
- 使用异步传输(
2.3.3 并行执行配置
- 网格与块维度:
dim3 blockSize(256, 1);
dim3 gridSize((N + blockSize.x - 1) / blockSize.x, 1);
kernel<<<gridSize, blockSize>>>(d_data);
多维计算:对于二维数据(如图像),可使用二维块和网格:
dim3 blockSize(16, 16);
dim3 gridSize((width + 15)/16, (height + 15)/16);
kernel<<<gridSize, blockSize>>>(d_image);
第三章 CUDA编程核心知识点和编程实践
3.1 核心知识点详解
3.1.1 内存层次结构与优化策略
CUDA的内存层次结构是性能优化的核心,需结合不同内存类型的特点设计代码:
-
全局内存(Global Memory)
- 特点:容量大(可达数十GB),但访问延迟高(约几百个时钟周期)。
- 优化策略:
- 内存合并(Memory Coalescing):确保线程块内线程按顺序访问连续的全局内存地址。例如,线程
i
访问A[i]
时,若线程ID连续,访问会被合并为一个请求。
- 内存合并(Memory Coalescing):确保线程块内线程按顺序访问连续的全局内存地址。例如,线程
__global__ void vectorAdd(float* A, float* B, float* C, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) C[i] = A[i] + B[i]; // 合并访问:线程i访问连续地址
}
-
-
- 数据对齐:确保数据按128-bit边界对齐(如使用
__align__
或cudaMalloc
自动对齐)。
- 数据对齐:确保数据按128-bit边界对齐(如使用
-
-
共享内存(Shared Memory)
- 特点:块内线程共享,访问速度极快(约10-30时钟周期),但容量有限(通常每个SM 96KB)。
- 典型应用:减少全局内存访问次数,如矩阵乘法的Tile方法:
__global__ void matrixMulShared(float* A, float* B, float* C, int N) {
__shared__ float tileA[TILE_WIDTH][TILE_WIDTH];
__shared__ float tileB[TILE_WIDTH][TILE_WIDTH];
// 加载数据到共享内存,后续计算基于共享内存
...
}
-
- 资源限制:需控制共享内存的使用量。例如,若每个线程块使用4KB共享内存,而SM的共享内存总量为96KB,则最多可同时驻留24个块。
-
纹理内存(Texture Memory)
- 特点:只读,带有硬件缓存,适用于具有空间局部性的数据(如图像处理)。
- 使用示例:
// 绑定数据到纹理对象
cudaArray* cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);
cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);
cudaBindTextureToArray(texRef, cuArray, channelDesc);
// 核函数中访问纹理内存
__global__ void textureKernel(...) {
float value = tex2D(texRef, x, y); // 纹理拾取
...
}
常量内存(Constant Memory)
- 特点:只读,独立缓存,适合存储不变的参数(如算法常数)。
- 示例
__constant__ float constData[1024]; // 设备端常量内存
cudaMemcpyToSymbol(constData, h_constData, size); // 主机到设备
3.1.2 线程模型与并行规模
-
线程索引与ID计算
一维线程
int tid = blockIdx.x * blockDim.x + threadIdx.x;
- 二维线程:
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
-
并行规模与SM资源限制
- SM驻留线程数:每个SM的理论最大线程数由其资源(寄存器、共享内存)决定。例如,若SM有64KB共享内存,且每个线程块需4KB共享内存,则最多可驻留16个块。
- 优化策略:
- 调整线程块大小:选择与SM资源匹配的块大小(如256或512线程)。
- 减少资源占用:避免过度使用共享内存或寄存器,防止SM因资源不足而减少驻留线程数。
3.1.3 核函数设计与算术强度
-
算术强度(Arithmetic Intensity)
- 定义:计算操作时间与访存操作时间的比值。高算术强度意味着更少的访存开销,适合GPU加速。
- 提升方法:
- 复用数据:通过共享内存缓存数据,减少全局内存访问(如矩阵乘法的Tile方法)。
- 减少分支:避免条件判断导致的Warp Divergence。
-
核函数设计原则
- 细粒度计算:确保每个线程执行足够多的计算,避免“细粒度”任务(如每个线程仅执行简单加法)。
- 避免冲突:共享内存访问需对齐到bank边界,防止bank冲突。
3.1.4 流(Stream)与异步执行
-
流的定义与作用
- 流:GPU操作的队列,允许多流并行执行,提升设备利用率。
- 特性:
- 同一流内操作按FIFO顺序执行。
- 不同流的操作可重叠,例如一边计算一边传输数据。
2. 流的使用示例
// 创建三个流
cudaStream_t stream0, stream1, stream2;
cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 异步数据传输与计算
cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, stream0);
kernel<<<grid, block, 0, stream1>>>(d_A, d_B);
cudaMemcpyAsync(h_C, d_C, size, cudaMemcpyDeviceToHost, stream2);
// 等待所有流完成
cudaStreamSynchronize(stream0);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
3.2 编程实践与案例
3.2.1 向量加法优化
__global__ void vectorAdd(float* A, float* B, float* C, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) C[i] = A[i] + B[i];
}
// 主机代码
int main() {
int N = 1<<20; // 1M元素
float* d_A, *d_B, *d_C;
cudaMalloc(&d_A, N * sizeof(float));
// ... 初始化数据 ...
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
cudaDeviceSynchronize();
// ... 回收资源 ...
}
3.2.2 矩阵乘法(共享内存优化)
#define TILE_WIDTH 16
__global__ void matrixMulShared(float* A, float* B, float* C, int N) {
__shared__ float s_A[TILE_WIDTH][TILE_WIDTH];
__shared__ float s_B[TILE_WIDTH][TILE_WIDTH];
int bx = blockIdx.x, by = blockIdx.y;
int tx = threadIdx.x, ty = threadIdx.y;
int Row = by * TILE_WIDTH + ty;
int Col = bx * TILE_WIDTH + tx;
float sum = 0;
for (int m = 0; m < (N + TILE_WIDTH - 1)/TILE_WIDTH; m++) {
// 加载数据到共享内存
if (Row < N && (m*TILE_WIDTH + tx) < N)
s_A[ty][tx] = A[Row * N + m*TILE_WIDTH + tx];
else s_A[ty][tx] = 0;
if ((m*TILE_WIDTH + ty) < N && Col < N)
s_B[ty][tx] = B[(m*TILE_WIDTH + ty)*N + Col];
else s_B[ty][tx] = 0;
__syncthreads();
// 计算部分内积
for (int k = 0; k < TILE_WIDTH; k++)
sum += s_A[ty][k] * s_B[k][tx];
__syncthreads();
}
if (Row < N && Col < N)
C[Row * N + Col] = sum;
}
3.2.3 纹理内存加速图像处理
texture<float, 2> texRef; // 定义纹理对象
__global__ void blurKernel(float* out, 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 sum = 0;
for (int dx = -1; dx <= 1; dx++) {
for (int dy = -1; dy <= 1; dy++) {
sum += tex2D(texRef, x+dx, y+dy); // 纹理拾取
}
}
out[y * width + x] = sum / 9.0f;
}
}
// 主机代码
cudaArray* cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);
cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);
cudaBindTextureToArray(texRef, cuArray, channelDesc);
3.3 性能调优技巧
3.3.1 内存优化
- 减少全局内存带宽压力:
- 使用共享内存缓存频繁访问的数据。
- 通过循环展开复用数据:
for (int i = 0; i < N; i += 4) {
sum += A[i] + A[i+1] + A[i+2] + A[i+3];
}
- L2缓存利用:确保数据访问具有空间或时间局部性。
3.3.2 线程配置优化
- 线程块尺寸选择:
- 根据SM资源选择块大小(如256或512线程)。
- 使用
cudaOccupancyMaxPotentialBlockSize
计算最优块大小。
- 网格粒度:确保网格足够大以充分利用所有SM。
3.3.3 算法适配
- SIMD指令:利用CUDA内置函数(如
__shfl_sync
)实现线程块内数据共享。 - 分支优化:避免条件判断,或确保同一warp内线程执行相同分支。
3.4 错误检测与调试
-
运行时错误检查
#define cudaCheckError() { \
cudaError_t e = cudaGetLastError(); \
if (e != cudaSuccess) { \
printf("CUDA Error: %s\n", cudaGetErrorString(e)); \
exit(-1); \
} \
}
// 在关键API调用后检查错误:
cudaMalloc(&d_A, size); cudaCheckError();
CUDA-MEMCHECK工具
- 检测内存越界访问、未初始化内存等:
cuda-memcheck --leak-check full ./your_program
3.5 最新架构特性与实践
3.5.1 Blackwell架构优化
- FP4数据格式:适用于大语言模型(LLM)推理,减少显存占用。
- 多芯片封装(MCM):通过高速互联(10TB/s)提升计算密度,需设计分布式计算任务。
3.5.2 CUDA Graphs
- 静态计算图:将计算流程编译为图,减少API调用开销
cudaGraph_t graph;
cudaGraphExec_t graphExec;
// 创建图并捕获操作
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernel<<<...>>>(...);
cudaMemcpyAsync(...);
cudaStreamEndCapture(stream, &graph);
// 执行图
cudaGraphLaunch(graphExec, stream);