CUDA从入门到精通（三）—

CUDA 编程简介

CUDA（Compute Unified Device Architecture）是由 NVIDIA 提供的一种并行计算平台和编程模型。它允许开发者利用 NVIDIA GPU 的并行计算能力，编写可以在 GPU 上高效运行的代码，从而加速计算密集型任务。

CUDA 通过扩展标准的 C/C++ 语言，提供了编译工具链和库，使程序员可以轻松编写并行代码。

1. CUDA 的核心概念

GPU 和 CPU 的协同计算

CPU：负责运行主程序（主机代码，Host Code），将任务分配给 GPU。
GPU：负责运行数据并行的任务（设备代码，Device Code），执行大量并行计算。

CUDA 核函数 (Kernel)

Kernel 是 CUDA 程序中的核心函数，它运行在 GPU 上的所有线程中。

通过 __global__ 关键字定义：

__global__ void myKernel(int *a) {
    int i = threadIdx.x; // 当前线程索引
    a[i] = a[i] * 2;     // 并行计算
}

线程、线程块和网格 (Grid)

CUDA 提供了一种分层的 线程模型，用于在 GPU 上组织并行计算：

线程 (Thread)：执行计算的基本单元。
线程块 (Thread Block)：由一组线程组成，可以共享数据。
网格 (Grid)：由多个线程块组成。

线程的组织结构可以用一维、二维或三维方式定义：

dim3 blockSize(256);       // 每个线程块 256 个线程
dim3 gridSize((N + 255)/256); // 网格大小 (任务总数 / 线程数)
myKernel<<<gridSize, blockSize>>>(d_data);

内存模型

CUDA 提供多种内存空间，分为 主机内存（CPU） 和 设备内存（GPU）：

全局内存（Global Memory）：GPU 所有线程都可以访问，访问速度慢但容量大。
共享内存（Shared Memory）：线程块内的线程共享的数据，速度快但容量小。
寄存器（Registers）：线程私有的内存，速度最快。
常量内存（Constant Memory）：只读内存，所有线程共享。
纹理内存 / 表面内存：用于特殊优化的数据访问场景。

2. CUDA 程序结构

一个典型的 CUDA 程序分为以下几个步骤：

主机代码初始化（CPU）：
- 分配和初始化内存。
从CPU内存中拷贝数据到GPU内存
- 将数据传输到 GPU（cudaMemcpy）。
内核函数执行（GPU）：
- 调用核函数（<<<gridSize, blockSize>>>），GPU 执行并行任务。
结果传回主机：
- 将 GPU 上的数据传回 CPU。
释放资源：
- 释放主机和设备内存。

3. 简单 CUDA 示例

3.1 helloworld示例

#include<stdio.h>
__global__ void helloWorldFromGPU(){
    printf("hello world from GPU!\n\n");
}

int main(int argc, char const *argv[])
{
    helloWorldFromGPU<<<1,10>>>();
    return 0;
}

执行：

nvcc helloworld.cu -o main
main

在这里插入图片描述

cudaDeviceReset函数

cudaDeviceReset 是一个用于清理 CUDA 设备状态的函数，它会释放与当前 CUDA 上下文（Context）相关联的所有资源，并将设备状态重置为初始状态。如果程序不调用它，CUDA 运行时可能会保留一些状态或资源，直到程序完全结束。

调用 cudaDeviceReset 的情况：

释放 GPU 上的资源：cudaDeviceReset 会确保设备上的内存和上下文被完全释放，设备状态恢复为初始状态。
干净退出：CUDA 上下文释放后，程序会正常退出，不会留有残余资源。

移除 cudaDeviceReset 的情况：

资源未完全释放：在程序结束时，GPU 上的资源会保留，直到操作系统自动回收。
影响多次执行的程序：如果后续还有其他 CUDA 程序，设备资源状态可能会受到前一个程序的影响。
调试环境行为：在调试工具（如 cuda-memcheck）中，未调用 cudaDeviceReset 可能会产生警告，指出设备资源未被显式释放。

