news 2026/7/2 2:54:07

CUDA系统学习教程

作者头像

张小明

前端开发工程师

1.2k 24
文章封面图
CUDA系统学习教程

课程大纲

课次主题重点内容
1CUDA 基础概念GPU 架构、异构计算模型
2线程层级结构Grid、Block、Thread
3内核函数__global__、__device__、启动语法
4内存管理cudaMalloc、cudaMemcpy、cudaFree
5线程索引计算blockIdx、threadIdx、多维索引
6并行计算模式向量加法、矩阵乘法
7同步与共享内存__shared__、__syncthreads()
8性能优化内存合并、避免分支分歧

第一课：CUDA 基础概念

知识点

什么是 CUDA？

CUDA= Compute Unified Device Architecture（统一计算设备架构）

  • NVIDIA 开发的并行计算平台和编程模型
  • 让程序员可以使用 C/C++ 直接编程 GPU
  • 实现 CPU（主机）和 GPU（设备）的协同工作
异构计算模型
┌─────────────────────────────────────────────────┐
│ 程序执行流程 │
├─────────────────────────────────────────────────┤
│ │
│ CPU（主机） GPU（设备） │
│ ├─ 串行代码 ├─ 并行代码 │
│ ├─ 逻辑控制 ├─ 大量计算 │
│ └─ 内存管理 └─ 数据并行 │
│ │
│ CPU 准备数据 ──→ 传输到 GPU ──→ GPU 计算 │
│ CPU 接收结果 ←── 传回 CPU ←── GPU 完成 │
│ │
└─────────────────────────────────────────────────┘
CPU vs GPU 对比
特性CPUGPU
核心数少（4-64）多（数千）
时钟频率高（3-5 GHz）低（1-2 GHz）
缓存大小
适用场景串行、逻辑控制并行、数值计算
内存主机内存（Host Memory）设备内存（Device Memory）
CUDA 程序基本结构
int main() {
// 1. 分配主机内存
float *h_data = (float*)malloc(size);
// 2. 分配设备内存
float *d_data;
cudaMalloc(&d_data, size);
// 3. 拷贝数据到设备
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
// 4. 启动内核（并行计算）
myKernel<<<grid, block>>>(d_data);
// 5. 拷贝结果回主机
cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);
// 6. 释放内存
cudaFree(d_data);
free(h_data);
return 0;
}

练习题 1

题号问题正确答案
1CUDA 的全称是什么？Compute Unified Device Architecture（统一计算设备架构）
2CPU 和 GPU 分别适合什么类型的任务？CPU：串行、逻辑控制、复杂决策
GPU：并行、数值计算、大规模数据处理
3CUDA 程序的基本执行流程？1. 分配内存（主机+设备）
2. 数据传输到设备
3. 启动内核并行计算
4. 结果传回主机
5. 释放内存

第二课：线程层级结构

知识点

三级线程层级

CUDA 使用三级线程层级来组织并行执行：

┌─────────────────────────────────────────────────┐
│ Grid（网格） │
│ 一个内核启动产生一个 Grid │
│ │
│ ┌─────────────┐ ┌─────────────┐ ┌─────────┐ │
│ │ Block 0 │ │ Block 1 │ │ Block N │ │
│ │ 线程块 │ │ 线程块 │ │ 线程块 │ │
│ │ │ │ │ │ │ │
│ │ T0 T1 T2... │ │ T0 T1 T2... │ │ T0 T1.. │ │
│ │ 线程 │ │ 线程 │ │ 线程 │ │
│ └─────────────┘ └─────────────┘ └─────────┘ │
│ │
│ Grid = 所有 Block 的集合 │
│ Block = 一组 Thread 的集合 │
│ Thread = 最小执行单元 │
└─────────────────────────────────────────────────┘
层级说明
层级说明数量限制
Grid网格，一个内核启动的所有线程最多 2^31-1 个 Block
Block线程块，一组可协作的线程最多 1024 个 Thread
Thread线程，最小执行单元执行一个内核函数实例
Block 和 Grid 的维度

可以是一维、二维或三维：

