教育平台网站建设,手机网站模版,制作网站教程,淘宝做网站的店2.编程模型
更多精彩内容#xff0c;请扫描下方二维码或者访问https://developer.nvidia.com/zh-cn/developer-program 来加入NVIDIA开发者计划 本章通过概述CUDA编程模型是如何在c中公开的#xff0c;来介绍CUDA的主要概念。
编程接口中给出了对 CUDA C 的广泛描述。
本章…2.编程模型
更多精彩内容请扫描下方二维码或者访问https://developer.nvidia.com/zh-cn/developer-program 来加入NVIDIA开发者计划 本章通过概述CUDA编程模型是如何在c中公开的来介绍CUDA的主要概念。
编程接口中给出了对 CUDA C 的广泛描述。
本章和下一章中使用的向量加法示例的完整代码可以在 vectorAddCUDA示例中找到。
2.1 内核
CUDA C 通过允许程序员定义称为Kernel的 C 函数来扩展 C当调用内核时由 N 个不同的 CUDA 线程并行执行 N 次而不是像常规 C 函数那样只执行一次。
使用 __global__ 声明说明符定义内核并使用新的 ... 执行配置语法指定内核调用的 CUDA 线程数请参阅 C语言扩展。 每个执行内核的线程都有一个唯一的线程 ID可以通过内置变量在内核中访问。
作为说明以下示例代码使用内置变量 threadIdx 将两个大小为 N 的向量 A 和 B 相加并将结果存储到向量 C 中 // Kernel definition__global__ void VecAdd(float* A, float* B, float* C){int i threadIdx.x;C[i] A[i] B[i];}int main(){...// Kernel invocation with N threadsVecAdd1, N(A, B, C);...}这里执行 VecAdd() 的 N 个线程中的每一个线程都会执行一个加法。
2.2 线程层次
为方便起见threadIdx 是一个 3 分量向量因此可以使用一维、二维或三维的线程索引来识别线程形成一个一维、二维或三维的线程块称为block。这提供了一种跨域的元素例如向量、矩阵或体积调用计算的方法。
线程的索引和它的线程 ID 以一种直接的方式相互关联对于一维块它们是相同的 对于大小为(Dx, Dy)的二维块索引为(x, y)的线程的线程ID为(x yDx) 对于大小为 (Dx, Dy, Dz) 的三维块索引为 (x, y, z) 的线程的线程 ID 为 (x yDx zDxDy)。
例如下面的代码将两个大小为NxN的矩阵A和B相加并将结果存储到矩阵C中: // Kernel definition__global__ void MatAdd(float A[N][N], float B[N][N],float C[N][N]){int i threadIdx.x;int j threadIdx.y;C[i][j] A[i][j] B[i][j];}int main(){...// Kernel invocation with one block of N * N * 1 threadsint numBlocks 1;dim3 threadsPerBlock(N, N);MatAddnumBlocks, threadsPerBlock(A, B, C);...}每个块的线程数量是有限制的因为一个块的所有线程都应该驻留在同一个处理器核心上并且必须共享该核心有限的内存资源。在当前的gpu上一个线程块可能包含多达1024个线程。
但是一个内核可以由多个形状相同的线程块执行因此线程总数等于每个块的线程数乘以块数。
块被组织成一维、二维或三维的线程块网格(grid)如下图所示。网格中的线程块数量通常由正在处理的数据的大小决定通常超过系统中的处理器数量。 ... 语法中指定的每个块的线程数和每个网格的块数可以是 int或 dim3 类型。如上例所示可以指定二维块或网格。
网格中的每个块都可以由一个一维、二维或三维的惟一索引标识该索引可以通过内置的blockIdx变量在内核中访问。线程块的维度可以通过内置的blockDim变量在内核中访问。
扩展前面的MatAdd()示例来处理多个块代码如下所示。 // Kernel definition__global__ void MatAdd(float A[N][N], float B[N][N],float C[N][N]){int i blockIdx.x * blockDim.x threadIdx.x;int j blockIdx.y * blockDim.y threadIdx.y;if (i N j N)C[i][j] A[i][j] B[i][j];}int main(){...// Kernel invocationdim3 threadsPerBlock(16, 16);dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);MatAddnumBlocks, threadsPerBlock(A, B, C);...}线程块大小为16x16(256个线程)尽管在本例中是任意更改的但这是一种常见的选择。网格是用足够的块创建的这样每个矩阵元素就有一个线程来处理。为简单起见本例假设每个维度中每个网格的线程数可以被该维度中每个块的线程数整除尽管事实并非如此。
程块需要独立执行必须可以以任何顺序执行它们并行或串行。这种独立性要求允许跨任意数量的内核以任意顺序调度线程块如下图所示使程序员能够编写随内核数量扩展的代码。 块内的线程可以通过一些共享内存共享数据并通过同步它们的执行来协调内存访问来进行协作。 更准确地说可以通过调用 __syncthreads()内部函数来指定内核中的同步点 __syncthreads() 充当屏障块中的所有线程必须等待然后才能继续。 Shared Memory 给出了一个使用共享内存的例子。 除了__syncthreads()之外Cooperative Groups API 还提供了一组丰富的线程同步示例。
为了高效协作共享内存是每个处理器内核附近的低延迟内存很像 L1 缓存并且 __syncthreads() 是轻量级的。
2.3 存储单元层次
CUDA 线程可以在执行期间从多个内存空间访问数据如下图所示。每个线程都有私有的本地内存。每个线程块都具有对该块的所有线程可见的共享内存并且具有与该块相同的生命周期。 所有线程都可以访问相同的全局内存。 还有两个额外的只读内存空间可供所有线程访问常量和纹理内存空间。全局、常量和纹理内存空间针对不同的内存使用进行了优化请参阅设备内存访问。纹理内存还为某些特定数据格式提供不同的寻址模式以及数据过滤请参阅纹理和表面内存。
全局、常量和纹理内存空间在同一应用程序的内核启动中是持久的。
2.4 异构编程
如下图所示CUDA 编程模型假定 CUDA 线程在物理独立的设备上执行该设备作为运行 C 程序的主机的协处理器运行。例如当内核在 GPU 上执行而 C 程序的其余部分在 CPU 上执行时就是这种情况。 CUDA 编程模型还假设主机(host)和设备(device)都在 DRAM 中维护自己独立的内存空间分别称为主机内存和设备内存。因此程序通过调用 CUDA 运行时在编程接口中描述来管理内核可见的全局、常量和纹理内存空间。这包括设备内存分配和释放以及主机和设备内存之间的数据传输。
统一内存提供托管内存来桥接主机和设备内存空间。托管内存可从系统中的所有 CPU 和 GPU 访问作为具有公共地址空间的单个连贯内存映像。此功能可实现设备内存的超额订阅并且无需在主机和设备上显式镜像数据从而大大简化了移植应用程序的任务。有关统一内存的介绍请参阅统一内存编程。
注:串行代码在主机(host)上执行并行代码在设备(device)上执行。
2.5 异步SIMT编程模型
在 CUDA 编程模型中线程是进行计算或内存操作的最低抽象级别。 从基于 NVIDIA Ampere GPU 架构的设备开始CUDA 编程模型通过异步编程模型为内存操作提供加速。 异步编程模型定义了与 CUDA 线程相关的异步操作的行为。
异步编程模型为 CUDA 线程之间的同步定义了异步屏障的行为。 该模型还解释并定义了如何使用 cuda::memcpy_async 在GPU计算时从全局内存中异步移动数据。
2.5.1 异步操作
异步操作定义为由CUDA线程发起的操作并且与其他线程一样异步执行。在结构良好的程序中一个或多个CUDA线程与异步操作同步。发起异步操作的CUDA线程不需要在同步线程中.
这样的异步线程as-if 线程总是与发起异步操作的 CUDA 线程相关联。异步操作使用同步对象来同步操作的完成。这样的同步对象可以由用户显式管理例如cuda::memcpy_async或在库中隐式管理例如cooperative_groups::memcpy_async。
同步对象可以是 cuda::barrier 或 cuda::pipeline。这些对象在Asynchronous Barrier 和 Asynchronous Data Copies using cuda::pipeline.中进行了详细说明。这些同步对象可以在不同的线程范围内使用。作用域定义了一组线程这些线程可以使用同步对象与异步操作进行同步。下表定义了CUDA C中可用的线程作用域以及可以与每个线程同步的线程。
Thread ScopeDescriptioncuda::thread_scope::thread_scope_threadOnly the CUDA thread which initiatedasynchronous operations synchronizes.cuda::thread_scope::thread_scope_blockAll or any CUDA threads within thesame thread block as the initiating thread synchronizes.cuda::thread_scope::thread_scope_deviceAll or any CUDA threads in the sameGPU device as the initiating thread synchronizes.cuda::thread_scope::thread_scope_systemAll or any CUDA or CPU threads in thesame system as the initiating thread synchronizes.
这些线程作用域是在CUDA标准c库中作为标准c的扩展实现的。
2.6 Compute Capability
设备的Compute Capability由版本号表示有时也称其“SM版本”。该版本号标识GPU硬件支持的特性并由应用程序在运行时使用以确定当前GPU上可用的硬件特性和指令。
Compute Capability包括一个主要版本号X和一个次要版本号Y用X.Y表示
主版本号相同的设备具有相同的核心架构。设备的主要修订号是8为NVIDIA Ampere GPU的体系结构的基础上,7基于Volta设备架构,6设备基于Pascal架构,5设备基于Maxwell架构,3基于Kepler架构的设备,2设备基于Fermi架构,1是基于Tesla架构的设备。
次要修订号对应于对核心架构的增量改进可能包括新特性。
Turing是计算能力7.5的设备架构是基于Volta架构的增量更新。
CUDA-Enabled GPUs 列出了所有支持 CUDA 的设备及其计算能力。Compute Capabilities给出了每个计算能力的技术规格。
注意:特定GPU的计算能力版本不应与CUDA版本(如CUDA 7.5、CUDA 8、CUDA 9)混淆CUDA版本指的是CUDA软件平台的版本。CUDA平台被应用开发人员用来创建运行在许多代GPU架构上的应用程序包括未来尚未发明的GPU架构。尽管CUDA平台的新版本通常会通过支持新的GPU架构的计算能力版本来增加对该架构的本地支持但CUDA平台的新版本通常也会包含软件功能。
从CUDA 7.0和CUDA 9.0开始不再支持Tesla和Fermi架构。
