CUDA编程学习 (1)——CUDA C介绍
1. 内存分配和数据移动 API 函数
CUDA编程模型是一个异构模型,需要CPU和GPU协同工作。在CUDA中,host 和 devic e 是两个重要的概念,我们用host指代CPU及其内存,而用device指代GPU及其内存。

CUDA程序中既包含host程序,又包含device程序,它们分别在CPU和GPU上运行。同时,host与device之间可以进行通信,这样它们之间可以进行数据拷贝。典型的CUDA程序的执行流程如下:
- 分配host内存,并进行数据初始化;
- 分配device内存,并从host将数据拷贝到device上;
- 调用CUDA的核函数在device上完成指定的运算;
- 将device上的运算结果拷贝到host上;
- 释放device和host上分配的内存。
#include <cuda.h>
void vecAdd(float *h_A, float *h_B, float *h_C, int n){
int size = n* sizeof(float);
float *d_A, *d_B, *d_C;
// Part 1
// 为 A,B,C 分配 device 内存
// copy A 和 B 到 device 内存
// Part 2
// 内核启动代码 - 设备执行实际的矢量加法
// Part 3
// 从 device 内存 copy C
}

Device 代码可以:
- 写/读 per-thread register
- 写/读 all-shared global memory
Host 代码可以:
- 将数据传输到每个 grid global memory / 从每个 grid global memory 传输数据
1.1 Device Memory 管理 API

cudaMalloc()
-
在 device global memory 中分配一个对象
-
两个参数
- 分配对象指针的地址
- 分配对象的大小(以字节为单位)
cudaFree()
- 从 device global memory 中释放对象
- 一个参数
- 释放对象的指针
1.2 Host-Device 数据传输 API

cudaMemCpy()
-
内存数据传输
-
4个参数
- 目的地指针
- 源指针
- 复制的字节数
- 传输类型/方向
-
与 host 同步向 device 传输数据
1.3 矢量加法-显式内存管理

1.4 统一内存(Unified Memory)

cudaMallocManaged(void** ptr, size_t size)
-
统一内存是 CPU 和 GPU 共享的相同内存空间。这意味着,CPU 和 GPU 使用相同的地址空间,无需像传统的 CUDA 编程那样手动管理内存传输(如
cudaMemcpy()
)。- 它维持数据的单一副本,即不再需要独立的主机内存和设备内存副本。
-
CUDA-managed 数据
- 其中
ptr
是指向分配内存的指针,size
是所需内存的大小。 - 按需页面迁移:在GPU执行代码时,所需的数据会按需从主机(Host)迁移到设备(Device)。
- 其中
-
与
cudaMalloc()
和cudaFree()
兼容 -
可优化
- cudaMemAdvise(): 给CUDA提供建议,提示内存的使用模式(例如,数据主要在哪个处理器上使用)。
- cudaMemPrefetchAsync(): 允许程序员预取数据到设备或主机,减少内存访问延迟。
- cudaMemcpyAsync(): 在异步传输数据时,可以优化性能。
1.5 矢量加法-统一内存