对于小型程序，两种情况下的运行结果几乎没有差异，程序可以正常执行。但在以下情况下，差异会更加明显：

长时间运行的程序：不调用 cudaDeviceReset 可能导致设备资源持续占用。
反复调用 CUDA 程序：如果 GPU 设备状态未重置，后续程序可能会遇到资源分配失败或初始化冲突。
工具检测：使用 CUDA 调试工具（如 cuda-memcheck）时，不调用 cudaDeviceReset 可能会报告内存泄漏。

#pragma execution_character_set("utf-8")
#include <cuda_runtime.h>
#include <iostream>

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    const int N = 1000000; // 数据规模
    size_t size = N * sizeof(int);

    // 分配主机内存
    int *h_a = (int *)malloc(size);
    int *h_b = (int *)malloc(size);
    int *h_c = (int *)malloc(size);

    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_a[i] = i;
        h_b[i] = 2 * i;
    }

    // 分配设备内存
    int *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    // 将数据从主机传输到设备
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    // 启动核函数
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, N);

    // 将结果传回主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 验证结果
    for (int i = 0; i < 10; i++) {
        std::cout << h_c[i] << " ";
    }
    std::cout << std::endl;

    // 释放资源
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    // cudaDeviceReset 调用 (可以移除)
    cudaDeviceReset();

    return 0;
}

在这里插入图片描述

cudaDeviceSynchronize函数

`cudaDeviceSynchronize` 与 `cudaDeviceReset` 的区别

cudaDeviceSynchronize：
- 作用：确保所有之前提交到 GPU 的任务（如核函数、数据传输）都执行完成，起到同步的作用。
- 执行后：程序继续执行后续的代码，但不会清理 GPU 上的资源。
cudaDeviceReset：
- 作用：除了同步任务外，还会将 GPU 设备状态重置并释放所有资源（如设备内存、上下文等）。
- 执行后：GPU 状态重置为初始状态，所有分配的资源会被释放。

为什么不能直接用 `cudaDeviceSynchronize` 替换 `cudaDeviceReset`？

cudaDeviceSynchronize 仅起到同步作用，它不会释放 GPU 资源。
如果程序在多次运行 CUDA 任务时，没有调用 cudaDeviceReset，GPU 上的资源可能无法及时释放，导致 资源泄漏 或 初始化冲突。

替换示例

以下是代码中用 cudaDeviceSynchronize 替换 cudaDeviceReset 的示例：

#pragma execution_character_set("utf-8")
#include <cuda_runtime.h>
#include <iostream>

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    const int N = 1000000; // 数据规模
    size_t size = N * sizeof(int);

    // 分配主机内存
    int *h_a = (int *)malloc(size);
    int *h_b = (int *)malloc(size);
    int *h_c = (int *)malloc(size);

    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_a[i] = i;
        h_b[i] = 2 * i;
    }

    // 分配设备内存
    int *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    // 将数据从主机传输到设备
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    // 启动核函数
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, N);

    // 将结果传回主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 验证结果
    for (int i = 0; i < 10; i++) {
        std::cout << h_c[i] << " ";
    }
    std::cout << std::endl;

    // 释放资源
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    // cudaDeviceReset 调用 (可以移除)
    // cudaDeviceReset();
    cudaDeviceSynchronize();

    return 0;
}

在这里插入图片描述

打印线程索引

#include<stdio.h>
__global__ void helloWorldFromGPU(){
    //  需要添加线程索引
    int threadIndex = threadIdx.x;
    printf("hello world from GPU! threadIndex: %d\n\n", threadIndex);
}

int main(int argc, char const *argv[])
{
    helloWorldFromGPU<<<1,10>>>();  // 在GPU上执行helloWorldFromGPU函数
    return 0;
}

在这里插入图片描述

调用 cudaDeviceReset：
- GPU 资源会被完全释放。
- 在多次运行程序时，GPU 设备状态重新初始化。
调用 cudaDeviceSynchronize：
- 程序会等待所有任务完成，但不会释放 GPU 资源。
- 如果不手动释放内存（如 cudaFree），资源会在程序结束时才被操作系统回收。