// 一维
dim3 grid(10); // 10 个 Block
dim3 block(256); // 每个 Block 256 个 Thread
// 二维
dim3 grid(10, 10); // 10x10 = 100 个 Block
dim3 block(16, 16); // 每个 Block 16x16 = 256 个 Thread
// 三维
dim3 grid(10, 10, 10); // 10x10x10 = 1000 个 Block
dim3 block(8, 8, 8); // 每个 Block 8x8x8 = 512 个 Thread
为什么使用多维？
  • 一维：处理数组、向量
  • 二维：处理图像、矩阵
  • 三维：处理体积数据、3D 模型

练习题 2

题号问题正确答案
1CUDA 的三级线程层级是什么？Grid（网格）→ Block（线程块）→ Thread（线程）
2一个 Block 最多能有多少个 Thread？1024 个
3为什么需要多维的 Block 和 Grid？不同维度的数据结构：一维处理数组，二维处理图像/矩阵，三维处理体积数据

第三课：内核函数

知识点

函数修饰符

CUDA 使用特殊修饰符区分不同类型的函数：

修饰符执行位置调用位置说明
__global__GPUCPU 或 GPU内核函数，启动并行执行
__device__GPUGPU设备函数，只能被 GPU 调用
__host__CPUCPU主机函数，普通 C/C++ 函数
内核函数定义
// __global__ 内核函数
__global__ void myKernel(int *data, int n) {
int idx = threadIdx.x; // 获取线程索引
if (idx < n) {
data[idx] = data[idx] * 2; // 每个线程处理一个元素
}
}
内核启动语法
// 启动内核
myKernel<<<gridSize, blockSize>>>(参数列表);
// gridSize: Block 数量（可以是 int 或 dim3）
// blockSize: 每个 Block 的 Thread 数量（可以是 int 或 dim3）
示例
// 一维启动
int gridSize = 10; // 10 个 Block
int blockSize = 256; // 每个 Block 256 个 Thread
myKernel<<<gridSize, blockSize>>>(d_data, n);
// 二维启动
dim3 grid(10, 10); // 10x10 个 Block
dim3 block(16, 16); // 每个 Block 16x16 个 Thread
myKernel<<<grid, block>>>(d_matrix, width, height);
内核函数限制
  • 不能有返回值（必须是 void）
  • 不能使用递归
  • 不能使用静态变量
  • 不能使用可变参数

练习题 3

题号问题正确答案
1__global__修饰符的作用是什么？定义内核函数，在 GPU 上执行，可从 CPU 或 GPU 调用
2写出启动内核addKernel的语法，使用 5 个 Block，每个 Block 128 个 ThreadaddKernel<<<5, 128>>>(参数);
3内核函数有哪些限制？不能有返回值、不能递归、不能使用静态变量、不能使用可变参数

第四课：内存管理

知识点

CUDA 内存类型
内存类型位置作用
主机内存（Host Memory）CPU存储输入数据和接收结果
设备内存（Device Memory）GPU存储 GPU 计算数据
共享内存（Shared Memory）GPU Block 内Block 内线程共享，高速
常量内存（Constant Memory）GPU只读，缓存优化
全局内存（Global Memory）GPU所有线程可访问
基本内存操作函数
// 1. 分配设备内存
cudaMalloc(void **ptr, size_t size);
// 2. 拷贝数据
cudaMemcpy(void *dst, void *src, size_t size, cudaMemcpyKind kind);
// 3. 释放设备内存
cudaFree(void *ptr);
cudaMemcpyKind 类型
类型方向
cudaMemcpyHostToDevice主机 → 设备
cudaMemcpyDeviceToHost设备 → 主机
cudaMemcpyDeviceToDevice设备 → 设备
完整示例
int main() {
int n = 1000;
size_t size = n * sizeof(float);
// 1. 分配主机内存
float *h_a = (float*)malloc(size);
float *h_b = (float*)malloc(size);
float *h_c = (float*)malloc(size);
// 初始化数据
for (int i = 0; i < n; i++) {
h_a[i] = i;
h_b[i] = i * 2;
}
// 2. 分配设备内存
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
// 3. 拷贝数据到设备
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
// 4. 启动内核
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
// 5. 拷贝结果回主机
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
// 6. 释放内存
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
free(h_a);
free(h_b);
free(h_c);
return 0;
}
错误检查
// 检查 CUDA 错误
cudaError_t err = cudaMalloc(&d_data, size);
if (err != cudaSuccess) {
printf("CUDA 错误: %s\n", cudaGetErrorString(err));
}

练习题 4

