CUDA编程之SPMD

杜子源源
5 阅读24分钟

对于CUDA编程来说，有一个至关重要的概念——Single-Program Multiple-Data（SPMD）。

SPMD：CPU与GPU协同的编程哲学

SPMD，单程序多数据流，这听起来似乎有些抽象，但它恰恰是理解CUDA如何组织和执行并行任务的核心。

“CUDA Integrated CPU+GPU application C program”

这意味着，我们通常所说的CUDA程序，并非是一个单纯运行在GPU上的孤立实体，而是一个异构程序，它巧妙地融合了在CPU上运行的串行代码和在GPU上运行的并行代码。

  • Serial C code executes on CPU: 这部分我们称之为主机代码（Host Code），负责程序的整体流程控制、数据准备、内存管理以及启动GPU任务等串行逻辑。
  • Parallel Kernel C code executes on GPU thread blocks: 这部分是设备代码（Device Code），也就是我们常说的核函数（Kernel）。同一个核函数会被GPU上的大量线程以并行的方式执行，每个线程处理数据集中的一部分。

何谓Single Program？

这里的“单程序”，并非指整个系统只有一个执行绪或一个简单的可执行文件。更准确地说，从开发者的视角看，我们通常在同一个源代码文件（例如.cu文件）中编写和管理CPU（主机）代码和GPU（设备）代码。

NVCC编译器扮演了关键角色，它能够智能地解析这些.cu文件：

  • 主机代码部分，它会交给系统标准的C/C++编译器（如GCC, MSVC）处理，编译成CPU可执行的指令。
  • 设备代码（Kernel函数），它会编译成GPU可执行的PTX中间代码或特定架构的机器码。

所以，“单程序”指的是开发者维护的一份统一的逻辑视图，其中包含了串行和并行的两个部分。

CPU扮演“指挥官”的角色，调度和管理任务；GPU则扮演“大规模并行处理器”的角色，高效执行计算密集型任务。SPMD模型使得开发者能够使用一份统一的内核代码，在GPU的数千个核心上同时处理海量数据，从而实现显著的性能加速。

理解SPMD模型，特别是其在Grid、Block、Thread层次上的组织方式以及CPU与GPU的协同工作流程，是高效进行CUDA编程的基础。

我们继续深入了解CUDA的编程模型。

现在是时候深入研究GPU上这些并行任务是如何被组织和管理的呢？这引出CUDA编程中两个核心的组织概念：网格(Grids)和线程块(Blocks)

GPU的执行层级结构

image.png

1. 内核执行单位：网格Grid

A Kernel is executed as a grid of thread blocks

当我们从CPU启动一个CUDA内核函数时，这个kernel实际上是在GPU上以一个Grid的形式执行的。一个Grid代表一次内核调用的全部并行工作量。你可以将其视为本次内核执行所涉及到的所有线程块的几何。

Host通过Kernel1的调用，在Device上创建了一个Grid1；类似地，Kernel2的调用创建了Grid2，每一个Grid都是一次独立的内核执行实例。

Grid可以是1D、2D或者3D的，这允许我们将线程组织映射到多维数据上。

当我们使用<<<Dg, Db, Ns, S>>>语法启动内核时：

/*
Dg: 定义了Grid的维度和大小 (多少个Block)
Db: 定义了每个Block的维度和大小 (每个Block多少个Thread)
Ns: 可选参数，动态分配的共享内存大小 (bytes per block)
S:  可选参数，关联的CUDA流
*/

// 示例1: 启动一个一维Grid，包含 numBlocks 个线程块
// 假设我们有 N 个元素需要处理，每个线程块处理 threadsPerBlock 个元素
int N = 1024 * 1024;
int threadsPerBlock = 256;
int numBlocks = (N + threadsPerBlock - 1) / threadsPerBlock; // 向上取整确保所有元素被处理

dim3 gridDim(numBlocks);         // 一维Grid，包含 numBlocks 个Block
dim3 blockDim(threadsPerBlock);  // 一维Block，每个Block包含 threadsPerBlock 个Thread

myKernel<<<gridDim, blockDim>>>(/* kernel arguments */);