1.6 检查 host 代码中的 API 错误
cudaError_t err = cudaMalloc((void **) &d_A, size);
if(err != cudaSuccess){
printf("%s in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__);
exit(EXIT_FAILURE);
}
2. 线程和内核函数
2.1 Kernel 线程层次结构
首先GPU上很多并行化的轻量级线程,kernel 在 device 上执行时实际上是启动很多线程。
一个 kernel 所启动的所有线程称为一个网格(grid),同一个网格上的线程共享相同的 global memory,grid 是线程结构的第一层次,而网格又可以分为很多 线程块(block),一个线程块里面包含很多线程,这是第二个层次。

2.2 数据并行 - 向量加法示例

2.3 CUDA 执行模型
- 异构 host(CPU)+ device(GPU)应用 C 程序
- host C 代码中的串行部件
- device SPMD 内核代码中的并行部分

2.4 并行 thread 阵列
- CUDA kernel 由 threads 组成的一个 grid(array)执行
- grid 中的所有 thread 运行相同的 kernel 代码(单程序多数据)
- 每个 thread 都有索引,用于计算内存地址和做出控制决定

2.4.1 线程块:可扩展的协作

- 将 thread array 划分为多个 block
- block 内的 thread 通过 共享内存(shared memory)、原子操作(atomic operations)和障碍同步(barrier synchronization)进行协作
- 不同 block 中的 thread 不会相互影响
2.4.2 block 索引和 thread 索引

-
每个 thread 使用索引来决定处理哪些数据
- blockIdx:1D, 2D or 3D (CUDA 4.0)
- threadIdx:1D, 2D or 3D
-
处理多维数据时简化内存寻址
- 图像处理
- 解决体积上的 PDEs
对于一个 2-dim 的
3. CUDA 工具包简介
3.1 NVCC Compiler
- 英伟达™(NVIDIA®)提供 CUDA-C 编译器,简单的编译命令为:
nvcc vectorAdd.cu -o vectorAdd
这将使用默认参数编译 vectorAdd.cu
,并输出一个叫 vectorAdd
的可执行文件。
-
nvcc
编译设备代码,然后将代码转发给 host 编译器(如 g++)。 -
可用于编译和链接 host 应用程序
两种源文件:
对于 CUDA 程序,我们通常有:
- .cu 文件:包含 CUDA C/C++ 代码,包括主程序代码和 CUDA kernels,需要
nvcc
编译。 - .cpp/.c 文件:只包含主程序代码,不含 CUDA kernels,可以用
gcc/g++
编译。
编译 CUDA 程序的一般步骤是:
- 用
nvcc
编译 .cu 文件,生成 .o 文件。 - 用
g++
编译 .cpp 文件,也生成 .o 文件。 - 再用
g++/nvcc
链接所有 .o 文件生成最终可执行文件。
nvcc -c demo.cu -o demo.o
g++ -c main.cpp -o main.o
nvcc demo.o main.o -o demo -lcudart
3.2 调试工具(Debugger)

- NVIDIA 提供的调试工具
- Nsight:这是一个用于 GPU 编程调试和优化的工具,专门设计用来分析 GPU 代码的性能问题,并进行代码调试。支持对 CUDA 和 OpenGL 等代码进行深入的分析。
- Nsight Systems:一个面向系统级性能分析的工具,可以帮助开发者查找瓶颈,了解程序在 CPU 和 GPU 之间的负载分布。
- CUDA-GDB:这是一个基于 GDB 的调试器,专门用于 CUDA 编程的调试工作,支持查看 GPU 上的内存状态、线程执行情况等。
- CUDA MEMCHECK:用于检测 CUDA 代码中的内存错误,包括非法内存访问、未初始化内存使用等,类似于 CPU 上的 valgrind 工具。
- 第三方调试工具
- ARM Forge:支持跨平台的高性能计算(HPC)调试工具,适用于调试大规模并行程序,包括CUDA应用。
- TotalView:另一个广泛使用的高性能并行调试器,适用于调试复杂的CUDA程序,提供对多核、多线程代码的深入支持。
3.3 性能分析工具(Profilers)

-
NVIDIA 提供的性能分析工具
- NSIGHT:与调试工具一样,Nsight 还可用于性能分析,帮助开发者分析 GPU 代码的瓶颈、内存使用、指令执行效率等。它支持 CUDA、OpenGL 和 Vulkan 等多种 API。
- NVVP(NVIDIA Visual Profiler):专门为 CUDA 程序设计的图形化性能分析工具,提供可视化的性能瓶颈分析,帮助开发者定位代码中导致性能下降的具体部分。
- NVPROF:是一个命令行性能分析工具,专注于 CUDA 程序的性能数据收集。它可以通过收集时间信息、内存传输、计算效率等指标,提供详细的性能分析结果。
-
第三方性能分析工具
- TAU:一款适用于并行程序的全面性能分析工具,支持多种硬件架构和 API,帮助用户从全局上评估程序的性能表现。
- VampirTrace:一个可视化的性能分析工具,用于追踪并行程序的执行情况,帮助开发者分析程序的时间线、通信延迟等。
4. Nsight 计算和 Nsight 系统

-
Phasing Out(逐步淘汰的工具)
- NVPROF
- NVVP(NVIDIA Visual Profiler)
-
Current(当前主流的工具)
- Nsight Compute
- Nsight Systems
参考文献
本文作者:Astron_fjh
本文链接:https://www.cnblogs.com/Astron-fjh/p/18503824
版权声明:本作品采用知识共享署名-非商业性使用-禁止演绎 2.5 中国大陆许可协议进行许可。
【推荐】编程新体验,更懂你的AI,立即体验豆包MarsCode编程助手
【推荐】凌霞软件回馈社区,博客园 & 1Panel & Halo 联合会员上线
【推荐】抖音旗下AI助手豆包,你的智能百科全书,全免费不限次数
【推荐】博客园社区专享云产品让利特惠,阿里云新客6.5折上折
【推荐】轻量又高性能的 SSH 工具 IShell:AI 加持,快人一步