CUDA架构介绍与设计模式解析-易微帮

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>>>。
核函数中的每个线程都将执行相同的加法操作，但操作不同的数据元素，也就是向量中不同维度的数据。

生产-消费者模式

CUDA 中的生产者-消费者模式用于解决主机（CPU）和设备（GPU）之间的数据传输和计算任务之间的协作。主机（CPU）充当生产者，负责生成指令和数据并将其传输到设备上，而设备（GPU）充当消费者，负责获取数据并执行计算任务。

生产者-消费者模式的主要特点包括：

数据传输：生产者负责生成数据，并将数据传输到设备的内存中。这可以通过 CUDA 的内存复制函数（如 cudaMemcpy）来完成。
计算任务：消费者负责从设备的内存中获取数据，并执行相应的计算任务，通常是通过启动核函数来实现的。
数据同步：CUDA 使用流（Stream）来管理并行执行的操作序列，进行数据同步操作，以确保数据传输和计算任务的正确执行顺序。

示例代码：

#include <iostream>
__global__ void kernel(int *data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        data[idx] *= 2;
    }
}
int main() {
    const int N = 1000;
    int *hostData, *deviceData;
    
    // 分配主机内存并初始化数据
    hostData = new int[N];
    for (int i = 0; i < N; ++i) {
        hostData[i] = i;
    }
    // 分配设备内存
    cudaMalloc(&deviceData, N * sizeof(int));
    // 将数据从主机复制到设备
    cudaMemcpy(deviceData, hostData, N * sizeof(int), cudaMemcpyHostToDevice);
    // 启动核函数，在设备上执行计算任务
    const int blockSize = 256;
    const int gridSize = (N + blockSize - 1) / blockSize;
    kernel<<<gridSize, blockSize>>>(deviceData, N);
    // 将结果从设备复制回主机
    cudaMemcpy(hostData, deviceData, N * sizeof(int), cudaMemcpyDeviceToHost);
    // 打印结果
    for (int i = 0; i < N; ++i) {
        std::cout << hostData[i] << " ";
    }
    std::cout << std::endl;
    // 释放内存
    delete[] hostData;
    cudaFree(deviceData);
    return 0;
}

源码分析：

主机（生产者）生成了一个包含 1000 个整数的数据，并将数据传输到设备上。
设备（消费者）执行了核函数，进行具体的处理。
最后，设备将处理后的数据传输回主机，主机打印出结果。
通过这种方式，实现了主机和设备之间的生产者-消费者模式，充分利用了 GPU 的并行计算能力。

工作池模式

CUDA 中的工作池模式用于管理和调度大量的任务，以实现高效的并行计算。在工作池模式中，任务被组织成一个池子（工作池），并由多个工作线程并行执行。这种模式使得可以在 GPU 上同时处理多个任务，从而充分利用 GPU 的并行计算资源。

工作池模式的关键组件：

任务队列（Task Queue）：用于存储待执行的任务。任务可以是任何需要在 GPU 上执行的操作，例如核函数调用、数据传输等。
工作线程（Worker Threads）：负责从任务队列中获取任务，并执行这些任务。通常，工作线程会持续地从任务队列中获取任务，并在 GPU 上执行这些任务，直到任务队列为空或者达到某个终止条件。
任务调度器（Task Scheduler）：负责调度工作线程的执行。任务调度器通常会根据任务队列的状态和工作线程的负载情况，动态地分配任务给工作线程，以实现任务的均衡分配和高效执行。

示例代码：

// 任务类型
typedef void (*Task)(int *, int);
// 任务队列
std::queue<Task> taskQueue;
std::mutex queueMutex;
std::condition_variable queueCondition;
// 工作线程函数
void workerThread(int *array, int N) {
    while (true) {
        Task task;
        {
            std::unique_lock<std::mutex> lock(queueMutex);
            // 等待任务队列非空
            queueCondition.wait(lock, [](){ return !taskQueue.empty(); });
            // 从任务队列中获取任务
            task = taskQueue.front();
            taskQueue.pop();
        }
        // 执行任务
        task(array, N);
    }
}
// 核函数，将数组中的每个元素乘以2
__global__ void kernel(int *array, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        array[idx] *= 2;
    }
}
// 添加任务到任务队列
void addTask(int *array, int N) {
    taskQueue.push(kernel);
    queueCondition.notify_one(); // 唤醒一个工作线程
}
int main() {
    const int N = 1000;
    int *d_array;
    cudaMalloc(&d_array, N * sizeof(int));
    // 定义并行执行的任务数量
    const int numTasks = 10;
    // 创建并启动工作线程
    std::thread threads[numTasks];
    for (int i = 0; i < numTasks; ++i) {
        threads[i] = std::thread(workerThread, d_array, N);
    }
    // 添加多个任务到任务队列
    for (int i = 0; i < numTasks; ++i) {
        addTask(d_array, N);
    }
    // 等待所有工作线程执行完成
    for (int i = 0; i < numTasks; ++i) {
        threads[i].join();
    }
    // 处理执行结果...
    cudaFree(d_array);
    return 0;
}