// 示例2: 启动一个二维Grid，例如处理图像
int imageWidth = 1920, imageHeight = 1080;
int TILE_DIM_X = 16, TILE_DIM_Y = 16; // 每个Block处理16x16的tile
dim3 gridDim_2D( (imageWidth + TILE_DIM_X - 1) / TILE_DIM_X,
                 (imageHeight + TILE_DIM_Y - 1) / TILE_DIM_Y );
dim3 blockDim_2D(TILE_DIM_X, TILE_DIM_Y);
imageProcessingKernel<<<gridDim_2D, blockDim_2D>>>(/* ... */);

2. 协作单位：线程块Thread Block

A thread block is a batch of threads that can cooperate with each other by

线程块是Grid的组成单元，一个Grid由1个或者多个线程块组成。Grid1中被划分为多个黄色的Block，这些就是线程块。与Grid类似，Block内部的线程也可以组织成1D、2D、3D。

在同一个Thread Block内的线程具有特殊的协作能力： Synchronizing their execution using barrier使用屏障同步执行。

线程块内所有的线程可以通过调用__syncthreads()内建函数进行同步。当一个线程到达__syncthreads()时，它会暂停执行，直到该块内所有其他线程也都到达这个同步点，他们才一起执行。对于需要分阶段计算，且后一阶段依赖于前一阶段所有线程结果的场景非常关键。

__global__ void cooperativeKernel(float* data) {
    // ... 每个线程执行一些独立计算 ...
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    // data[tid] = some_initial_value;

    // 确保块内所有线程都完成了上述操作
    __syncthreads();

    // 现在可以安全地读取块内其他线程写入data的更新值（如果data在共享内存中）
    // 或者执行依赖于块内所有线程已完成某阶段任务的操作
    // if (threadIdx.x == 0) { /* 块内某个线程执行汇总操作 */ }
}

Efficiently sharing data through a low latency shared memory(通过低延迟的共享内存高效共享数据)。

每个线程块都拥有一块私有、低延时的片上共享内存（Shared Memory）。这块内存对于该块内所有线程都是可见的，并且访问速度远远快于全局内存。

块内线程可以将数据从全局内存加载到共享内存，进行高频次的读写操作，然后再将最终结果写回全局内存，从而大幅提升性能。__syncthreads()经常与共享内存配合使用，以确保数据在被其他线程读取之前被正确写入。

__global__ void sumReductionBlock(float *g_idata, float *g_odata, unsigned int n) {
    // 每个线程块处理一部分数据
    extern __shared__ float sdata[]; // 动态分配或静态声明共享内存

    unsigned int tid_in_block = threadIdx.x;
    unsigned int global_idx = blockIdx.x * blockDim.x + threadIdx.x;

    // 从全局内存加载数据到共享内存
    if (global_idx < n) {
        sdata[tid_in_block] = g_idata[global_idx];
    } else {
        sdata[tid_in_block] = 0; // padding for reduction
    }
    __syncthreads(); // 确保所有数据已加载到共享内存

    // 在共享内存中进行并行归约操作
    for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid_in_block < s) {
            sdata[tid_in_block] += sdata[tid_in_block + s];
        }
        __syncthreads(); // 确保每一轮归约完成
    }

    // 块内的第一个线程将结果写回全局内存
    if (tid_in_block == 0) {
        g_odata[blockIdx.x] = sdata[0];
    }
}

Two threads from two different blocks cannot cooperate directly.（来自不同块的两个线程不能直接协作）

这是一个非常重要的限制，__syncthreads()仅仅在块内有效，共享内存也仅对块内线程可见。不同线程块之间的线程不能直接访问彼此的共享内存。

如果他们需要同步或者数据交换，通常需要通过全局内存，并且可能需要将内核拆分成多个，或者依赖于原子操作等更复杂的机制。这种独立性使得线程块可以被GPU以任意顺序、在任意可用的流多处理器SM上调度执行，这是CUDA实现大规模可伸缩并行性的关键。

内存可见性

All tgreads share global memory space.

