CUDA 第一个Hello World程序

• 1. 源码及解析
  • 1.1. 源代码
  • 1.2. 编译和运行
• 2. 代码解析
  • 2.1. 核函数定义
  • 2.2. 内核启动
  • 2.3. 同步操作
• 3. 执行空间限定符
  • 3.1. __global__
  • 3.2. __device__
  • 3.3. __host__
  • 3.4. 组合使用
  • 3.5. 总结与对比表
• 4. 线程层次结构
  • 4.1. Grid、Block、Thread
    • 4.1.1. 线程 (Thread) - 最基本的执行单元
    • 4.1.2. 线程块 (Block) - 线程的集合
    • 4.1.3. 网格 (Grid) - 线程块的集合
  • 4.2. 线程模型示意图
  • 4.3. 内置坐标变量
    • 4.3.1. 坐标变量说明
    • 4.3.2. Demo示例
• 5. 总结

《CUDA开发环境搭建》一章中我们写了一个Hello World程序来验证开发环境，现在我们从这一个Demo开始入手，开启我们的CUDA编程之旅吧！

1. 源码及解析

1.1. 源代码

#include <stdio.h>

__global__ void helloFromGPU()
{
    printf("Hello World from GPU!\n");
}

int main()
{
    printf("Hello World from CPU!\n");
    helloFromGPU<<<1, 2>>>();
    cudaDeviceSynchronize();
    return 0;
}

1.2. 编译和运行

# 编译命令
nvcc hello.cu -o hello

# 运行命令
./hello
Hello World from CPU!
Hello World from GPU!
Hello World from GPU!

2. 代码解析

会发现整个代码和C语言几乎一样，只有几个特殊的地方，这里展开讲解。

2.1. 核函数定义

__global__ void helloFromGPU()

• __global__关键字修饰的函数叫内核函数(kernel)，也经常简称核函数。核函数是CUDA并行计算的入口点，是GPU上大量线程并行执行的代码。
• 核函数可以从主机(Host)调用并在设备(Device)上执行。
• 返回类型必须是void
• 核函数在GPU上并行执行，每个线程都会执行相同的代码
• __global__在CUDA中也叫执行空间限定符，除__global__外，还有其他一些关键字，参见后面的内容。

2.2. 内核启动

helloFromGPU<<<1, 2>>>();

• <<<1, 2>>> 是CUDA特有的语法，称为"执行配置"
• 第一个参数 1 表示使用1个线程块
• 第二个参数 2 表示每个线程块有2个线程

2.3. 同步操作

cudaDeviceSynchronize();

• CPU和GPU是异步执行的，这条语句让CPU等待GPU上的所有操作完成。
• 确保所有GPU输出都能在程序结束前显示。

3. 执行空间限定符

3.1. __global__

• 含义： “全局” 函数。这是一个在 GPU 上执行但可以从 CPU 调用的函数。函数调用的空间横跨了CPU和GPU，所以称为__global__(全局的)。
• 功能： 用于定义内核(Kernel)。内核是 CUDA 并行计算的入口点，是 GPU 上大量线程并行执行的代码。
• 调用方式：
  • 使用特殊的尖括号语法 <<<grid_dim, block_dim>>> 从主机调用。
  • 返回值必须为 void。
  • 无法直接获取其返回值（需要通过设备内存传递结果）。

3.2. __device__

• 含义： “设备” 函数。它只能在 GPU 上执行，并且只能被 GPU 上的其他函数调用。
• 功能： 充当设备代码的辅助函数。用于封装一些在 __global__ 内核中重复使用的逻辑，使代码更模块化、更清晰。
• 调用方式：
  • 只能在其他 __device__ 或 __global__ 函数内部调用。
  • 无法从主机 CPU 直接调用。

// 一个设备端辅助函数，计算平方
__device__ float square(float x) {
    return x * x;
}

__global__ void anotherKernel(float *a, float *b) {
    int idx = ...; // 计算索引
    // 在内核中调用设备函数
    b[idx] = square(a[idx]); 
}

// 错误示例：不能在主机函数中调用 __device__ 函数
int main() {
    // square(2.0f); // 这行代码无法编译！
}

3.3. __host__

• 含义： “主机” 函数。它只能在 CPU 上执行，并且只能被 CPU 上的其他函数调用。
• 功能： 这就是标准的 C/C++ 函数。如果一个函数没有任何限定符，它默认就是 __host__ 函数。显式地写上 __host__ 通常是在使用组合限定符的场景，或者为了强调说明。
• 调用方式： 像普通函数一样调用。

// 以下两个函数是等价的
void hostFunction1() { ... }       // 默认是 __host__
__host__ void hostFunction2() { ... } // 显式声明

3.4. 组合使用

CUDA 允许将 __host__ 和 __device__ 限定符组合在一起使用。这告诉编译器：请为这个函数同时生成主机版本和设备版本。

• 功能： 用于编写既可以在 CPU 上运行，也可以在 GPU 上运行的函数。这对于代码复用和避免写两套相同的函数非常有用。
• 限制： 函数不能有复杂的标准库调用（如 printf, malloc），因为设备版本无法支持这些。只能使用 CUDA 支持的 C++ 语法和内置函数。

// 编译器会为该函数生成两个版本：一个在CPU上运行，一个在GPU上运行。
__host__ __device__ float calculateSomething(float a, float b) {
    return (a + b) * (a - b);
}

