文章 标签 分类

目录

GPU Kernel 是什么？

 收录于 ai
   约 2450 字   预计阅读 5 分钟 

GPU Kernel

1. GPU Kernel 是什么？

核心概念： GPU Kernel（GPU 内核）是一个在 GPU 设备上并行执行的函数。

可以把它想象成一个你为 GPU 编写的“小程序”或“函数”。与普通在 CPU 上运行的函数不同，当你“调用”或“启动”（Launch）一个 GPU Kernel 时，它会被成千上万个 GPU 线程（Threads）同时执行。每个线程都执行完全相同的代码，但通常处理不同的数据。

一个生动的比喻：

  • CPU (中央处理器)：像一个拥有几位（比如 8 或 16 位）能力超强的专家（核心）。他们可以处理非常复杂的任务，但人数少，不适合处理成千上万个简单重复的任务。
  • GPU (图形处理器)：像一个拥有成千上万名（比如 10240 名）熟练工人（核心）的工厂。每个工人能力不如专家强，但他们可以同时做同样一件简单、重复的工作（例如，给成千上万个零件拧螺丝）。

GPU Kernel 就是你给这成千上万名工人下达的“工作指令”。 这个指令告诉每个工人具体要做什么。

2. 关键特性与术语 (以 NVIDIA CUDA 为例)

要理解 Kernel，需要了解它的执行模型：

  1. Host (主机) vs. Device (设备)：

    • Host: 指的是 CPU 及其内存（RAM）。你的主程序（例如 main 函数）运行在 Host 上。
    • Device: 指的是 GPU 及其显存（VRAM）。Kernel 函数运行在 Device 上。

  2. Kernel 函数的声明：

    • 在 CUDA C++ 中，Kernel 函数使用 __global__ 关键字来声明。这告诉编译器，这个函数将从 Host 调用，并在 Device 上执行。

  3. 线程层次结构 (Thread Hierarchy)：

    • Thread (线程)：执行 Kernel 的最小单位。每个线程都有一个唯一的 ID。
    • Block (线程块)：一组线程的集合。同一个 Block 内的线程可以相互协作，例如通过共享内存（Shared Memory）快速交换数据。
    • Grid (网格)：一组 Block 的集合。当你启动一个 Kernel 时，你实际上是在启动一个 Grid 的线程。

    /posts/gpu_kernel%E6%98%AF%E4%BB%80%E4%B9%88/grid-of-thread-blocks.png

    这个层次结构非常重要，因为它让你可以将问题分解并映射到 GPU 的硬件上。

  4. Kernel 启动：

    • 在 Host 端，使用特殊的语法 <<<...>>> 来启动 Kernel。
    • 语法：kernel_name<<<Grid_Size, Block_Size>>>(argument1, argument2, ...);
      • Grid_Size：指定启动多少个 Block。
      • Block_Size：指定每个 Block 中包含多少个线程。
      • 总线程数 = Grid_Size * Block_Size。

3. 详细例子：向量加法 (Vector Addition)

这是一个 GPU 编程的 “Hello, World!"。我们要计算两个大向量 A 和 B 的和，结果存入向量 C。即 C[i] = A[i] + B[i]。

a. CPU 上的串行实现 (用于对比)

void vectorAdd_cpu(float* a, float* b, float* c, int n) {
    for (int i = 0; i < n; i++) {
        c[i] = a[i] + b[i];
    }
}

这个实现很简单，就是一个 for 循环，一次处理一个元素，非常慢。

b. GPU 上的并行实现 (使用 CUDA)

第1步：编写 GPU Kernel

这个 Kernel 函数就是每个 GPU 线程要执行的任务。我们的策略是让每个线程负责计算结果向量 C 中的一个元素。

// __global__ 关键字表示这是一个将在 GPU 上运行的 Kernel 函数
__global__ void vectorAdd_gpu(float* a, float* b, float* c, int n) {
    // 计算当前线程的全局唯一索引 (ID)
    // blockIdx.x: 当前是第几个 Block
    // blockDim.x: 每个 Block 有多少个线程
    // threadIdx.x: 当前线程在 Block 内的索引
    int index = blockIdx.x * blockDim.x + threadIdx.x;

    // 边界检查：防止线程访问越界内存
    // (因为总线程数可能大于向量的实际大小)
    if (index < n) {
        c[index] = a[index] + b[index];
    }
}

代码解读：

  • int index = ...：这是 GPU 编程中最核心的一行代码。它让每个线程都能计算出自己应该处理哪个数据。例如，第0个线程计算 C[0] = A[0] + B[0]，第100个线程计算 C[100] = A[100] + B[100]。
  • if (index < n)：这是一个安全检查。如果我们有 1000 个元素，但为了凑整启动了 1024 个线程，那么多出来的 24 个线程就不应该执行任何操作。

