CUDA编程入门：从基础概念到实战示例

CUDA编程入门：从基础概念到实战示例

引用

CSDN

1.

https://blog.csdn.net/weixin_48835186/article/details/144987515

CUDA（Compute Unified Device Architecture）是由NVIDIA开发的一种并行计算平台和编程模型，旨在利用NVIDIA GPU的强大计算能力进行通用计算（GPGPU）。CUDA使开发者能够通过C、C++、Fortran等语言编写程序，将计算任务分配给GPU来执行，从而实现计算加速。本文将详细介绍CUDA的基本概念、编程模型以及具体的编程实例。

CUDA的含义及作用

CUDA（Compute Unified Device Architecture）是由NVIDIA开发的一种并行计算平台和编程模型，旨在利用NVIDIA GPU（图形处理单元）的强大计算能力进行通用计算（GPGPU，General-Purpose computing on Graphics Processing Units）。CUDA使开发者能够通过 C、C++、Fortran 等语言编写程序，将计算任务分配给 GPU 来执行，而不是仅限于 CPU。

CUDA 编程包括以下几个方面：

• 内核函数（Kernel）：这是程序中可以在 GPU 上并行执行的核心函数。开发者需要编写内核函数，并将它们发送到 GPU 上执行。
• 线程（Thread）和线程块（Block）：CUDA 编程模型以线程为基本单位，多个线程按块（block）和网格（grid）的结构组织，可以同时在 GPU 上并行运行。
• 内存管理：GPU 拥有多种类型的内存（如全局内存、共享内存、常量内存等），开发者需要优化内存访问策略，以提高程序性能。

使用CUDA编程有以下好处：

• 并行计算加速：GPU 拥有大量的计算核心（CUDA 核心），可以在同一时间内执行成千上万个计算任务。这使得它在处理大规模并行计算时远远超过 CPU 的性能。例如，科学计算、图像处理、视频编码等任务，尤其是在机器学习和深度学习中的训练和推理任务中，CUDA 提供了极大的加速。
• 性能提升：对于计算密集型应用，CUDA 能够充分利用 GPU 的高并行度和高速内存，极大地提升计算效率。例如，深度学习框架（如 TensorFlow、PyTorch）利用 CUDA 实现了模型训练的加速。
• 广泛应用领域：CUDA 被广泛应用于科学计算、机器学习、人工智能、图像和视频处理、金融建模、物理仿真、基因组学等领域，几乎涉及所有需要高计算性能的领域。
• 开发工具和库：NVIDIA 提供了丰富的工具集和库来简化 CUDA 编程，例如：
• cuBLAS、cuFFT：用于线性代数和快速傅里叶变换的高效库。
• cuDNN：用于深度学习的 GPU 加速库。
• Thrust：CUDA 提供的一个并行算法库，简化了并行编程。
• 兼容性：CUDA 编程不需要特别的硬件要求，现代的 NVIDIA GPU（如 Tesla、Quadro、GeForce 等）都可以使用 CUDA 进行加速计算。同时，CUDA 可以与多种编程语言和开发环境兼容，提供了灵活的开发体验。

CUDA 编程模型

CUDA 编程模型有几个关键组成部分：

• 线程（Thread）：每个线程执行相同的代码，但是可能操作不同的数据。
• 线程块（Block）：线程块是线程的集合。线程块中的线程在同一块内共享内存，并可以同步执行。
• 网格（Grid）：网格是由多个线程块组成的。网格是一个执行的整体单元，包含了所有的线程块。
• 内存层次结构：CUDA 的内存结构分为不同的层次，包括全局内存、共享内存、常量内存等，分别适用于不同的数据访问模式。

展示一个简单的C＋＋进行CUDA编程的代码：

#include <iostream>
#include <cuda_runtime.h>

// CUDA内核函数：计算两个向量的加法
__global__ void vectorAdd(const float* A, const float* B, float* C, int N) {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if (index < N) {
        C[index] = A[index] + B[index];
    }
}

int main() {
    int N = 100000;  // 向量大小
    size_t size = N * sizeof(float);

    // 分配主机内存
    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] = i * 1.0f;
        h_B[i] = i * 2.0f;
    }

    // 分配设备内存
    float* d_A, * d_B, * d_C;
    cudaMalloc((void**)&d_A, size);
    cudaMalloc((void**)&d_B, size);
    cudaMalloc((void**)&d_C, size);

    // 将数据从主机复制到设备
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 设置CUDA内核的执行配置
    int blockSize = 256;  // 每个线程块中的线程数
    int numBlocks = (N + blockSize - 1) / blockSize;  // 计算块的数量

    // 启动内核
    vectorAdd << <numBlocks, blockSize >> > (d_A, d_B, d_C, N);

    // 检查内核执行是否有错误
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        std::cerr << "CUDA error: " << cudaGetErrorString(err) << std::endl;
        return -1;
    }

    // 将结果从设备复制回主机
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 验证结果
    for (int i = 0; i < 10; i++) {  // 输出前10个元素
        std::cout << "C[" << i << "] = " << h_C[i] << std::endl;
    }

    // 释放设备内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // 释放主机内存
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

