CUDA 3.0 编程模型分析和总结.docxVIP

本章引入了CUDA 本章引入了CUDA 编程模型背后的主要概念，方式是概述它们是怎样使用C 语言表示的。 CUDA3.0 编程模型 更多的关于 更多的关于CUDA C 的描述在 3.2 节。 本章使用的向量相加例子的完整代码和下一个例子可在SDK 中的vectorAdd代码样本中 本章使用的向量相加例子的完整代码和下一个例子可在SDK 中的vectorAdd代码样本中 找到。 找到。 2.1 内核 2.1 内核 CUDA 通过允许程序员定义称为内核的C 函数扩展了C，内核调用时会被N 个 CUDA 线程 CUDA 通过允许程序员定义称为内核的C 函数扩展了C，内核调用时会被N 个 CUDA 线程 执行 执行 N 次(注：这句话要好好理解，其实每个线程只执行了一次)，这和普通的 C 函数只执行 一次不同。 一次不同。 内核使 内核使用 global 声明符定义，使用一种新...执行配置语法指定执行某一指 定内核调用的线程数(参看附录B.13)。每个执行内核的线程拥有一个独一无二的线程ID， 定内核调用的线程数(参看附录B.13)。每个执行内核的线程拥有一个独一无二的线程ID， 可以通过内置的threadIdx 变量在内核中访问。 可以通过内置的threadIdx 变量在内核中访问。 下面的样本代码将两个长度为N 的向量A 和 B 相加，并将结果存入向量C 中。 下面的样本代码将两个长度为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 threads 这里，N 个线程中的每一个执行VecAdd()的一次成对加法。 这里，N 个线程中的每一个执行VecAdd()的一次成对加法。 2.2 线程层次 2.2 线程层次 为简便起见， 为简便起见，threadIdx 是一个有 3 个部件的向量，所以线程可以使用一维，二维，三 维索引标识，形成一维，二维，三维的线程块。这提供了一种自然的方式来调用作用在域内 维索引标识，形成一维，二维，三维的线程块。这提供了一种自然的方式来调用作用在域内 元素上的计算，如向量，矩阵，体元(volume)。 元素上的计算，如向量，矩阵，体元(volume)。 线程索引和线程 线程索引和线程ID 直接相关：对于一维的块，它们相同;对于二维长度为(Dx,Dy)的块， 线程索引为(x,y)的线程ID 是(x+yDx);对于三维长度为(Dx,Dy,Dz)的块，索引为(x,y,z)的 线程索引为(x,y)的线程ID 是(x+yDx);对于三维长度为(Dx,Dy,Dz)的块，索引为(x,y,z)的 线程 线程 ID 为(x+yDx+zDxDy)。 下面的例子代码将两个长度为N*N 的矩阵A 和B 相加，然后将结果写入矩阵C。 下面的例子代码将两个长度为N*N 的矩阵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() { 由于块内的所有线程必须存在于同一个处理器核心中且共享该核心有限的存储器资源， int main() { ... // Kernel invocation with one block of N * N * 1 threads int numBlocks = 1; dim3 threadsPerBlock(N, N); 因此，一个块内的线程数目是有限的。在目前的 因此，一个块内的线程数目是有限的。在目前的 GPU 上，一个线程块可以包含多达 512 个线 程。 程。 然而，一个内核可被多个同样大小的线程块执行，所以总的线程数等于每个块内的线程 然而，一个内核可被多个同样大小的线程块执行，所以总的线程数等于每个块内的线程 数乘以线程块数。 数乘以线程块数。 线程块被组织成一维或二维的线程网格，如图 线程块被组织成一维或二维的线程网格，如图 2-1 所示。一个网格内的线程块数往往由 被处理的数据量而不是系统的处理器数决定，前者往往远超后者。 被处理的数据量而不是系统的处理