尽管块内有shared memory，但grid中所有线程（无论属于哪个块）都可以访问全局内存Global memory，全局内存通常是GPU显存的主要部分，容量较大，但是访问延迟也较高。

CPU通过CUDA API分配和传输数据到全局内存，内核函数也主要从全局内存读取输入数据，并将最终结果写回全局内存。

揭秘<<<...>>>

我们知道GPU拥有成百上千个核心，其优势在于同时处理海量数据，那么我们如何在设备上运行并行代码呢？

这是每一个CUDA初学者都会遇到的问题，我们已经定义了内核函数(例如 __global__ void myKernel(...))，但是如何指示GPU以并行的方式执行这个内核，而不是像传统CPU一样只执行一次？

解决方案就在于三个尖括号之间的参数，这正是CUDA C对C/C++的扩展，也是启动并行内核的“芝麻开门”咒语。

这个<<<...>>>语法，我们称之为执行配置Execution Configuration

从单次执行到N次并行

add<<<1, 1>>>(dev_a, dev_b, dev_c);

这里的 add 是一个我们定义的内核函数名，它接受三个设备指针 dev_a、dev_b、dev_c 作为参数。关键在于 <<< 1, 1 >>>：

  • 第一个 1 指定了要启动的线程块 (Blocks) 的数量。这里是1个线程块。
  • 第二个 1 指定了每个线程块中线程 (Threads) 的数量。这里是每个块1个线程。

因此，add<<<1, 1>>> 总共启动了 1 (Block) * 1 (Thread/Block) = 1 个线程。这意味着 add 内核函数实际上只会被执行一次，其行为与一个普通的串行函数调用类似（尽管它是在GPU上执行的）。

add<<< N, 1 >>>( dev_a, dev_b, dev_c );

与上一个例子相比，唯一的改变是执行配置中的第一个参数从 1 变成了 N（由绿色箭头高亮指出）。 <<< N, 1 >>>：

  • 第一个 N 指定了要启动 N 个线程块。
  • 第二个 1 仍然指定每个线程块包含1个线程。

因此，这次内核启动总共会创建 N (Blocks) * 1 (Thread/Block) = N 个线程。

核心变化：由于启动了 N 个线程，并且根据SPMD模型，每个线程都会独立执行 add 内核函数的代码，所以 add 内核实际上会并行地执行N次。

执行配置<<<Dg, Db, Ns, S>>>详解

想象一下，你是一位大将军CPU，要指挥一支庞大的军队GPU去攻克一个城池(完成一项计算任务)，需要明确告诉他们：

  1. 总共要分成多少部队(Blocks)
  2. 每个部队有多少士兵(Threads)

CUDA的执行配置<<<...>>>就是下达指令的方式。

标准的内核启动语法如下：

kernel_name <<< Dg, Db, Ns, S >>> (argument_list);

我们重点关注前两个参数Dg、Db，因此它们是控制并行度的核心，Ns和S是可选的，初学可以暂时使用默认值。

Dg：网格维度

Dg(DimGrid)指定了我们要启动的线程块Thread Blocks的总数量，以及它们在网格中的排列方式。

数据类型通常是一个dim3类型的变量。dim3是CUDA提供的一个结构体，可以用来表示1D、2D或者3D的维度。

  • dim3(x): 定义一个一维的网格，包含 x 个线程块。
  • dim3(x, y): 定义一个二维的网格，总共包含 x * y 个线程块，排列成 x 列 y 行。
  • dim3(x, y, z): 定义一个三维的网格，总共包含 x * y * z 个线程块。

逻辑意义: 这个参数决定了你的整个任务被划分成多少个“独立的工作包”（线程块）。每个线程块可以被GPU调度到任何一个可用的流多处理器（SM）上执行。

Db：块维度

Db (DimBlock) 指定了每一个线程块内部包含的线程 (Threads)  的数量以及它们在块内的排列方式。

数据类型: 同样通常是一个 dim3 类型的变量。

  • dim3(x): 定义每个块包含 x 个线程，一维排列。
  • dim3(x, y): 定义每个块包含 x * y 个线程，二维排列。
  • dim3(x, y, z): 定义每个块包含 x * y * z 个线程，三维排列。

