2855 字
14 分钟
CUDA 编程模型
2026-07-17
CUDA-C
无标签

GPU#

GPU 架构的特有功能#

  • 一种通过层次结构在GPU中组织线程的方法
  • 一种通过层次结构在GPU中访问内存的方法

从以下角度看并行计算#

  • 领域层（算法设计）
  • 逻辑层
  • 硬件层

在 C语言的并行编程中，要使用 pthreads 或 OpenMP 来显式地管理线程。

内存管理#

cudaMalloc 只在 GPU 内存中分配内存。充分保持和 C 语言接口的一致性。 cudaMemcpy 负责主机与设备之间的数据传输，包括 4 种 kind：

该函数以同步方式执行，在 cudaMemcpy 返回及传输操作完成前，主机应用程序处于阻塞状态。除内核启动外的 CUDA 调用均返回 cudaError_t 类型的错误枚举值：若 GPU 内存分配成功则返回 cudaSuccess，否则返回 cudaErrorMemoryAllocation。此外，可调用 cudaGetErrorString(cudaError_t error) 将错误代码转换为可读的错误消息，该函数功能与 C 语言中的 strerror 类似。

GPU 内存层次结构#

-Xcompiler用于指定命令行选项是指向C编译器还是预处理器。

CPU 和 GPU 上的代码差异#

#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <stdio.h>
#include <math.h>


void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx < N; idx++)
    {
        C[idx] = A[idx] + B[idx];
    }
}


void initialData(float *ip, int size)
{
    time_t t;
    srand((unsigned int)time(&t));
    for (int i = 0; i < size; i++)
    {
        ip[i] = (float)(rand() & 0xFF) / 10.0f;
    }
}


__global__ void sumArraysOnGPU(float *A, float *B, float *C, const int N)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N)
    {
        C[idx] = A[idx] + B[idx];
    }
}


int main(int argc, char **argv)
{
    int nElem = 1024;
    printf("Vector size %d\n", nElem);
    size_t nBytes = nElem * sizeof(float);


    // === 1. 分配内存 ===
    float *h_A, *h_B, *h_C_gpu, *h_C_cpu;
    h_A = (float *)malloc(nBytes);
    h_B = (float *)malloc(nBytes);
    h_C_gpu = (float *)malloc(nBytes);  // 存 GPU 结果
    h_C_cpu = (float *)malloc(nBytes);  // 存 CPU 结果


    // === 2. 初始化数据（先初始化！）===
    initialData(h_A, nElem);
    initialData(h_B, nElem);


    // === 3. GPU 内存分配 ===
    float *d_A, *d_B, *d_C;
    cudaMalloc((float **)&d_A, nBytes);
    cudaMalloc((float **)&d_B, nBytes);
    cudaMalloc((float **)&d_C, nBytes);


    // === 4. 拷贝 CPU -> GPU ===
    cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);


    // === 5. GPU 计算 ===
    int blockSize = 256;
    int gridSize = (int)ceil((float)nElem / blockSize);
    printf("Launching kernel: <<<gridSize=%d, blockSize=%d>>>\n", gridSize, blockSize);


    sumArraysOnGPU<<<gridSize, blockSize>>>(d_A, d_B, d_C, nElem);
    cudaDeviceSynchronize();


    // === 6. 拷贝 GPU -> CPU ===
    cudaMemcpy(h_C_gpu, d_C, nBytes, cudaMemcpyDeviceToHost);


    // === 7. CPU 计算（对比验证）===
    sumArraysOnHost(h_A, h_B, h_C_cpu, nElem);


    // === 8. 结果对比 ===
    float maxError = 0.0f;
    for (int i = 0; i < nElem; i++)
    {
        float diff = fabs(h_C_gpu[i] - h_C_cpu[i]);
        if (diff > maxError) maxError = diff;
    }
    printf("Max error between GPU and CPU: %e\n", maxError);


    // === 9. 释放资源 ===
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    cudaDeviceReset();


    free(h_A);
    free(h_B);
    free(h_C_gpu);
    free(h_C_cpu);


    return 0;
}

统一寻址#

为了避免 CPU 和 GPU 存储空间内的不恰当引用，引入了统一寻址，使用一个指针来访问 GPU 和 CPU 的内存。

线程管理#

CUDA 采用网格-线程块-线程的三级并行执行模型：一次内核启动产生一个网格，网格由多个线程块组成，每个线程块包含一组线程；同一网格内的所有线程共享全局内存，而同一块内的线程可以通过同步和共享内存进行协作，不同块内的线程无法协作；每个线程通过 blockIdx（线程块在网格中的索引）和 threadIdx（线程在线程块中的索引）这两个坐标变量来唯一标识自己，从而实现对数据的分块并行处理。

上述的索引都是 unsigned int 类型的。

版本优点缺点
(nElem + block.x - 1) / block.x纯整数运算，无浮点转换开销需确保不会溢出
ceil((float)nElem / block.x)直观易读浮点运算稍慢

