CUDA架构介绍与设计模式解析

看不见的罗辑

已于 2024-04-29 20:44:04 修改

阅读量4.6k

点赞数 12

CC 4.0 BY-SA版权

文章标签： cuda 设计模式 cuda系统架构 cuda中的设计模式

于 2024-04-29 20:39:53 首次发布

本文链接：https://blog.csdn.net/a284365/article/details/138324736

CUDA架构介绍与设计模式解析

1 CUDA 介绍

CUDA发展历程

CUDA 是 NVIDIA 于 2006 年推出的通用并行计算架构，经过多年发展，已成为 GPU 计算的重要标准。

CUDA主要特性

CUDA 提供强大的并行计算能力，具备以下特性：

线程层次结构
扩展了 C/C++ 语言，引入了 GPU 扩展如 kernel 函数
显式的内存层次结构，允许根据计算需求访问内存数据
跨平台支持

CUDA系统架构

CUDA 系统架构由以下三部分构成，为开发者提供丰富的资源和接口：

开发库
驱动
运行期环境

CUDA应用场景

CUDA 最初用于图形渲染，现已广泛应用于：

高性能计算
游戏开发
深度学习等领域

编程语言支持

CUDA 支持多种编程语言，包括：

C
C++
Fortran
Java
Python等

但本次源码阅读和分析以 C/C++ 为主。

CUDA计算过程

分配 host 内存并进行数据初始化。
分配 device 内存，并将数据从 host 拷贝到 device 上。
调用 CUDA 核函数在 device 上完成指定的运算。
将 device 上的运算结果拷贝到 host 上。
释放 device 和 host 上分配的内存。

线程层次

Grid：最高级的并行计算单位，由多个线程块（blocks）组成，可同时在 GPU 上执行。
Block：线程块是网格的组成部分，包含多个线程，可协作和共享资源。
Thread：最基本的并行执行单元，存在于线程块内部，每个线程可以独立执行计算任务，根据自身的线程索引执行。

优势：

有效管理并行计算资源。
优化数据访问和内存使用。
提高任务并行度和负载均衡。
适应不同的并行计算需求。

存储层次

全局内存（Global Memory）：GPU 中所有线程都可访问的主要存储区域，用于存储大量数据和程序状态，是主机与设备之间数据传输的主要桥梁。
常量内存（Constant Memory）：用于存储只读数据，如常数和预计算的表格，具有较高的缓存性能。
纹理内存（Texture Memory）：存储图像和多维数据，提供高效、灵活的访问方式。
共享内存（Shared Memory）：位于线程块内部的低延迟、高带宽的存储区域，用于线程块内部的数据共享和通信。

2 CUDA 系统架构

分层架构

从体系结构的组成来说，CUDA 包含了三个部分：

开发库 (Libraries)：它是基于 CUDA 技术所提供的应用开发库。CUDA 包含两个重要的标准数学运算库——CUFFT（离散快速傅立叶变换）和 CUBLAS（基本线性代数子程序库）。这两个数学运算库解决的是典型的大规模的并行计算问题，也是在密集数据计算中非常常见的计算类型。
运行期环境 (Runtime)：提供了应用开发接口和运行期组件，包括基本数据类型的定义和各类计算、类型转换、内存管理、设备访问和执行调度等函数。例如：
- 运行时 API：提供了一系列函数，用于执行设备内存分配、内核启动、事件管理、流控制等操作。
- 核函数：这是在 GPU 上并行执行的函数，由开发者编写，用于执行具体的并行计算任务。
驱动 (Driver)：驱动程序层位于硬件层之上，是 CUDA 架构中的中间件，负责硬件抽象、资源管理、错误处理、CPU 和 GPU 通信。驱动程序层通过 CUDA 运行时 API 与主机代码交互。

并行计算模式

在这里插入图片描述

SIMT (Single Instruction Multiple Threads)：单指令多线程是 CUDA 架构中的重要执行模型，它允许一组线程（称为线程束或 warp）同时执行相同的指令，但处理不同的数据。
- 每个 warp 包含一定数量的 core 用于执行，warp 是硬件调度的并行的最小单位。
- 与 CPU 不同的是，GPU 的执行调度单位为 warp，而 CPU 调度可以精细化为单个线程。因为 CPU 同时执行的线程数目比较少，而 GPU 要同时管理成千上万个线程，因此 GPU 一次性调度的是多个线程。
Grids, Blocks, Threads：具体到编程层面，程序员可以通过 Grids、Blocks 和 Threads 的编程模型来组织执行的线程：
- Threads (线程)：是 CUDA 中最小的执行单位。每个线程执行核函数中的指令，并且可以独立处理数据。
- Blocks (块)：一个 Block 是一组线程的集合，这些线程可以共享一个内存空间。Block 内的线程通过 threadIdx 进行索引。
- Grids (网格)：Grid 是多个 Block 的集合，构成了核函数调用的全部执行范围。Grid 内的 Block 通过 blockIdx 进行索引。
- 当启动一个核函数时，CUDA 运行时环境会创建一个 Grid，该 Grid 由多个 Blocks 组成，每个 Block 又包含多个 Threads。

示例代码：

__global__ void vectorAdd(int *A, int *B, int *C, int N) {
   
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
   
        C[idx] = A[idx] + B[idx];
    }
}
int main() {
   
    int N = 1000;
    int *A, *B, *C;
    int block_size = 256;
    int grid_size = (N + block_size - 1) / block_size;
    // 分配和初始化设备内存
    cudaMalloc(&A, N * sizeof(int));
    cudaMalloc(&B, N * sizeof(int));
    cudaMalloc(&C, N * sizeof(int));
    // 将数据从主机复制到设备
    cudaMemcpy(A, host_A, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(B, host_B, N * sizeof(int), cudaMemcpyHostToDevice);
    // 启动核函数
    vectorAdd<<<grid_size, block_size>>>(A, B, C, N);
    // 等待核函数执行完成
    cudaDeviceSynchronize();
    // 将结果从设备复制回主机
    cudaMemcpy(host_C, C, N * sizeof(int),  cudaMemcpyDeviceToHost);
    // 释放设备内存
    cudaFree(A);
    cudaFree(B);
    cudaFree(C);
    return 0;}

源码分析：

核函数 vectorAdd() 定义了向量加法的操作；
int idx = blockIdx.x * blockDim.x + threadIdx.x;：计算出了当前线程的全局索引。blockIdx.x 表示当前线程所在的块的索引，blockDim.x 表示当前块中线程的数量，threadIdx.x 表示当前线程在块内的索引。
在 main 函数中，定义了 Blocks 的数量（grid_size）和每个 Block 的线程数量（block_size），并启动核函数 vectorAdd<<<grid_size, block_size>>>。
核函数中的每个线程都将执行相同的加法操作，但操作不同的数据元素，也就是向量中不同维度的数据。