题号问题正确答案
1写出分配 1000 个 float 的设备内存的代码float *d_data;
cudaMalloc(&d_data, 1000 * sizeof(float));
2cudaMemcpy 的四个参数分别是什么？目标地址、源地址、数据大小、传输方向
3为什么需要 cudaFree？释放设备内存，避免内存泄漏

第五课：线程索引计算

知识点

内置变量

CUDA 提供内置变量来获取线程索引：

变量说明类型
blockIdx当前 Block 在 Grid 中的索引dim3
threadIdx当前 Thread 在 Block 中的索引dim3
blockDimBlock 的维度（Thread 数量）dim3
gridDimGrid 的维度（Block 数量）dim3
一维索引计算
// 全局线程索引（一维）
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// 示例：
// blockIdx.x = 2, blockDim.x = 256, threadIdx.x = 10
// idx = 2 * 256 + 10 = 522
二维索引计算
// 全局线程索引（二维）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
// 转换为一维索引（处理矩阵）
int idx = y * width + x;
三维索引计算
// 全局线程索引（三维）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
// 转换为一维索引
int idx = z * height * width + y * width + x;
边界检查
__global__ void myKernel(float *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// 必须检查边界！
if (idx < n) {
data[idx] = data[idx] * 2;
}
}

为什么需要边界检查？

  • 数据大小可能不是 blockSize 的整数倍
  • 多余的线程不应该访问无效内存
计算 Grid 大小
// 确保 Grid 大小足够覆盖所有数据
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize; // 向上取整
// 示例：
// n = 1000, blockSize = 256
// gridSize = (1000 + 255) / 256 = 4
// 总线程数 = 4 * 256 = 1024 > 1000 ✓

练习题 5

题号问题正确答案
1写出一维全局线程索引的计算公式int idx = blockIdx.x * blockDim.x + threadIdx.x;
2blockIdx.x=3, blockDim.x=128, threadIdx.x=50，全局索引是多少？idx = 3 * 128 + 50 = 434
3为什么需要边界检查？数据大小可能不是 blockSize 的整数倍，多余线程不应访问无效内存

第六课：并行计算模式

知识点

向量加法
// 内核函数
__global__ void vectorAdd(float *a, float *b, float *c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
// 主函数
int main() {
int n = 1000;
size_t size = n * sizeof(float);
// 分配和初始化主机内存...
// 分配设备内存
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
// 拷贝数据
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
// 启动内核
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
// 拷贝结果
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
// 释放内存...
}
矩阵乘法
// 内核函数（简单版本）
__global__ void matrixMul(float *A, float *B, float *C, int width) {
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < width && col < width) {
float sum = 0.0f;
for (int k = 0; k < width; k++) {
sum += A[row * width + k] * B[k * width + col];
}
C[row * width + col] = sum;
}
}
// 启动内核
dim3 block(16, 16);
dim3 grid((width + 15) / 16, (width + 15) / 16);
matrixMul<<<grid, block>>>(d_A, d_B, d_C, width);
并行计算模式总结
模式特点应用场景
元素级并行每个线程处理一个元素向量加法、标量乘法
行/列并行每个线程处理一行/列矩阵操作
归约多线程协作合并结果求和、最大值、最小值
扫描计算前缀和排序、压缩

练习题 6

题号问题正确答案
1向量加法中，每个线程做什么？处理一个元素：c[idx] = a[idx] + b[idx]
2矩阵乘法中，row 和 col 如何计算？row = blockIdx.y * blockDim.y + threadIdx.y
col = blockIdx.x * blockDim.x + threadIdx.x
3为什么矩阵乘法使用二维 Block？矩阵是二维结构，二维 Block 更直观地映射到矩阵元素

第七课：同步与共享内存

知识点

共享内存

共享内存是 Block 内所有线程共享的高速内存：

__global__ void myKernel(float *data) {
// 声明共享内存
__shared__ float sharedData[256];
int idx = threadIdx.x;
// 每个线程加载一个元素到共享内存
sharedData[idx] = data[idx];
// 同步，确保所有线程都完成加载
__syncthreads();
// 使用共享内存计算
data[idx] = sharedData[idx] * 2;
}
共享内存特点
特性说明
位置GPU 片上，位于 Block 内
速度比全局内存快约 100 倍
大小每个 Block 最多 48KB（可配置）
生命周期Block 执行期间
可见性仅 Block 内线程可见
同步函数
__syncthreads(); // Block 内所有线程同步