CUDA 的两层线程层次结构（网格与块）通过将内核启动的线程组织为逻辑上的多维结构，使程序员能够根据 GPU 计算资源（如寄存器、共享内存等）灵活调整维度以优化性能，同时确保同一段核函数代码可在不同硬件配置的设备上高效执行。

核函数#

编写核函数#

核函数的限制#

CUDA核函数具有多项重要限制：它只能访问设备内存而非主机内存，必须声明为void返回类型且不支持返回值，不允许使用可变数量参数（如printf风格的变参函数），不支持静态变量声明，并且核函数的执行表现为异步行为——即调用后立即返回，实际的GPU计算在后台进行，需要通过同步机制（如cudaDeviceSynchronize）来确保执行完成。这些限制源于GPU大规模并行架构的设计需求，旨在保证线程间的独立性和执行效率。

限制原因
禁止可变参数GPU 需要提前知道每个线程的寄存器用量和内存布局，变参无法在编译期确定
禁止静态变量静态变量本质是全局共享状态，而 GPU 有成千上万个线程同时执行，共享可变状态会导致竞争条件和不可预测的结果

核函数的验证#

void checkResult(float *hostRef, float *gpuRef, const int N) {
    double epsilon = 1.0E-8;
    for (int i = 0; i < N; i++) {
        if (fabs(hostRef[i] - gpuRef[i]) > epsilon) {
            printf("Results do not match at index %d: host %f, gpu %f\n", i, hostRef[i], gpuRef[i]);
            return;
        }
    }
    printf("Results match!\n");
}

处理错误#

fprintf 是 C 标准库中的格式化输出函数，用于将格式化的数据写入指定的文件流。

常见用途#

场景代码示例
写入日志文件fprintf(log_fp, "[ERROR] 文件未找到\n");
输出错误信息到 stderrfprintf(stderr, "错误: 内存分配失败\n");
生成格式化报表fprintf(fp, "%-10s %8.2f\n", name, price);

与相关函数的区别#

函数输出目标示例
printf标准输出 (stdout)printf("hello");
fprintf指定文件流fprintf(fp, "hello");
sprintf字符串缓冲区sprintf(buf, "hello");
fscanf从文件流读取fscanf(fp, "%d", &n);
#define CHECK(call) \
{ \
    const cudaError_t error = call; \
    if (error != cudaSuccess) \
    { \
        fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__); \
        fprintf(stderr, "code: %d, reason: %s\n", error, cudaGetErrorString(error)); \
        exit(1); \
    } \
}

exit(1) 表示立即结束程序 call 是 CUDA API 调用参数，传入一个函数声明

该宏在检测到内存拷贝或异步操作错误时，会输出错误代码及可读信息并终止程序；而 CHECK(cudaDeviceSynchronize()) 则通过阻塞主机线程直到设备完成所有任务，来检查核函数启动后的错误。但需注意，后者仅适用于调试场景，因为它会在核函数启动后形成全局同步屏障，严重影响程序性能，不应在生产环境中使用。

全局同步屏障#

在 CUDA 编程中，全局同步屏障（global synchronization barrier）是指一个强制所有执行单元（线程、流、处理器核心等）在继续执行之前都必须到达并等待某一点的同步机制。

具体到 cudaDeviceSynchronize() 的场景：

它之所以被称为”全局屏障”，是因为：

  1. 阻塞整个设备：cudaDeviceSynchronize() 会暂停主机端（CPU）线程，直到设备端（GPU）上所有流（stream）中的所有已提交任务全部执行完毕。这意味着无论你的核函数是在哪个流中启动的，都会被强制同步等待。
  2. 破坏流水线并行：CUDA 的核心优势之一是允许主机端在提交一个核函数后继续向设备提交其他任务，实现 CPU 与 GPU 的流水线并行。插入 cudaDeviceSynchronize() 后，主机线程必须干等 GPU 全部做完，CPU 在此期间无法做其他有用工作（比如准备下一批数据、启动其他流的核函数等），整个流水线被迫中断。
  3. 粒度粗：它不是针对某个特定流或某个特定核函数的同步，而是”一刀切”地同步整个设备上的所有活动。

对比更细粒度的同步方式：

  • cudaStreamSynchronize(stream) —— 只同步指定流中的任务，其他流不受影响
  • cudaEventSynchronize(event) —— 只等待到某个特定事件点，更灵活
  • 核函数内部的 __syncthreads() —— 只同步同一个线程块（block）内的线程

总结一句话：
cudaDeviceSynchronize() 被称为全局同步屏障，是因为它强制整个 GPU 设备的所有工作全部完成后才放行主机线程，粗粒度的全局等待会彻底打断 CPU-GPU 流水线并行，因此只应在调试错误定位时使用，正式代码中应避免。

组织并行线程#

checkResult#

// checkresult函数用于验证GPU计算结果是否正确
void checkResult(float *hostRef, float *gpuRef, int nxy) {
    double epsilon = 1.0E-8;
    bool match = 1;
    for (int i = 0; i < nxy; i++) {
        if (fabs(hostRef[i] - gpuRef[i]) > epsilon) {
            match = 0;
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i], gpuRef[i], i);
            break;
        }
    }
    if (match) printf("Arrays match.\n\n");
}

