CUDA编程学习 (1)——CUDA C介绍

1. 内存分配和数据移动 API 函数

CUDA编程模型是一个异构模型，需要CPU和GPU协同工作。在CUDA中，host 和 devic e 是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。

image-20240921213601913

CUDA程序中既包含host程序，又包含device程序，它们分别在CPU和GPU上运行。同时，host与device之间可以进行通信，这样它们之间可以进行数据拷贝。典型的CUDA程序的执行流程如下：

  1. 分配host内存，并进行数据初始化；
  2. 分配device内存，并从host将数据拷贝到device上；
  3. 调用CUDA的核函数在device上完成指定的运算；
  4. 将device上的运算结果拷贝到host上；
  5. 释放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
}
image-20240921221415897

Device 代码可以：

  • 写/读 per-thread register
  • 写/读 all-shared global memory

Host 代码可以：

  • 将数据传输到每个 grid global memory / 从每个 grid global memory 传输数据

1.1 Device Memory 管理 API

image-20240921223550002

cudaMalloc()

  • 在 device global memory 中分配一个对象

  • 两个参数

    • 分配对象指针的地址
    • 分配对象的大小（以字节为单位）

cudaFree()

  • 从 device global memory 中释放对象
  • 一个参数
    • 释放对象的指针

1.2 Host-Device 数据传输 API

image-20240921225324676

cudaMemCpy()

  • 内存数据传输

  • 4个参数

    • 目的地指针
    • 源指针
    • 复制的字节数
    • 传输类型/方向

  • 与 host 同步向 device 传输数据

1.3 矢量加法-显式内存管理

image-20240921230232238

1.4 统一内存（Unified Memory）

image-20240921231200588

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 矢量加法-统一内存

image-20240921234243047

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），一个线程块里面包含很多线程，这是第二个层次。

img

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

image-20240922003424983

2.3 CUDA 执行模型

  • 异构 host（CPU）+ device（GPU）应用 C 程序
    • host C 代码中的串行部件
    • device SPMD 内核代码中的并行部分
image-20240922003839725

2.4 并行 thread 阵列

  • CUDA kernel 由 threads 组成的一个 grid（array）执行
    • grid 中的所有 thread 运行相同的 kernel 代码（单程序多数据）
    • 每个 thread 都有索引，用于计算内存地址和做出控制决定
image-20240922124715191

2.4.1 线程块：可扩展的协作

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

2.4.2 block 索引和 thread 索引

image-20240922125856835

  • 每个 thread 使用索引来决定处理哪些数据

    • blockIdx：1D, 2D or 3D (CUDA 4.0)
    • threadIdx：1D, 2D or 3D

  • 处理多维数据时简化内存寻址

    • 图像处理
    • 解决体积上的 PDEs

对于一个 2-dim 的 \(block(D_x,D_y)\)，线程 \((x,y)\) 的 ID 值为 \((x+y∗D_x)\) ，如果是 3-dim 的 \(block(D_x,D_y,D_z)\)，线程 \((x,y,z)\) 的ID值为 \((x+y∗D_x+z∗D_x∗D_y)\) 。另外线程还有内置变量 gridDim，用于获得网格块各个维度的大小。

3. CUDA 工具包简介

3.1 NVCC Compiler

  • 英伟达™（NVIDIA®）提供 CUDA-C 编译器，简单的编译命令为：
nvcc vectorAdd.cu -o vectorAdd

这将使用默认参数编译 vectorAdd.cu，并输出一个叫 vectorAdd 的可执行文件。

  • nvcc 编译设备代码，然后将代码转发给 host 编译器（如 g++）。

  • 可用于编译和链接 host 应用程序

两种源文件：

对于 CUDA 程序，我们通常有：

  1. .cu 文件：包含 CUDA C/C++ 代码，包括主程序代码和 CUDA kernels，需要 nvcc 编译。
  2. .cpp/.c 文件：只包含主程序代码，不含 CUDA kernels，可以用 gcc/g++ 编译。

编译 CUDA 程序的一般步骤是：

  1. 用 nvcc 编译 .cu 文件，生成 .o 文件。
  2. 用 g++ 编译 .cpp 文件，也生成 .o 文件。
  3. 再用 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）

image-20240922172325240
  • 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）

image-20240922173202824

  • NVIDIA 提供的性能分析工具

    • NSIGHT：与调试工具一样，Nsight 还可用于性能分析，帮助开发者分析 GPU 代码的瓶颈、内存使用、指令执行效率等。它支持 CUDA、OpenGL 和 Vulkan 等多种 API。
    • NVVP（NVIDIA Visual Profiler）：专门为 CUDA 程序设计的图形化性能分析工具，提供可视化的性能瓶颈分析，帮助开发者定位代码中导致性能下降的具体部分。
    • NVPROF：是一个命令行性能分析工具，专注于 CUDA 程序的性能数据收集。它可以通过收集时间信息、内存传输、计算效率等指标，提供详细的性能分析结果。

  • 第三方性能分析工具

    • TAU：一款适用于并行程序的全面性能分析工具，支持多种硬件架构和 API，帮助用户从全局上评估程序的性能表现。
    • VampirTrace：一个可视化的性能分析工具，用于追踪并行程序的执行情况，帮助开发者分析程序的时间线、通信延迟等。

4. Nsight 计算和 Nsight 系统

image-20240922174403339

  • Phasing Out（逐步淘汰的工具）

    • NVPROF
    • NVVP（NVIDIA Visual Profiler）

  • Current（当前主流的工具）

    • Nsight Compute
    • Nsight Systems

参考文献

[1] CUDA编程入门极简教程 - 知乎 (zhihu.com)

[2] nvcc 编译 .cu 源代码的基本用法_多个文件 nvcc怎么编译-CSDN博客

posted @ 2024-10-26 11:44  Astron_fjh  阅读(22)  评论(0)  编辑  收藏  举报