做网站有必要吗,网站推广营销策划,wordpress 使用jquery,个人型网站开站费用本博客参考官方文档进行介绍#xff0c;全网仅此一家进行中文翻译#xff0c;走过路过不要错过。
官方网址#xff1a;https://docs.nvidia.com/cuda/cuda-c-programming-guide/
本文档分成多个博客进行介绍#xff0c;在本人专栏中含有所有内容#xff1a;
https://bl…本博客参考官方文档进行介绍全网仅此一家进行中文翻译走过路过不要错过。
官方网址https://docs.nvidia.com/cuda/cuda-c-programming-guide/
本文档分成多个博客进行介绍在本人专栏中含有所有内容
https://blog.csdn.net/qq_33345365/category_12610860.html
CUDA 12.4为2024年3月2日发表本专栏开始书写日期2024/4/8当时最新版本4.1
本人会维护一个总版本一个小章节的版本总版本会持续更新小版本会及时的调整错误和不合理的翻译内容大部分使用chatGPT 4翻译部分内容人工调整 开始编辑时间2024/4/8最后编辑时间2024/4/10
5.1 内核Kernels
CUDA C通过允许程序员定义C函数称为内核当被调用时这些函数由N个不同的CUDA线程并行执行N次而不是像常规的C函数那样只执行一次。
内核使用 global 声明说明符定义并且给定内核调用的执行该内核的CUDA线程数使用新的 ...执行配置语法指定。执行内核的每个线程都被赋予一个在内核内通过内置变量可以访问的唯一线程ID。
作为示例以下样本代码使用内置变量 threadIdx将两个大小为N的向量A和B相加并将结果存储到向量C中
__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个线程中的每一个都执行一次成对的加法。
5.2 线程架构
为了方便threadIdx是一个三组件向量因此可以使用一维、二维或三维的线程索引来标识线程形成一个一维、二维或三维的线程块称为线程块。这为在像向量、矩阵或体积等域中的元素之间调用计算提供了一种自然的方式。
线程的索引和线程ID之间的关系很直接对于一维的块它们是相同的对于大小为(Dx, Dy)的二维块索引为(x, y)的线程的线程ID为(x y Dx)对于大小为(Dx, Dy, Dz)的三维块索引为(x, y, z)的线程的线程ID为(x y Dx z Dx Dy)。
下面的代码就是一个例子它将两个NxN大小的矩阵A和B加在一起并将结果存储到矩阵C中
__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个线程。
然而一个内核可以被多个形状相同的线程块执行这样总的线程数量等于每个块的线程数量乘以块的数量。
块被组织成一个一维、二维或三维的线程块网格如图4所示。网格中线程块的数量通常由被处理的数据的大小决定这通常超过了系统中的处理器数量。 图4线程块的网格
在 ...语法中指定的每个块的线程数量和每个网格的块数量可以是int类型或dim3类型。可以像上面的例子那样指定二维的块或网格。
网格内的每个块都可以通过一个一维、二维或三维的唯一索引来标识这个索引在内核中可以通过内置的blockIdx变量访问。线程块的维度在内核中可以通过内置的blockDim变量访问。
下面的代码是将前面的MatAdd()示例扩展为处理多个块的情况
// 内核定义
__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);...
}线程块的大小为16x16256个线程尽管在这种情况下是任意的但这是一个常见的选择。如前所述创建足够的块每个矩阵元素有一个线程。为了简单起见这个例子假设每个维度的网格线程数可以被该维度的块线程数整除尽管并非一定是这种情况。
线程块需要独立执行它们必须能够按任何顺序执行可以并行也可以串行。这种独立性要求允许线程块按任何顺序在任何数量的核心上调度如图3所示使程序员能写出随核心数量扩展的代码。
块内的线程可以通过共享一些共享内存和同步它们的执行来协调内存访问来合作。更精确地说可以在内核中通过调用__syncthreads()内在函数来指定同步点__syncthreads()充当一个屏障所有的块内线程都必须在这里等待直到允许任何线程继续。共享内存章节给出了一个使用共享内存的例子。除了__syncthreads()协作组API章节还提供了一整套丰富的线程同步原语。
为了有效的合作预计共享内存是接近每个处理器核心的低延迟内存很像L1缓存并且__syncthreads()预计是轻量级的。
5.2.1 线程块集群 Thread Block Clusters 设计线程数共享内存上同步 线程块集群内同步 全局同步 随着NVIDIA计算能力9.0的引入CUDA编程模型引入了一个叫做线程块群集的可选等级层次这些都是由线程块构成的。与线程块中的线程保证在流式多处理器上被并行调度类似线程块群集中的线程块也保证在GPU处理集群GPC上进行并行调度。
与线程块类似群集也以一维、二维或三维的方式组织如图5所示。群集中的线程块数量可以由用户定义CUDA中支持以8个线程块为单位的群集大小作为最大限制。注意在GPU硬件或MIG配置中如果太小以致不能支持8个多处理器那么最大群集大小将相应减小。识别这些较小的配置以及支持线程块群集大小超过8的较大配置是架构特定的并可以使用cudaOccupancyMaxPotentialClusterSize API进行查询。 图5线程块集群的网格 在使用群集支持启动的内核中为了兼容性gridDim变量仍然表示线程块数量的大小。可以通过使用Cluster Group API找到群集中块的等级。 线程块群集可以通过使用__cluster_dims__(X,Y,Z)的编译器时间内核属性或使用CUDA内核启动API cudaLaunchKernelEx在内核中启用。下面的例子展示了如何使用编译器时间内核属性启动一个群集。使用内核属性的群集大小在编译时固定然后可以使用经典的来启动内核。如果内核使用编译时群集大小那么在启动内核时群集大小不能被修改。
∕∕ 内核定义, Compile time cluster size 2 in X-dimension and 1 in Y and Z dimension
__global__ void __cluster_dims__(2, 1, 1) cluster_kernel(float *input, float* output)
{ }
int main()
{float *input, *output;∕∕ Kernel invocation with compile time cluster sizedim3 threadsPerBlock(16, 16);dim3 numBlocks(N ∕ threadsPerBlock.x, N ∕ threadsPerBlock.y);∕∕ The grid dimension is not affected by cluster launch, and is still enumerated∕∕ using number of blocks.∕∕ The grid dimension must be a multiple of cluster size.cluster_kernelnumBlocks, threadsPerBlock(input, output);
}线程块群集大小也可以在运行时设置并通过CUDA内核启动API cudaLaunchKernelEx启动内核。下面的代码示例展示了如何使用可扩展API启动一个群集内核。
∕∕ 内核定义
∕∕ No compile time attribute attached to the kernel
__global__ void cluster_kernel(float *input, float* output)
{ }
int main()
{float *input, *output;dim3 threadsPerBlock(16, 16);dim3 numBlocks(N ∕ threadsPerBlock.x, N ∕ threadsPerBlock.y);∕∕ Kernel invocation with runtime cluster size{cudaLaunchConfig_t config {0};∕∕ The grid dimension is not affected by cluster launch, and is still enumerated∕∕ using number of blocks.∕∕ The grid dimension should be a multiple of cluster size.config.gridDim numBlocks;config.blockDim threadsPerBlock;cudaLaunchAttribute attribute[1];attribute[0].id cudaLaunchAttributeClusterDimension;attribute[0].val.clusterDim.x 2; ∕∕ Cluster size in X-dimensionattribute[0].val.clusterDim.y 1;attribute[0].val.clusterDim.z 1;config.attrs attribute;config.numAttrs 1;cudaLaunchKernelEx(config, cluster_kernel, input, output);}
}在具有9.0计算能力的GPU中群集中的所有线程块都保证在单个GPU处理集群GPC上进行协同调度并允许群集中的线程块使用Cluster Group API cluster.sync()执行硬件支持的同步。群集组还提供成员函数使用num_threads()和num_blocks()API分别查询群集组大小以线程数量或块数量表示。群集组中的线程或块的排名可以使用dim_threads()和dim_blocks()API分别进行查询。
属于群集的线程块可以访问分布式共享内存。群集中的线程块具有读取、写入和执行分布式共享内存中任何地址的原子操作的能力。分布式共享内存章节给出了在分布式共享内存中执行直方图的示例。