第2步：编写 Host 端主程序来调用 Kernel

这个主程序负责：

  1. 准备数据 (在 CPU 内存中)。
  2. 在 GPU 显存中分配空间。
  3. 将数据从 CPU 拷贝到 GPU。
  4. 启动 Kernel 执行计算。
  5. 将结果从 GPU 拷贝回 CPU。
  6. 清理内存。
#include <iostream>
#include <cuda_runtime.h>

// (将上面的 Kernel 函数 vectorAdd_gpu 放在这里)
__global__ void vectorAdd_gpu(float* a, float* b, float* c, int n) { ... }

int main() {
    int N = 1000000; // 向量大小，一百万个元素
    size_t size = N * sizeof(float);

    // 1. 在 Host (CPU) 上分配内存并初始化数据
    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] = 1.0f;
        h_b[i] = 2.0f;
    }

    // 2. 在 Device (GPU) 上分配内存
    float* d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    // 3. 将数据从 Host 拷贝到 Device
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    // 4. 定义 Kernel 的执行配置并启动 Kernel
    int threadsPerBlock = 256;
    // 计算需要的 Block 数量，(N + threadsPerBlock - 1) / threadsPerBlock 是一个向上取整的常用技巧
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock; 
    
    std::cout << "Launching Kernel..." << std::endl;
    std::cout << "Grid size: " << blocksPerGrid << ", Block size: " << threadsPerBlock << std::endl;
    
    // <<<...>>> 语法启动 Kernel！
    vectorAdd_gpu<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N);

    // 5. 将计算结果从 Device 拷贝回 Host
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 6. 验证结果 (简单抽查第一个和最后一个)
    if (h_c[0] == 3.0f && h_c[N-1] == 3.0f) {
        std::cout << "Calculation successful!" << std::endl;
    } else {
        std::cout << "Calculation failed!" << std::endl;
    }

    // 7. 释放 Host 和 Device 上的内存
    free(h_a);
    free(h_b);
    free(h_c);
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    return 0;
}

总结这个例子： vectorAdd_gpu 函数本身就是 GPU Kernel。它定义了单个线程的行为。而 main 函数通过 <<<...>>> 语法，一次性启动了 blocksPerGrid * threadsPerBlock 个线程，让它们在 GPU 上并行地执行这个 Kernel，从而在极短的时间内完成了百万次加法运算。

4. 相关的 GitHub 项目

这里有一些从入门到实践的好项目，可以帮助你理解 GPU Kernel 的应用。

  1. 入门学习 & 官方示例

    • NVIDIA/cuda-samples
      • 简介：NVIDIA 官方的 CUDA 示例代码库。这是学习 CUDA 和 Kernel 编程的 最佳起点。它包含了从基础（如向量加法）到高级（如并行归约、矩阵乘法、FFT）的各种例子。
      • 看点：0_Introduction/vectorAdd 就是我们上面例子的完整版。

  2. 经典书籍配套代码

    • cuda-by-example
      • 简介：经典入门书籍《CUDA by Example》（中文版《GPU高性能编程CUDA实战》）的配套代码。代码组织清晰，与书本章节对应，非常适合初学者。
      • 看点：包含了从基本内存操作到使用共享内存优化 Kernel 性能的完整过程。

  3. 深度学习领域的 Kernel 应用

    • pytorch/pytorch 或 tensorflow/tensorflow
      • 简介：主流的深度学习框架。它们的底层大量使用了 CUDA Kernel 来加速神经网络的各种运算（如卷积、矩阵乘法、激活函数等）。
      • 看点：代码非常复杂，但可以让你看到工业级的 Kernel 是如何编写的。例如，可以在 PyTorch 的 aten/src/ATen/native/cuda 目录下找到大量 .cu 文件，里面就是各种运算的 Kernel 实现。
    • karpathy/llm.c
      • 简介：Andrej Karpathy 用纯 C/CUDA 从零开始实现一个 GPT-2 模型的训练和推理。这个项目非常火，因为它代码量小，可读性强。
      • 看点：llm.c 的 CUDA 版本（例如 train_gpt2.cu）中包含了实现 Transformer 模型关键组件（如 LayerNorm, MatMul, Softmax）的自定义 Kernel，是学习在实际复杂应用中编写 Kernel 的绝佳案例。

  4. 计算机视觉和图形学

    • opencv/opencv
      • 简介：著名的计算机视觉库。OpenCV 的 cuda 模块提供了许多函数的 GPU 加速版本。
      • 看点：可以查看 modules/cudaarithm/src/cuda 等目录，看到图像处理（如滤波、颜色空间转换）的 Kernel 是如何实现的。
更新于 2025-06-30
 GPU, ai
返回 | 主页