源码分析：

在 main 中创建一组任务和工作线程。
每个工作线程不断地从任务队列中获取任务，并在 GPU 上执行这些任务，直到任务队列为空。

异步编程模式

CUDA 中的异步编程模式是指，在主机和设备之间并行执行多个任务，从而提高了整体的计算和数据传输效率。

CUDA 中异步编程模式的特点：

并行执行：允许主机和设备之间同时执行多个任务，从而避免了在等待一个任务完成时的空闲时间，提高了整体的计算和数据传输效率。
任务调度：CUDA 运行时会自动管理和调度异步任务的执行顺序，以最大程度地利用 GPU 的并行计算资源，减少任务之间的等待时间。
流控制：开发者可以通过 CUDA 中的流（Stream）对异步任务进行组织和管理。通过流控制，开发者可以控制任务的执行顺序，以及任务之间的依赖关系，从而实现更复杂的并行计算和数据传输操作。

示例代码：

#include <iostream>
__global__ void kernel(int *array, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        array[idx] *= 2;
    }
}
int main() {
    const int N = 1000;
    int *d_array;
    cudaMalloc(&d_array, N * sizeof(int));
    // 创建一个流
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    // 异步启动核函数
    kernel<<<(N + 255) / 256, 256, 0, stream>>>(d_array, N);
    // 主机代码继续执行其他任务...
    // 等待流中的核函数执行完成
    cudaStreamSynchronize(stream);
    // 处理核函数执行结果...
    // 释放流
    cudaStreamDestroy(stream);
    cudaFree(d_array);
    return 0;
}

源码分析：

主机创建一个流，并把核函数绑定到这个流实现异步启动，这样主机可以继续执行其他任务。
然后，主机通过 cudaStreamSynchronize 函数等待流中的核函数执行完成，以确保可以安全地处理核函数的执行结果。
通过异步编程模式，主机和设备之间实现了并行执行，提高了整体的计算和数据传输效率。

3 CUDA 中的设计模式

工厂模式

工厂模式（Factory Pattern）是一种常用的软件设计模式，属于创建型模式。其核心思想是通过工厂类中的一个共用方法来处理对象的创建，将对象的创建和对象的使用分离开来，以便于达到扩展性高、耦合度低的目的。

工厂模式的基本组成：

产品 (Product)：定义了产品的接口。
具体产品 (Concrete Product)：实现了产品接口的具体类。
工厂 (Factory)：提供了一个创建产品的方法，可以返回不同类型的产品对象。

内存分配函数 cudaMalloc

示例代码：

#include <stdio.h>
#include <cuda_runtime.h>
int main() {
  int size = 1024 * sizeof(int);
  int *d_array;
  // 使用工厂模式的 cudaMalloc 在设备上分配内存
  cudaMalloc((void**)&d_array, size);
  // 内存分配成功后，可以在设备上使用 d_array

  // ...
  cudaFree(d_array); // 释放设备内存
  return 0;
}

源码分析：

cudaMalloc 函数是 CUDA Runtime API 提供的内存分配函数，其定义如下：

cudaError_t cudaMalloc(void** devPtr, size_t size);

devPtr 是一个指向设备内存指针的指针，在函数调用后将指向分配的设备内存。
size 是要分配的内存大小（字节数）。
函数返回 cudaError_t 类型的错误码，表示函数执行的状态。

策略模式

策略模式（Strategy Pattern）是一种定义一系列算法的软件设计模式，它属于行为型模式。策略模式允许客户端在运行时根据需要选择不同的算法，而不需要修改客户端代码。其目的是将算法的使用从算法的实现中分离出来，让两者之间不相互依赖，从而提高程序的灵活性和可扩展性。

策略模式的基本组成：

策略 (Strategy)：定义所有支持的算法的公共接口。
具体策略 (Concrete Strategy)：实现了策略接口的具体算法类。
环境 (Context)：持有一个策略类的引用，用该策略的接口来与之交互。

示例代码：

#include <iostream>
#include <cuda_runtime.h>
class ParallelizationStrategy {
public:
    virtual void execute(int* data, int size) const = 0;
    virtual ~ParallelizationStrategy() {}
};
class ThreadPerElementStrategy : public ParallelizationStrategy {//基于线程的并行化
public:
    virtual void execute(int* data, int size) const override {
        // Execute algorithm with one thread per element
        int idx = threadIdx.x + blockIdx.x * blockDim.x;
        if (idx < size) {
            // Do computation
        }
    }
};
class ThreadBlockStrategy : public ParallelizationStrategy {//基于线程块的并行化
public:
    virtual void execute(int* data, int size) const override {
        // Execute algorithm with one thread per block
        int idx = blockIdx.x;
        if (idx < size) {
            // Do computation
        }
    }
};
class Kernel {
private:
    ParallelizationStrategy* strategy;

public:
    Kernel(ParallelizationStrategy* strat) : strategy(strat) {}