逻辑意义: 这个参数决定了每个“独立工作包”内部有多少个“士兵”协同工作。同一个块内的线程可以高效地通过共享内存交换数据，并通过 __syncthreads()进行同步。

示例

让我们用一个非常简单的例子来演示：假设我们有8个数据元素需要处理，我们希望每个线程处理1个元素。

目标：启动8个线程来并行处理这两个元素。

思路：8个Block，每个Blkock1个thread；4Block2thred、2Block4thread、1Block8thread。

他们有什么区别呢？

稍后我们再讨论，我们先选择第四种方案——1个Block，每个Block8个thread。

#include <iostream>

// 内核函数：非常简单，每个线程打印自己的全局ID
__global__ void simpleKernel() {
    // 计算当前线程的全局唯一ID
    // 在这个简单例子中，由于只有一个Block，blockIdx.x 总是0
    // blockDim.x 是每个Block中的线程数 (这里是8)
    // threadIdx.x 是当前线程在Block内的ID (0到7)
    int globalThreadId = blockIdx.x * blockDim.x + threadIdx.x;

    printf("Hello from GPU! My Global ID is: %d (Block ID: %d, Thread ID in Block: %d)\n",
           globalThreadId, blockIdx.x, threadIdx.x);
}

int main() {
    // ----- 执行配置参数 -----
    // Dg: Grid Dimensions
    // 我们需要1个线程块
    dim3 numBlocks(1); // Dg.x = 1, Dg.y = 1, Dg.z = 1 (因为只传了一个参数)

    // Db: Block Dimensions
    // 每个线程块包含8个线程
    dim3 threadsPerBlock(8); // Db.x = 8, Db.y = 1, Db.z = 1

    printf("Host: Launching kernel with %u blocks and %u threads per block.\n",
           numBlocks.x, threadsPerBlock.x);
    printf("Host: Total threads to be launched: %u\n",
           numBlocks.x * threadsPerBlock.x);

    // 启动内核
    simpleKernel<<<numBlocks, threadsPerBlock>>>();

    // 因为内核中的printf是异步的，并且输出到主机控制台可能需要时间
    // 我们需要同步设备，以确保所有printf都执行完毕并且内容被刷新
    cudaError_t err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        printf("CUDA Error after kernel launch: %s\n", cudaGetErrorString(err));
    }

    printf("Host: Kernel execution finished.\n");

    return 0;
}

把这个代码放在以cu为结尾的文件中，之后使用nvcc test.cu进行编译，最终罪行代码。

image.png

输出结果如图所示。

现在我们同样可以改变配置

int main() {
    dim3 numBlocks(2);       // 2个线程块
    dim3 threadPerBlock(4);  // 4个线程
}

image.png 我们可以发现block内部的线程可以保证顺序一致，但是block之间是随机的。

Block和Thread如何设置？

回答我们刚刚的问题，这几种不同配置方案（8块x1线程、4块x2线程、2块x4线程、1块x8线程）在总线程数（都是8个）  和 每个线程完成的工作（假设内核逻辑相同，每个线程处理一个数据单元）  这两点上是相同的。然而，它们在GPU的执行方式、资源利用、以及潜在的性能特性上存在显著的区别。