作用：

  • 确保所有线程都执行到同一位置
  • 用于共享内存读写同步

注意：

  • 只能在 Block 内同步
  • 不能在条件分支中调用（可能导致死锁）
共享内存应用：矩阵乘法优化
__global__ void matrixMulShared(float *A, float *B, float *C, int width) {
__shared__ float As[16][16];
__shared__ float Bs[16][16];
int row = blockIdx.y * 16 + threadIdx.y;
int col = blockIdx.x * 16 + threadIdx.x;
float sum = 0.0f;
// 分块计算
for (int k = 0; k < width; k += 16) {
// 加载到共享内存
As[threadIdx.y][threadIdx.x] = A[row * width + k + threadIdx.x];
Bs[threadIdx.y][threadIdx.x] = B[(k + threadIdx.y) * width + col];
__syncthreads();
// 计算部分结果
for (int i = 0; i < 16; i++) {
sum += As[threadIdx.y][i] * Bs[i][threadIdx.x];
}
__syncthreads();
}
C[row * width + col] = sum;
}

练习题 7

题号问题正确答案
1共享内存用什么关键字声明？shared
2共享内存比全局内存快多少倍？约 100 倍
3__syncthreads() 的作用是什么？Block 内所有线程同步，确保都执行到同一位置
4为什么 __syncthreads() 不能在条件分支中调用？可能导致部分线程跳过同步，造成死锁

第八课：性能优化

知识点

1. 内存合并（Memory Coalescing）

概念：相邻线程访问相邻内存地址时，GPU 可以合并为一次内存访问。

// 好的访问模式（合并）
int idx = blockIdx.x * blockDim.x + threadIdx.x;
data[idx] = value; // 线程 0 访问地址 0，线程 1 访问地址 1...
// 坏的访问模式（不合并）
int idx = threadIdx.x * stride; // 线程 0 访问地址 0，线程 1 访问地址 stride...
data[idx] = value;
2. 避免分支分歧（Branch Divergence）

概念：同一个 Warp（32 个线程）内的线程执行不同分支时，会串行执行。

// 坏的分支（分歧）
if (threadIdx.x % 2 == 0) {
// 偶数线程执行
} else {
// 奇数线程执行
}
// 好的分支（无分歧）
if (threadIdx.x < 16) {
// 前 16 个线程执行
} else {
// 后 16 个线程执行
}
3. 使用共享内存减少全局内存访问
// 不使用共享内存（多次访问全局内存）
for (int i = 0; i < N; i++) {
sum += data[idx + i]; // 每次都访问全局内存
}
// 使用共享内存（一次加载，多次使用）
__shared__ float sharedData[BLOCK_SIZE];
sharedData[threadIdx.x] = data[idx];
__syncthreads();
for (int i = 0; i < N; i++) {
sum += sharedData[i]; // 访问共享内存
}
4. 合适的 Block 大小
  • 通常是 128、256 或 512
  • 需要考虑：
    • 寄存器使用量
    • 共享内存使用量
    • Warp 数量（建议至少 2 个 Warp = 64 个线程）
5. 使用 CUDA 事件计时
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
myKernel<<<grid, block>>>(...);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("执行时间: %.2f ms\n", milliseconds);
cudaEventDestroy(start);
cudaEventDestroy(stop);

性能优化总结

优化方法说明效果
内存合并相邻线程访问相邻地址提高内存带宽利用率
避免分支分歧Warp 内线程执行相同分支减少串行执行
使用共享内存减少全局内存访问大幅提升速度
合适 Block 大小128/256/512平衡资源使用
减少数据传输最小化 Host-Device 传输减少传输延迟

练习题 8

题号问题正确答案
1什么是内存合并？相邻线程访问相邻内存地址时，GPU 合并为一次访问
2Warp 的大小是多少？32 个线程
3为什么分支分歧会影响性能？Warp 内线程执行不同分支时会串行执行，降低并行效率
4共享内存为什么能提高性能？比全局内存快约 100 倍，减少全局内存访问次数

综合练习：完整 CUDA 程序

向量加法完整代码