    void run(int* data, int size) const {
        strategy->execute(data, size);
    }
};
int main() {
    const int dataSize = 1000;
    int* data;
    cudaMalloc(&data, dataSize * sizeof(int));

    ThreadPerElementStrategy perElementStrategy;
    Kernel kernel1(&perElementStrategy);
    kernel1.run(data, dataSize);

    ThreadBlockStrategy blockStrategy;
    Kernel kernel2(&blockStrategy);
    kernel2.run(data, dataSize);

    cudaFree(data);
    return 0;
}

源码分析：

在主函数中创建了两个不同的Kernel对象，分别使用不同的策略来执行内核函数。这样就实现了在运行时选择不同的算法行为，而不需要改变客户端代码的逻辑。

观察者模式

观察者模式（Observer Pattern）是一种行为型设计模式，它定义了对象之间的一对多依赖关系。当一个对象（被观察者）的状态发生变化时，所有依赖它的对象（观察者）都会收到通知并自动更新。这种模式常用于事件驱动的系统，以实现松耦合。

观察者模式的基本组成：

主题 (Subject)：维护了一个观察者列表，并能够触发观察者列表中的所有观察者。
观察者 (Observer)：定义了一个更新接口，用以更新状态。
具体主题 (Concrete Subject)：实现了主题，具体通知观察者。
具体观察者 (Concrete Observer)：实现了观察者的更新接口，以反映主题的状态。

CUDA 事件可以用来监视核函数执行的状态，并在完成时通知主机。

示例代码：

cudaEvent_t event;
cudaEventCreate(&event); // 创建事件
myKernel<<<numBlocks, threadsPerBlock>>>();//核函数执行
cudaEventRecord(event); // 记录事件，核函数执行完成后触发事件
cudaEventSynchronize(event); // 等待事件完成
// 处理事件完成后的逻辑
…

适配者模式

迭代器模式（Iterator Pattern）是软件设计中的一种常用行为型模式，它允许客户端在不暴露其底层数据结构的情况下，顺序访问聚合对象中的元素。迭代器模式定义了一个数据访问接口，用于在元素之间进行导航，而不依赖于聚合对象的具体类。

迭代器模式的基本组成：

迭代器 (Iterator)：定义了访问和遍历元素的接口。
具体迭代器 (Concrete Iterator)：实现了迭代器接口，跟踪当前访问的位置。
聚合 (Aggregate)：定义了一个创建迭代器的方法。
具体聚合 (Concrete Aggregate)：实现了创建具体迭代器的方法。

有两个版本的核函数，它们执行相同的操作，但是实现方式不同。在主机代码中使用一个统一的接口来调用这些核函数，而不需要关心它们的具体实现细节。

示例代码：

// 核函数版本1

__global__ void kernelVersion1(int *data, int value) {
  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  if (tid < size) {
    data[tid] += value;
  }
}
// 核函数版本2
__global__ void kernelVersion2(int *data, int value) {
  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  atomicAdd(&data[tid], value); // 使用原子操作来更新数据
}
// 核函数版本1的适配器
class KernelLauncherV1 : public IKernelLauncher {
public:
    void launch(int *data, int value, int size) override {
        kernelVersion1<<<...>>>(data, value);
    }
};
// 核函数版本2的适配器
class KernelLauncherV2 : public IKernelLauncher {
public:
    void launch(int *data, int value, int size) override {
        kernelVersion2<<<...>>>(data, value);
    }
};
int main() {
    int *data; // 假设已经分配和初始化了设备内存
    int value = 10;
    int size = 1024;

    // 根据需要选择使用哪个版本的核函数
    std::unique_ptr<IKernelLauncher> launcher;
    if (useVersion1) {
        launcher = std::make_unique<KernelLauncherV1>();
    } else {
        launcher = std::make_unique<KernelLauncherV2>();
    }

    // 使用适配器接口调用核函数
    launcher->launch(data, value, size);

    cudaDeviceSynchronize(); // 等待GPU完成工作
    // ... 后续处理 ...

    return 0;
}

源码分析：

在上述CUDA适配器模式的示例中，其创建了一个抽象的IKernelLauncher接口，该接口定义了启动核函数的通用方法launch。这个接口允许我们通过统一的调用方式来执行不同的核函数实现，而不必在主机代码中硬编码每个核函数的具体调用。

迭代器模式

迭代器模式的基本组成：

迭代器 (Iterator)：定义了访问和遍历元素的接口，通常包含 next(), prev(), first(), last(), isDone() 和 currentItem() 等方法。
具体迭代器 (Concrete Iterator)：实现了迭代器接口，跟踪当前访问的位置，并在内部维护一个指向聚合对象中某个位置的引用。
聚合 (Aggregate)：定义了一个创建迭代器的方法，通常是一个 createIterator() 方法，用于返回一个迭代器的实例。
具体聚合 (Concrete Aggregate)：实现了创建具体迭代器的方法，返回一个具体迭代器的实例，该实例能够遍历具体聚合中的所有元素。

Ps:迭代器模式在CUDA中不是直接应用的，因为CUDA的编程模型主要关注于如何高效地映射并行计算到GPU上。迭代器模式在CUDA中的一种可能应用场景：遍历设备内存中的数组。

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

文章目录