这些区别主要源于线程块 (Thread Block)  在CUDA中的特殊地位和属性：

  1. 调度单位 (Scheduling Unit) :

    • 线程块是GPU流多处理器（SM）调度和分配的基本单位。GPU会将一个个线程块分配给可用的SM去执行。

    • 区别体现:

      • 8 Blocks x 1 Thread/Block: GPU会调度8个独立的块。如果SM数量少于8，一些块需要等待其他块执行完毕才能被调度。
      • 1 Block x 8 Threads/Block: GPU只需要调度1个块。这个块会被分配给一个SM。

  2. 资源分配单位 (Resource Allocation Unit) :

    • 共享内存 (Shared Memory)  是按线程块分配的。每个块有自己私有的共享内存。

    • 寄存器 (Registers)  也是一个关键资源，虽然是按线程分配，但SM上的总寄存器数量是有限的，一个SM能同时容纳的线程块数量会受到每个块内线程总数以及每个线程消耗寄存器数量的影响。

    • 区别体现:

      • 8 Blocks x 1 Thread/Block: 如果你的内核使用了共享内存，那么系统会为这8个块分别准备共享内存空间（尽管每个块内只有一个线程，可能用不上或用得很少）。如果每个块占用的共享内存很少，这可能不是问题。但如果每个块需要固定量的共享内存开销（即使只有一个线程），那么总的共享内存开销可能会比 1 Block x 8 Threads/Block 的情况高（尽管后者总的共享内存“需求”可能也高，但它集中在一个块内）。
      • 1 Block x 8 Threads/Block: 8个线程共享同一个块的共享内存。如果这8个线程需要协作，共享内存就非常有用。

  3. 协作范围 (Cooperation Scope) :

    • 同一个线程块内的线程可以进行同步 (__syncthreads()) 和高效的数据共享 (通过共享内存)。

    • 不同线程块之间的线程不能直接同步，也不能直接访问彼此的共享内存。

    • 区别体现:

      • 8 Blocks x 1 Thread/Block: 每个块只有一个线程，所以块内同步和共享内存的概念几乎没有意义。这8个线程是完全独立的，无法直接协作。
      • 1 Block x 8 Threads/Block: 这8个线程在同一个块内，它们可以非常方便地通过共享内存交换数据，并使用 __syncthreads() 来确保操作的顺序。如果你的算法需要这种紧密的协作（例如并行归约、卷积中的数据复用等），这种配置是必需的。

  4. 执行开销 (Overhead) :

    • 启动和管理每个线程块本身会有一些微小的开销。

    • 区别体现:

      • 8 Blocks x 1 Thread/Block: 启动8个块的开销可能会略高于启动1个块的开销，尽管这种差异通常很小，除非块的数量非常巨大且每个块的工作量极小。

  5. 占用率和并行度 (Occupancy and Parallelism) :

    • 占用率 (Occupancy)  是指一个SM上活跃的Warp（线程束，通常32个线程）数量与该SM理论上能支持的最大活跃Warp数量的比值。更高的占用率通常有助于隐藏访存延迟，提升性能，但并非绝对。

    • 每个SM能同时驻留的线程块数量和线程数量是有限的，受限于SM的物理资源（如寄存器数量、共享内存大小、最大线程数/块数限制）。

    • 区别体现:

      • 1 Block x 8 Threads/Block: 只有8个线程，这远少于一个Warp（32个线程）。这意味着即便这个块被调度到一个SM上，该SM的大部分计算能力也是空闲的。并行度非常低。

      • 8 Blocks x 1 Thread/Block: 类似地，每个块也只有一个线程，并行度也极低。即使8个块被同时调度到8个不同的SM（如果GPU有那么多空闲SM），每个SM也只运行一个线程。

      • 更实际的比较: 如果我们将总线程数扩大，比如处理1024个元素：

        • 1024 Blocks x 1 Thread/Block: 很多独立的块，可能无法充分利用每个SM内的并行能力。
        • 4 Blocks x 256 Threads/Block: 更为常见的配置。每个块有256个线程（即8个Warp），这有助于在一个SM内部实现较好的并行度，并能利用块内协作。GPU可以将这4个块分配到不同的SM上（或者同一个SM上分时执行，如果SM资源足够）。
        • 1 Block x 1024 Threads/Block: 如果算法允许并且资源足够（寄存器、共享内存），一个大块可以填满一个SM的一部分或全部Warp槽位。

  6. 可扩展性和灵活性 (Scalability and Flexibility) :

    • 使用更多的、规模适中的线程块通常能更好地适应不同规模的GPU硬件（有些GPU SM多，有些少）。

    • 区别体现:

      • 如果只有1个大块 (1 Block x 8 Threads/Block 的极端例子），那么你的并行性就受限于这一个块，它只能在一个SM上运行。如果你的问题规模远大于一个SM的处理能力，这种方式扩展性就很差。
      • 使用多个小块 (8 Blocks x 1 Thread/Block 的极端例子）虽然看似灵活，但每个块的效率太低。
      • 平衡是关键：通常选择每个线程块包含相当数量的线程（例如64、128、256、512个，通常是32的倍数以充分利用Warp），然后根据总任务量计算出合适的线程块数量。

总结表格对比

特性8 Blocks x 1 Thread/Block4 Blocks x 2 Threads/Block2 Blocks x 4 Threads/Block1 Block x 8 Threads/Block
调度单元8个独立调度4个独立调度2个独立调度1个独立调度
块内协作无意义 (单线程)2个线程可协作4个线程可协作8个线程可协作
共享内存利用若使用，每个块分配，可能浪费块内2线程共享块内4线程共享块内8线程共享，最适合块内数据共享算法
SM内并行度极低 (每个SM上最多跑1个活动线程，即使多个块在同一SM)较低较低较低 (总共也只有8个线程)
适应性可能无法充分利用现代GPU单个SM的强大计算能力同上同上同上，且只能利用一个SM

对于最初的8个线程的简单例子，这些配置在实际性能上可能差异不大，因为总工作量太小了。但是，当我们将问题规模放大时，这些选择的差异就会变得至关重要：

  • 如果你的算法是“高度并行且独立的任务”  (Embarrassingly Parallel)，比如每个线程独立计算 c[i] = a[i] + b[i]，那么使用适量线程块，每个线程块包含多个线程（比如 num_elements / threads_per_block 个块，每个块 threads_per_block 个线程，例如256）通常是好的。块太多、每个块线程太少（如 N Blocks x 1 Thread）会导致Warp利用率低。
  • 如果你的算法需要线程间的紧密协作（例如，在共享内存中进行局部数据的聚合、滤波等），那么你必须将协作的线程放在同一个线程块内。此时，1 Block x N Threads 或者 M Blocks x (N/M) Threads（其中 N/M 是协作单元的大小）会是更合适的选择。

GPU代码的“方言”

我们已经深入了解了SPMD的哲学，浏览了网格和块层次结构，甚至窥探了<<<..>>>的魔力.现在，让我们来谈谈细节：在编写注定要在强大的GPU上运行的代码时，我们到底能做什么，不能做什么？

GPU上执行的代码是“C/C++ with some restrictions”，这意味着你不能把所有在CPU上能跑的C++花活儿原封不动地搬到GPU上。为啥呢？因为GPU的架构和执行模型与CPU大相径庭，它为大规模并行而生，对某些串行或复杂特性支持有限。

C/C++ with some restrictions

1. can only access GPU memory (old generations)

在早期的CUDA版本中，GPU代码通常只能直接访问位于GPU显存的数据，CPU内存对它来说太远了，需要通过明确的cudaMemcpy等操作，将数据在CPU和GPU之间进行搬运。

值得注意的是，现在CPU通过统一虚拟寻址（Unified Virtual Addressing，UVA）和统一内存，也是我们之前的博客中提到的技术，极大地模糊了CPU和GPU的内存界限。在支持统一内存的系统上，PGU可以直接访问系统内存，或者由CUDA运行时透明地管理数据迁移。

2. No variable number of arguments 像C语言中 printf(const char* format, ...) 那样的可变参数函数，在GPU内核中是不支持的。内核函数的参数列表必须在编译时固定下来。

3. "No static variables" (没有静态变量)

在GPU内核函数内部或 __device__ 函数内部声明 static 局部变量是不允许的。这是因为静态变量通常具有跨函数调用保持其值或在整个编译单元内唯一的特性，这在GPU的大规模并行、多线程环境下难以有效管理和保证其行为的确定性。每个线程都是独立的执行路径。

注意: 文件作用域的 static __device__ 或 static __constant__ 变量是可以的，它们分别创建在全局内存或常量内存中，对Grid内的所有线程可见。这里主要指的是函数内部的局部静态变量。

4. "No recursion" (没有递归)

GPU内核不支持函数递归调用。GPU的硬件栈空间非常有限，递归很容易导致栈溢出。而且，递归的控制流对于SIMT（单指令多线程）执行模型来说效率不高。

5. "No dynamic polymorphism" (没有动态多态)*

C++中的虚函数 (virtual functions) 和动态类型转换 (like dynamic_cast) 这类依赖运行时类型信息（RTTI）的动态多态特性，在GPU内核中通常是不支持或支持受限的。这些特性会引入额外的开销和复杂的控制流，不适合GPU的执行模型。

GPU函数的“身份证”

仅仅遵守了上述“家规”还不够，你还需要给你的函数打上特殊的“标记”，告诉编译器这个函数是给谁（CPU还是GPU）用的，以及它应该在哪里执行。这就是“Must be declared with a qualifier”（必须用限定符声明）。

这些限定符就像函数的“身份证”，标明了它的“户籍”和“职能”：

  1. __global__: 这是我们最常打交道的内核函数。

    • 职能: "launched by CPU, cannot be called from GPU, must return void" (由CPU启动，不能从GPU内部调用，必须返回void)。
    • __global__ 函数是CPU和GPU沟通的桥梁的GPU端。CPU通过 <<<...>>> 语法启动它。它就像一个“入口点”，是并行任务的开始。它不能有返回值（结果通常通过指针参数写回GPU内存），也不能被其他 __global__ 或 __device__ 函数直接调用（在较新的计算能力中，通过动态并行特性可以从GPU启动其他__global__内核，但这是高级主题）。

  2. __device__: GPU上的工具函数。

    • 职能: "called from other GPU functions, cannot be called by the CPU" (从其他GPU函数调用，不能被CPU调用)。
    • __device__ 函数是纯粹在GPU设备上执行的函数。它可以被 __global__ 函数调用，也可以被其他的 __device__ 函数调用。CPU无法直接启动或调用它。你可以把它看作是内核函数中复用的代码块。

  3. __host__: 我们熟悉的CPU函数。

    • 职能: "can be called by CPU" (可以被CPU调用)。
    • 这是标准的C/C++函数，在CPU上执行。如果一个函数没有特别的CUDA限定符，它默认就是 __host__ 函数。

  4. __host__ and __device__ qualifiers can be combined (可以组合使用)

    • 职能: "sample use: overloading operators" (示例用途：重载操作符)。
    • 你可以为一个函数同时指定 __host__ 和 __device__ 限定符。这意味着NVCC编译器会为这个函数生成两个版本的代码：一个在CPU上运行，一个在GPU上运行。这在编写既希望能在主机代码中调用，又希望能在设备代码中调用的工具函数或操作符重载时非常有用，可以避免代码重复。
// GPU上的工具函数 (设备代码)
__device__ float add_on_gpu(float a, float b) {
    return a + b;
}

// 内核函数 (由CPU启动，在GPU执行)
__global__ void my_cuda_kernel(float* input_a, float* input_b, float* output, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        // 调用 __device__ 函数
        output[idx] = add_on_gpu(input_a[idx], input_b[idx]);
    }
}

// CPU上的函数 (主机代码)
__host__ void prepare_data_on_cpu(float* arr, int N) {
    for (int i = 0; i < N; ++i) {
        arr[i] = static_cast<float>(i);
    }
}

// 既能在CPU上运行，也能在GPU上运行的函数
__host__ __device__ float utility_square(float x) {
    return x * x;
}

__global__ void another_kernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        data[idx] = utility_square(data[idx]); // 在GPU上调用
    }
}

int main() {
    // ... (省略内存分配和数据拷贝) ...
    float h_data[10];
    prepare_data_on_cpu(h_data, 10); // 在CPU上调用 __host__ 函数

    for(int i=0; i<5; ++i) {
        h_data[i] = utility_square(h_data[i]); // 在CPU上调用 __host__ __device__ 函数
    }

    // my_cuda_kernel<<<grid, block>>>(d_a, d_b, d_c, N); // 启动内核
    // another_kernel<<<grid, block>>>(d_data, N);
    // ...
    return 0;
}

今天的内容就先讲到这里吧！关注杜子源源，AI、HPC新知，源源不断，带你冲向技术最前锋！

avatar
文章
阅读
粉丝