#include <stdio.h>
// 内核函数
__global__ void vectorAdd(float *a, float *b, float *c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
int main() {
int n = 10000;
size_t size = n * sizeof(float);
// 1. 分配主机内存
float *h_a = (float*)malloc(size);
float *h_b = (float*)malloc(size);
float *h_c = (float*)malloc(size);
// 初始化数据
for (int i = 0; i < n; i++) {
h_a[i] = (float)i;
h_b[i] = (float)(i * 2);
}
// 2. 分配设备内存
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
// 3. 拷贝数据到设备
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
// 4. 启动内核
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
// 5. 拷贝结果回主机
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
// 6. 验证结果
bool success = true;
for (int i = 0; i < n; i++) {
if (h_c[i] != h_a[i] + h_b[i]) {
success = false;
break;
}
}
printf("验证结果: %s\n", success ? "成功" : "失败");
// 7. 释放内存
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
free(h_a);
free(h_b);
free(h_c);
return 0;
}

学习检查清单

基础概念 ✅

  • 理解 CUDA 是什么
  • 理解 CPU-GPU 异构计算模型
  • 掌握 CUDA 程序基本结构

线程组织 ✅

  • 理解 Grid-Block-Thread 层级
  • 掌握 Block 和 Grid 维度设置
  • 理解 Block 大小限制

内核函数 ✅

  • 掌握__global__、__device__修饰符
  • 掌握内核启动语法<<<grid, block>>>
  • 了解内核函数限制

内存管理 ✅

  • 掌握 cudaMalloc、cudaMemcpy、cudaFree
  • 理解 Host-Device 数据传输
  • 掌握错误检查方法

线程索引 ✅

版权声明: 本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若内容造成侵权/违法违规/事实不符，请联系邮箱：809451989@qq.com进行投诉反馈，一经查实，立即删除！
网站建设 2026/7/2 2:52:58

SpringBoot开发实践

SpringBoot开发实践&#xff1a;从“约定大于配置”到高效微服务在Java企业级开发的演进历程中&#xff0c;SpringBoot无疑是一道分水岭。它不仅仅是一个框架的升级&#xff0c;更代表了一种开发哲学的转变——从繁琐的XML配置地狱到“约定大于配置”的优雅实践。本文将深入探讨…

作者头像 李华
网站建设 2026/7/2 2:49:36

Python高级异步编程实战技巧与最佳实践

Python高级异步编程实战技巧与最佳实践在当今高并发的互联网应用场景中&#xff0c;异步编程已成为Python开发者必须掌握的核心技能。从Web服务到数据处理&#xff0c;从网络爬虫到实时通信&#xff0c;异步编程范式通过非阻塞I/O操作显著提升了程序性能。本文将深入探讨Python…

作者头像 李华
网站建设 2026/7/2 2:48:24

前端工程化构建工具链配置实战教程

前端工程化构建工具链配置实战教程前端工程化已成为现代Web开发的标配&#xff0c;它通过自动化流程提升开发效率、保障代码质量。本文将带你从零开始配置一套完整的前端工程化工具链&#xff0c;涵盖开发、构建、测试到部署的全流程。一、环境初始化与包管理首先确保已安装Nod…

作者头像 李华
网站建设 2026/7/2 2:45:48

那些年搞不懂的术语、概念：协变、逆变、不变体

简述什么是协变性、逆变性、不变性 协变性&#xff0c;如&#xff1a;string->object &#xff08;子类到父类的转换&#xff09;逆变性&#xff0c;如&#xff1a;object->string &#xff08;父类到子类的转换&#xff09;不变性&#xff0c;基于上面两种情况&#xf…

作者头像 李华
网站建设 2026/7/2 2:45:02

实战测评：PixsoAI与FigmaAI生成UI界面的优缺点与差异化

里&#xff0c;AI到底能不能帮我们提速&#xff1f; 正好最近项目不忙&#xff0c;我抽出时间上手跑了几个典型需求。这篇就主要聊聊&#xff0c;实操下来AI生成UI到底能做到什么程度&#xff1f;PixsoAI和Figma MakeDesign&#xff0c;到底有什么相同点和差异化。实测一&#…

作者头像 李华
网站建设 2026/7/2 2:44:26

堆排序实现教程

堆排序算法详解&#xff1a;从原理到实现一、什么是堆排序&#xff1f;堆排序&#xff08;Heap Sort&#xff09;是一种基于二叉堆数据结构的比较排序算法&#xff0c;由J.W.J. Williams于1964年发明。它结合了插入排序和归并排序的优点&#xff0c;具有原地排序&#xff08;只…

作者头像 李华