// 主机代码可以调用
void main() {
    float cpu_result = calculateSomething(5.0, 3.0); // 调用主机版本
}

// 设备代码也可以调用
__global__ void myKernel() {
    float gpu_result = calculateSomething(5.0, 3.0); // 调用设备版本
}

注意：你不能将 __global__ 与 __host__ 或 __device__ 组合使用，因为 __global__ 的定义（在 GPU 执行，从 CPU 调用）与它们的定义是冲突的。

3.5. 总结与对比表

为了更清晰地理解，我们可以从多个维度对它们进行对比：

特性	__global__	__device__	__host__
执行位置	GPU	GPU	CPU
调用者	CPU (通常) / GPU (动态并行)	GPU (其他设备函数)	CPU
返回值	必须为 void	任意类型	任意类型
主要用途	并行计算的入口点 (内核)	设备代码的辅助函数	普通的主机函数
调用语法	func<<<Dg, Db>>>(args)	func(args) (在设备代码中)	func(args)
可否从主机调用	可以 (使用 <<<>>>)	不可以 (编译错误)	可以
可否从设备调用	可以 (需CUDA动态并行)	可以	不可以 (编译错误)

说明： CUDA动态并行的概念会在后面的章节中讲解，暂时你可以先忽略他。

4. 线程层次结构

CUDA 的线程模型是其并行计算能力的核心，它采用了一种层次化的组织结构，使得程序员能够高效地管理和调度成千上万个并发线程。这个模型是理解 CUDA 编程的基础。

CUDA 的线程模型采用三层结构，从大到小分别是：网格 (Grid) → 线程块 (Block) → 线程 (Thread)。

4.1. Grid、Block、Thread

4.1.1. 线程 (Thread) - 最基本的执行单元

• 角色：最小的执行单位，每个线程独立执行相同的内核代码
• 特点：
  • 拥有自己的程序计数器、寄存器组
  • 执行相同的指令，但处理不同的数据（SIMT架构）
  • 通过内置变量区分彼此的身份

4.1.2. 线程块 (Block) - 线程的集合

• 角色：一组协同工作的线程的容器
• 重要特性：
  • 共享内存：块内的线程可以通过 __shared__ 内存高效通信
  • 同步机制：块内线程可以使用 __syncthreads() 进行同步
  • 线程限制：一个线程块中的所有线程都需驻留在同一个流式多处理器核心(SM)上，因此每个线程块中的线程数量是有限制的，限制大小因GPU的交互而异，目前最大不超过1024个线程。
  • 维度：可以组织成一维、二维或三维结构

4.1.3. 网格 (Grid) - 线程块的集合

• 角色：所有线程块的容器，构成一个完整的内核启动
• 特点：
  • 包含一个或多个线程块
  • 块间通常不需要同步（除非使用较新的特性）

4.2. 线程模型示意图

图1：线程块网格

在核函数调用时，一个核函数可以由多个形状相同的线程块执行。因此，线程总数 = 每个线程块中的线程数 * 线程块的数量。

如：

// 1个block, 每个block2个thread，线程总数 = 1 * 2 = 1
helloFromGPU<<<1, 2>>>();

// 2个block, 每个block5个thread，线程总数 = 2 * 5 = 10
helloFromGPU<<<2, 5>>>();

4.3. 内置坐标变量

4.3.1. 坐标变量说明

CUDA 提供了一系列内置变量用于标识线程的位置：

变量	描述	维度
threadIdx.x, .y, .z	线程在块内的索引	3D
blockIdx.x, .y, .z	块在网格内的索引	3D
blockDim.x, .y, .z	块的维度（每个方向的线程数）	3D
gridDim.x, .y, .z	网格的维度（每个方向的块数）	3D

说明：

• block的索引和thread的索引可以组织成一维、二维或三维的结构。
• 一维时只会使用threadIdx.x，二维时会使用threadIdx.x和threadIdx.y，三维时会使用threadIdx.x、threadIdx.y和threadIdx.z。
• 现在我们只需要先理解一维的场景，更复杂的场景会在后面的章节中讲解。

计算全局线程ID：

对于一维的网格和块，计算全局线程ID的公式为：

int globalId = blockIdx.x * blockDim.x + threadIdx.x;

4.3.2. Demo示例

我们修改一下上面的Demo如下：

// version 2.0
#include <stdio.h>

__global__ void helloFromGPU()
{
    int globalId = blockIdx.x * blockDim.x + threadIdx.x;
    printf("Hello World from GPU! blockIdx:%d, threadIdx:%d, globalId:%d\n", blockIdx.x, threadIdx.x, globalId);
}

int main()
{
    printf("Hello World from CPU!\n");
    helloFromGPU<<<2, 3>>>();
    cudaDeviceSynchronize();
    return 0;
}

执行结果如下：

# 编译
nvcc ./hello.cu -o ./hello

# 运行
./hello                   
Hello World from CPU!
Hello World from GPU! blockIdx:1, threadIdx:0, globalId:3
Hello World from GPU! blockIdx:1, threadIdx:1, globalId:4
Hello World from GPU! blockIdx:1, threadIdx:2, globalId:5
Hello World from GPU! blockIdx:0, threadIdx:0, globalId:0
Hello World from GPU! blockIdx:0, threadIdx:1, globalId:1
Hello World from GPU! blockIdx:0, threadIdx:2, globalId:2

5. 总结

到目前为止，我们通过Hello World的程序，了解了CUDA编程最核心的两个概念：执行空间限定符和线程层次结构模型。