TimeOnCPU#

double cpuSecond() {
    struct timeval tp;
    gettimeofday(&tp, NULL);
    return ((double)tp.tv_sec + (double)tp.tv_usec * 1.e-6);
}


int main() {
  ...
  double iStart = cpuSecond();
  sumArraysOnGPU<<<gridSize, blockSize>>>(d_A, d_B, d_C);
  double iElaps = cpuSecond() - iStart;
  ...
}

matrixSum#

// 二维矩阵求和核函数
__global__ void sumMatrix2D(float *A, float *B, float *C, int nx, int ny) {
    int ix = blockIdx.x * blockDim.x + threadIdx.x;
    int iy = blockIdx.y * blockDim.y + threadIdx.y;
    int idx = iy * nx + ix;


    if (ix < nx && iy < ny) {
        C[idx] = A[idx] + B[idx];
    }
}

设备管理#

CUDA运行时API函数#

#include <cuda_runtime.h>
#include <stdio.h>


int main() {
    // 先检查系统有多少 CUDA 设备
    int deviceCount = 0;
    cudaError_t err = cudaGetDeviceCount(&deviceCount);


    if (err != cudaSuccess) {
        printf("Failed to get device count: %s\n", cudaGetErrorString(err));
        return -1;
    }


    if (deviceCount == 0) {
        printf("No CUDA-capable device found!\n");
        return -1;
    }


    printf("Found %d CUDA device(s)\n", deviceCount);


    for (int device = 0; device < deviceCount; device++) {
        cudaDeviceProp prop;
        err = cudaGetDeviceProperties(&prop, device);


        if (err != cudaSuccess) {
            printf("Error getting properties for device %d: %s\n",
                   device, cudaGetErrorString(err));
            continue;
        }


        printf("\n=== Device %d: %s ===\n", device, prop.name);
        printf("  Compute capability: %d.%d\n", prop.major, prop.minor);
        printf("  Total global memory: %.2f GB\n", prop.totalGlobalMem / (1024.0*1024.0*1024.0));
        printf("  Shared memory per block: %zu bytes\n", prop.sharedMemPerBlock);
        printf("  Registers per block: %d\n", prop.regsPerBlock);
        printf("  Warp size: %d\n", prop.warpSize);
        printf("  Max threads per block: %d\n", prop.maxThreadsPerBlock);
        printf("  Max threads dimensions: (%d, %d, %d)\n",
               prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
        printf("  Max grid size: (%d, %d, %d)\n",
               prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
    }


    return 0;
}

nvidia-smi#

参数显示内容
MEMORYGPU 显存使用情况（总量、已用、空闲、预留等）
UTILIZATIONGPU 利用率（计算、显存、编码器、解码器等）
ECCECC（错误校验码）状态，用于检测和纠正显存错误
TEMPERATUREGPU 温度信息（当前温度、阈值等）
POWER功耗信息（当前功率、功率限制、默认限制等）
CLOCK时钟频率（SM、显存、视频等时钟当前/最大频率）
COMPUTE计算模式（Default/Exclusive Thread/Exclusive Process/Prohibited）
PIDS当前占用 GPU 的进程 ID 列表
PERFORMANCE性能状态（P-State，如 P0 最高性能、P12 最低功耗）
SUPPORTED_CLOCKS支持的时钟频率组合列表
PAGE_RETIREMENT显存页退休状态（记录因错误被禁用的显存页）
ACCOUNTING进程资源统计（计算每个进程使用的 GPU 时间、显存等）

确定最优 GPU#

在运行时设置设备#

CUDA_VISIBLE_DEVICES 是运行时指定 GPU 的环境变量，无需修改应用程序代码即可控制哪些 GPU 对 CUDA 程序可见。驱动程序会重新映射设备 ID：例如设置 CUDA_VISIBLE_DEVICES=2 后，物理设备 2 在程序中显示为设备 0；设置 CUDA_VISIBLE_DEVICES=2,3 后，物理设备 2 和 3 分别映射为程序中的设备 0 和 1，其余 GPU 被屏蔽。

CUDA 程序与 C 语言并行编程的核心区别在于其独有的两级线程层次结构（网格 → 线程块 → 线程），通过这一抽象可以控制大规模并行环境。网格和线程块的尺寸对内核性能有显著影响。针对同一问题，核函数的实现方式和执行配置有多种选择，通常需要在实现复杂度和性能之间权衡。

CUDA 编程模型
https://www.dongyanzhang.com/posts/cuda-c/cuda-c-002/
作者
阿东阿言
发布于
2026-07-17
许可协议
CC BY-NC-SA 4.0
并行计算与 CUDA 基础
CUDA 执行模型
© 2026 阿东阿言. All Rights Reserved. / RSS / Sitemap
Powered by Astro & Fuwari