CUDA执行模型

最新推荐文章于 2024-10-15 11:52:15 发布

A11en3

最新推荐文章于 2024-10-15 11:52:15 发布

阅读量1k

点赞数 21

分类专栏： CUDA编程文章标签： c++ CUDA

本文链接：https://blog.csdn.net/A11en3/article/details/138200917

版权

CUDA编程专栏收录该内容

17 篇文章 0 订阅

订阅专栏

本文详细介绍了GPU架构，特别是SM组件的功能，CUDA的核心原理，以及SIMT和SIMD的区别。重点讨论了开普勒架构的动态并行特性和Hyper-Q技术，以及如何避免线程束分化以优化性能。

摘要由CSDN通过智能技术生成

GPU架构

GPU的架构是围绕一个流式多处理器（SM）的可扩展阵列搭建的。可以通过复制这种架构的构建块来实现GPU的硬件并行。

SM是GPU中的一个重要组件，负责执行CUDA核函数中的线程块（Thread Block）。每个SM包含多个CUDA核心，可以同时执行多个线程块中的线程，以实现并行计算。

每个CUDA核心有自己的缓存器，也可以访问同一个SM中的共享内存。每个CUDA核心可以并行多个线程。简单来说，一个GPU可以并行的最大线程数=CUDA核心数*一个CUDA核心可以并行的线程数。

CUDA采用单指令多线程（SIMT）架构来管理和执行线程，每32个线程为一组，被称为线程束。线程束中所有线程同时执行相同的指令。每个SM都将分配给它的线程块划分到包含32个线程的线程束中，然后在可用的硬件资源上调度执行。

SIMT架构与SIMD（单指令多数据）架构相似。两者都是将相同的指令广播给多个执行单元来实现并行。一个关键的区别是SIMD要求同一个向量中的所有元素要在一个统一的同步组中一起执行，而SIMT允许属于一个线程束的多个线程独立执行。

一个线程束中的所有线程在相同的程序地址上同时开始执行，但是单独的线程仍有可能有不同的行为。

SIMT架构有三个SIMD所不具备的关键特征

每个线程都有自己的指令地址计数器
每个线程都有自己的寄存器状态
每个线程可以有一个独立的执行路径

一个线程块只能在一个SM上被调度，一旦开始调度，就会保存在该SM上直到执行完成。在同一时间，一个SM可以容纳多个线程块。

SM中的共享内存和寄存器

共享内存：
- 共享内存是SM中的一种存储器，用于在线程块中的线程之间共享数据。
- 共享内存的优势在于其访问速度非常快，比全局内存要快得多，因为共享内存是位于SM内部的，而全局内存则位于GPU芯片外部，访问速度较慢。
- 共享内存的大小是有限的，通常每个SM都有一定量的共享内存可用，线程块中的所有线程都可以访问这段共享内存。
- 开发者可以通过声明__shared__关键字来定义共享内存变量，例如：__shared__ float sharedData[256];。
寄存器：
- 寄存器是SM中的另一种存储器，用于存储每个线程的私有数据和状态。
- 每个线程都有自己的一组寄存器，用于存储局部变量和中间计算结果。
- 寄存器的访问速度非常快，但是寄存器的数量是有限的，如果每个线程使用的寄存器过多，可能会导致寄存器溢出，从而影响性能。
- 开发者无法直接控制线程使用的寄存器数量，但可以通过调整编译器的优化选项来尽量减少寄存器的使用量。

开普勒架构的动态并行

开普勒架构引入了动态并行的特性，允许GPU启动新的网格。也就是说，任一内核都可以启动其他的内核。这一特性使得我们更容易创建和优化递归及动态数据相关的执行模式。

https://blog.51cto.com/u_14125/9699692查看自己GPU的硬件信息

开普勒架构的Hyper-Q技术

开普勒架构（Kepler Architecture）引入了一项名为Hyper-Q的技术，旨在提高GPU的并行处理能力和效率。Hyper-Q技术允许多个CPU核心同时向GPU提交任务，这些任务可以独立地在GPU上执行，从而提高了CPU与GPU之间的并行性。

具体来说，Hyper-Q技术包括以下几个关键点：