在一般的C++程序中，声明CUDA函数：

__global__ void function1() {
    // ...
}

表示建立一个CUDA内核函数，运行此函数使用GPU。

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

这个公式的目的是计算当前线程在整个网格中的全局索引（index）。要理解这段代码，需要了解以下几个变量的含义：

1. threadIdx.x：表示当前线程在其所属的线程块中的索引。每个线程块内的线程都有一个唯一的threadIdx，而threadIdx.x表示当前线程在块中的横向索引（即X维度）。
2. blockIdx.x：表示当前线程块在整个网格中的索引。blockIdx.x是一个标量，表示线程块在网格中所在的横向位置（即X维度）。CUDA支持三维的块和网格，所以这里是x维。
3. blockDim.x：表示每个线程块中线程的总数。blockDim.x是一个标量，表示每个线程块中有多少个线程。通常在一维线程块的情况下，blockDim.x表示线程块的横向大小。

// 将数据从主机复制到设备
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);

• d_A：目标内存地址，表示设备（GPU）上的内存位置。这个变量通常是一个指向 GPU 内存的指针，类型通常为float*、int*等。
• h_A：源内存地址，表示主机（CPU）上的内存位置。这个变量通常是一个指向 CPU 内存的指针，类型通常为float*、int*等。
• size：要复制的数据大小（以字节为单位）。例如，如果你有一个数组float h_A[100];，并且每个元素是一个float（4字节），则size应为100 * sizeof(float)，即 400 字节。
• cudaMemcpyHostToDevice：指定复制方向，表示数据从主机内存复制到设备内存。

剩下的步骤都很容易理解！可以联系下矩阵相加的CUDA编程，相对于向量只是变成二维的了。这里给出代码：

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

// CUDA 核函数，用于矩阵加法
__global__ void MatAddKernel(float* A, float* B, float* C, int height, int width) {
    // 获取线程的全局ID
    int i = blockIdx.x * blockDim.x + threadIdx.x; // 计算全局行索引
    int j = blockIdx.y * blockDim.y + threadIdx.y; // 计算全局列索引

    // 确保索引在矩阵范围内
    if (i < width && j < height) {
        // 计算当前线程对应的元素索引
        int index = j * width + i;

        // 从矩阵 A 和 B 中读取数据
        float src_data_A = A[index];
        float src_data_B = B[index];

        // 执行加法运算
        float result = src_data_A + src_data_B;

        // 将结果写入矩阵 C
        C[index] = result;
    }
}

void MatAdd(int height, int width) {
    // 在主机内存中分配 A、B 和 C
    float* A = (float*)malloc(height * width * sizeof(float));
    float* B = (float*)malloc(height * width * sizeof(float));
    float* C = (float*)malloc(height * width * sizeof(float));

    // 初始化输入矩阵 A 和 B
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            A[i * width + j] = i + j; // 简单初始化，A的元素为行索引+列索引
            B[i * width + j] = i - j; // 简单初始化，B的元素为行索引-列索引
        }
    }

    // 第一步：在设备内存中为矩阵 A、B 和 C 分配内存
    float* d_A;
    cudaMalloc(&d_A, height * width * sizeof(float)); // 分配矩阵 A 的设备内存
    float* d_B;
    cudaMalloc(&d_B, height * width * sizeof(float)); // 分配矩阵 B 的设备内存
    float* d_C;
    cudaMalloc(&d_C, height * width * sizeof(float)); // 分配矩阵 C 的设备内存

    // 第二步：将矩阵 A 和 B 从主机内存复制到设备内存
    cudaMemcpy(d_A, A, height * width * sizeof(float), cudaMemcpyHostToDevice); // 复制 A
    cudaMemcpy(d_B, B, height * width * sizeof(float), cudaMemcpyHostToDevice); // 复制 B

    // 第三步：调用 CUDA 核函数
    dim3 threadsPerBlock(16, 16); // 定义每个块中的线程数
    dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x,
        (height + threadsPerBlock.y - 1) / threadsPerBlock.y); // 计算网格中的块数
    MatAddKernel << <numBlocks, threadsPerBlock >> > (d_A, d_B, d_C, height, width); // 启动 CUDA 核函数

    // 第四步：将结果从设备内存复制回主机内存
    cudaMemcpy(C, d_C, height * width * sizeof(float), cudaMemcpyDeviceToHost);

    // 输出结果矩阵 C
    printf("Matrix C (result of A + B):\n");
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            printf("%f ", C[i * width + j]);
        }
        printf("\n");
    }

    // 释放设备内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // 释放主机内存
    free(A);
    free(B);
    free(C);
}

int main() {
    int height = 3; // 矩阵的高度
    int width = 3;  // 矩阵的宽度

    // 调用矩阵加法函数
    MatAdd(height, width);

    return 0;
}