课程大纲课次主题重点内容1CUDA 基础概念GPU 架构、异构计算模型2线程层级结构Grid、Block、Thread3内核函数__global__、__device__、启动语法4内存管理cudaMalloc、cudaMemcpy、cudaFree5线程索引计算blockIdx、threadIdx、多维索引6并行计算模式向量加法、矩阵乘法7同步与共享内存__shared__、__syncthreads()8性能优化内存合并、避免分支分歧第一课CUDA 基础概念知识点什么是 CUDACUDA 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 MemoryCUDA 程序基本结构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. 启动内核并行计算myKernelgrid, 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 个 BlockBlock线程块一组可协作的线程最多 1024 个 ThreadThread线程最小执行单元执行一个内核函数实例Block 和 Grid 的维度可以是一维、二维或三维// 一维dim3 grid(10); // 10 个 Blockdim3 block(256); // 每个 Block 256 个 Thread// 二维dim3 grid(10, 10); // 10x10 100 个 Blockdim3 block(16, 16); // 每个 Block 16x16 256 个 Thread// 三维dim3 grid(10, 10, 10); // 10x10x10 1000 个 Blockdim3 block(8, 8, 8); // 每个 Block 8x8x8 512 个 Thread为什么使用多维一维处理数组、向量二维处理图像、矩阵三维处理体积数据、3D 模型练习题 2题号问题正确答案1CUDA 的三级线程层级是什么Grid网格→ Block线程块→ Thread线程2一个 Block 最多能有多少个 Thread1024 个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; // 每个线程处理一个元素}}内核启动语法// 启动内核myKernelgridSize, blockSize(参数列表);// gridSize: Block 数量可以是 int 或 dim3// blockSize: 每个 Block 的 Thread 数量可以是 int 或 dim3示例// 一维启动int gridSize 10; // 10 个 Blockint blockSize 256; // 每个 Block 256 个 ThreadmyKernelgridSize, blockSize(d_data, n);// 二维启动dim3 grid(10, 10); // 10x10 个 Blockdim3 block(16, 16); // 每个 Block 16x16 个 ThreadmyKernelgrid, block(d_matrix, width, height);内核函数限制不能有返回值必须是 void不能使用递归不能使用静态变量不能使用可变参数练习题 3题号问题正确答案1__global__修饰符的作用是什么定义内核函数在 GPU 上执行可从 CPU 或 GPU 调用2写出启动内核addKernel的语法使用 5 个 Block每个 Block 128 个 ThreadaddKernel5, 128(参数);3内核函数有哪些限制不能有返回值、不能递归、不能使用静态变量、不能使用可变参数第四课内存管理知识点CUDA 内存类型内存类型位置作用主机内存Host MemoryCPU存储输入数据和接收结果设备内存Device MemoryGPU存储 GPU 计算数据共享内存Shared MemoryGPU Block 内Block 内线程共享高速常量内存Constant MemoryGPU只读缓存优化全局内存Global MemoryGPU所有线程可访问基本内存操作函数// 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;vectorAddgridSize, 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 中的索引dim3threadIdx当前 Thread 在 Block 中的索引dim3blockDimBlock 的维度Thread 数量dim3gridDimGrid 的维度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.x3, blockDim.x128, threadIdx.x50全局索引是多少idx 3 * 128 50 4343为什么需要边界检查数据大小可能不是 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;vectorAddgridSize, 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);matrixMulgrid, 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.ycol blockIdx.x * blockDim.x threadIdx.x3为什么矩阵乘法使用二维 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共享内存用什么关键字声明shared2共享内存比全局内存快多少倍约 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概念同一个 Warp32 个线程内的线程执行不同分支时会串行执行。// 坏的分支分歧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);myKernelgrid, 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;vectorAddgridSize, 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 数据传输掌握错误检查方法线程索引 ✅