并发处理：Hyper-Q技术允许多个CPU核心同时向GPU提交任务，这些任务可以在GPU上并发执行，提高了CPU与GPU之间的并行处理能力。
任务隔离：Hyper-Q技术通过为每个CPU核心提供独立的任务队列，可以隔离不同CPU核心提交的任务，防止它们之间相互干扰，提高了任务的执行效率。
任务调度：Hyper-Q技术通过智能调度算法，可以有效地管理和调度GPU上的任务，确保任务之间的高效并行执行。

Hyper-Q技术的任务调度算法旨在提高GPU上的任务并行度和效率，确保任务能够以最佳方式在GPU上执行。以下是Hyper-Q技术的任务调度算法的一般工作流程：

任务提交：CPU上的多个核心可以同时向GPU提交任务。每个CPU核心有自己的任务队列，可以将任务添加到队列中。
任务分配：GPU中的多个SM（流多处理器）可以从不同的任务队列中获取任务。任务分配的方式可以根据任务队列的状态和GPU资源的可用性进行动态调整。
任务执行：SM在获取任务后，将任务分配给其内部的CUDA核心进行执行。SM可以同时执行多个任务，从而提高GPU的利用率和并行性。
任务完成：任务执行完成后，SM将结果返回给主机CPU或者将结果存储到GPU内存中。完成的任务会从任务队列中移除。
任务调度策略：Hyper-Q技术的任务调度策略包括多种因素，如任务的优先级、任务队列的长度、GPU资源的利用率等。任务调度算法会根据这些因素动态地调整任务的执行顺序和方式，以保证任务能够以最佳方式在GPU上执行。

线程束执行

线程束和线程块

线程束是SM中基本的执行单元。当一个线程块的网格被启动后，网格中的线程块分布在SM中。一旦线程块被调度到一个SM上，线程块中的线程会被进一步划分为线程束。

在一个线程束中，所有线程按照SIMT方式执行；所有线程执行相同的指令，每个线程在私有数据上进行操作。

每个线程都可以在线程块中被唯一标识出来，每个线程的独特标识符都可以由threadIdx与blockDim计算出来。

一维线程块：threadIdx.x
二维线程块：threadIdx.y * blockDim.x + threadIdx.x
三维线程块：threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x

GPU总是给一个线程块分配一定数量的线程束。线程束不会再线程块之间分离。比如，如果一个二维线程块，维度为40*2。那么，由于线程在硬件上是一维分配的，硬件会给这个线程块分配3个线程束，即32*3个硬件线程。由这96个硬件线程支持80个软件线程，即使存在某些不活跃的硬件线程，但是它们依旧占用SM的资源。

线程块的逻辑角度与硬件角度

从逻辑角度来看，线程块是线程的集合，可以被组织为一维，二维，三维。

从硬件角度来看，线程块是一维线程束的集合。在线程块中被组织为一维布局。

线程束分化

GPU是相对简单的设备，没有CPU那样的复杂的分支预测机制。一个线程束中的所有线程在同一周期中必须执行相同的命令，如果一个线程执行一个命令，那么所有线程都要执行该命令。

如果同一线程束中的不同线程使用不同的路径通过同一个应用程序，可能会产生问题。例如

if(cond){
    ...
}
else{
    ...
}

因为SIMT执行模式，在所有线程中，cond的值可能是不同的。一部分线程执行if语句块，另一部分线程执行else语句块。这种在同一线程束中的线程执行不同的指令，被称为线程束分化。

为了获得最佳性能，应该避免在同一线程束中有不同的执行路径。在编码时，对数据进行分区以确保同一个线程束中的所有线程在一个应用程序中使用同一个控制路径。

获取CUDA设备的线程束大小
    int warpSize;
    //cudaDeviceProp prop;

    //cudaGetDeviceProperties(&prop, 0);
    cudaDeviceGetAttribute(&warpSize, cudaDevAttrWarpSize, 0);

    cout << "CUDA wrap size is " << warpSize;
获取第0个设备的线程束大小，cudaGetDeviceProperties函数用来获取设备信息

使用线程束的方法来交叉存储数据，可以避免线程束分化，并且设备的利用率可以达到100%。

使用

__global__ void func(){
    int tid = threadIdx.y * blockDim.x + threadIdx.x;

    if((tid / warpSize) % 2==0){
        ...
    }
    else{
        ...
    }
}

使得分支粒度是线程束大小的倍数。