cuda编程_CUDA刷新器：CUDA编程模型

CUDA刷新器&＃xff1a;CUDA编程模型

CUDA Refresher: The CUDA Programming Model

CUDA&＃xff0c;CUDA刷新器&＃xff0c;并行编程

这是CUDA更新系列的第四篇文章&＃xff0c;它的目标是刷新CUDA中的关键概念、工具和初级或中级开发人员的优化。

CUDA编程模型提供了GPU体系结构的抽象&＃xff0c;它充当了应用程序与其在GPU硬件上的可能实现之间的桥梁。这篇文章概述了CUDA编程模型的主要概念&＃xff0c;概述了它如何在通用编程语言如C/C&＃43;&＃43;中暴露出来。

介绍一下CUDA编程模型中常用的两个关键词&＃xff1a;主机和设备。

主机是系统中可用的CPU。与CPU相关联的系统内存称为主机内存。GPU被称为设备&＃xff0c;GPU内存也被称为设备内存。

要执行任何CUDA程序&＃xff0c;有三个主要步骤&＃xff1a;

将输入数据从主机内存复制到设备内存&＃xff0c;也称为主机到设备传输。

加载GPU程序并执行&＃xff0c;在片上缓存数据以提高性能。

将结果从设备内存复制到主机内存&＃xff0c;也称为设备到主机传输。

CUDA内核和线程层次结构

图1显示了CUDA内核是一个在GPU上执行的函数。应用程序的并行部分由K个不同的CUDA线程并行执行k次&＃xff0c;而不是像常规C/C&＃43;&＃43;函数那样只进行一次。

Figure 1. The kernel is a function executed on the GPU.

每一个CUDA内核都以一个__global__声明说明符开头。程序员通过使用内置变量为每个线程提供唯一的全局ID。

图2. CUDA内核被细分为块。

一组线程称为CUDA块。CUDA块被分组到一个网格中。内核作为线程块的网格来执行&＃xff08;图2&＃xff09;。

每个CUDA块由一个流式多处理器&＃xff08;SM&＃xff09;执行&＃xff0c;不能迁移到GPU中的其他SMs&＃xff08;抢占、调试或CUDA动态并行期间除外&＃xff09;。一个SM可以根据CUDA块所需的资源运行多个并发CUDA块。每个内核在一个设备上执行&＃xff0c;CUDA支持一次在一个设备上运行多个内核。图3显示了GPU中可用硬件资源的内核执行和映射。

图3. 在GPU上执行内核。

CUDA为线程和块定义了内置的三维变量。线程使用内置的三维变量threadIdx编制索引。三维索引提供了一种自然的方法来索引向量、矩阵和体积中的元素&＃xff0c;并使CUDA编程更容易。类似地&＃xff0c;块也使用名为blockIdx的内置三维变量编制索引。

以下是几个值得注意的要点&＃xff1a;

CUDA架构限制每个块的线程数&＃xff08;每个块限制1024个线程&＃xff09;。

线程块的维度可以通过内置的blockDim变量在内核中访问。

syncu中的线程可以使用syncu函数同步。使用同步线程时&＃xff0c;块中的所有线程都必须等待&＃xff0c;然后才能继续。

在<<…>>>语法中指定的每个块的线程数和每个网格的块数可以是int或dim3类型。这些三角括号标记从主机代码到设备代码的调用。它也被称为内核启动。

下面用于添加两个矩阵的CUDA程序显示多维blockIdx和threadIdx以及blockDim等其他变量。在下面的例子中&＃xff0c;为了便于索引&＃xff0c;选择了一个2D块&＃xff0c;每个块有256个线程&＃xff0c;x和y方向各有16个线程。使用数据大小除以每个块的大小来计算块的总数。

// Kernel - Adding two matrices MatA and MatB__global__ void MatAdd(float MatA[N][N], float MatB[N][N],float MatC[N][N]){ int i &＃61; blockIdx.x * blockDim.x &＃43; threadIdx.x; int j &＃61; blockIdx.y * blockDim.y &＃43; threadIdx.y; if (i MatC[i][j] &＃61; MatA[i][j] &＃43; MatB[i][j];}int main(){ ... // Matrix addition kernel launch from host code dim3 threadsPerBlock(16, 16); dim3 numBlocks((N &＃43; threadsPerBlock.x -1) / threadsPerBlock.x, (N&＃43;threadsPerBlock.y -1) / threadsPerBlock.y); MatAdd<<>>(MatA, MatB, MatC); ...}
Memory hierarchy

支持CUDA的GPU有一个内存层次结构&＃xff0c;如图4所示。

图4. gpu中的内存层次结构。

以下内存由GPU架构公开&＃xff1a;

这些寄存器对每个线程都是私有的&＃xff0c;这意味着分配给线程的寄存器对其他线程不可见。编译器决定寄存器的利用率。

一级/共享内存&＃xff08;SMEM&＃xff09;-每个SM都有一个快速的片上草稿行内存&＃xff0c;可用作一级缓存和共享内存。CUDA块中的所有线程都可以共享共享内存&＃xff0c;在给定SM上运行的所有CUDA块都可以共享SM提供的物理内存资源。。

只读内存每个SM都有一个指令缓存、常量内存、纹理内存和对内核代码只读的RO缓存。 二级缓存二级缓存在所有SMs中共享&＃xff0c;因此每个CUDA块中的每个线程都可以访问该内存。nvidiaa100 GPU已经将二级缓存大小增加到40mb&＃xff0c;而v100gpu中只有6mb。 全局内存这是位于GPU中的GPU和DRAM的帧缓冲区大小。

NVIDIA CUDA编译器在优化内存资源方面做得很好&＃xff0c;但专家CUDA开发人员可以选择有效地使用这种内存层次结构来优化CUDA程序。

计算能力

GPU的计算能力决定了GPU硬件支持的通用规范和可用特性。此版本号可由应用程序在运行时使用&＃xff0c;以确定当前GPU上可用的硬件功能或指令。

每个GPU都有一个版本号&＃xff0c;表示为X.Y&＃xff0c;其中X包括主要修订号&＃xff0c;Y包含次要修订号。小版本号对应于架构的增量改进&＃xff0c;可能包括新特性。

有关任何支持CUDA的设备的计算能力的更多信息&＃xff0c;请参阅CUDA示例代码设备查询。此示例枚举系统中存在的CUDA设备的属性

摘要

CUDA编程模型提供了一种异构环境&＃xff0c;其中主机代码在CPU上运行C/C&＃43;&＃43;程序&＃xff0c;内核在物理上分离的GPU设备上运行。CUDA编程模型还假设主机和设备都保持各自独立的内存空间&＃xff0c;分别称为主机内存和设备内存。CUDA代码还通过PCIe总线提供主机和设备内存之间的数据传输。

CUDA还公开了许多内置变量&＃xff0c;并提供了多维索引的灵活性&＃xff0c;以简化编程。CUDA还管理不同的内存&＃xff0c;包括寄存器、共享内存和一级缓存、二级缓存和全局内存。高级开发人员可以有效地使用这些内存来优化CUDA程序。