主要内容

内核从散集类型操作

GPU编码器™也支持削减的概念—金宝app—一个重要的例外,必须独立循环迭代。减少变量积累的值取决于一起迭代,但独立于迭代顺序。减少变量出现在赋值语句的双方,如求和,点积和排序。下面的例子显示了减少变量的典型用法x:

x =…;% x的一些初始化为i = 1: n x = x + d(我);结束

的变量x在每个迭代中得到它的值在进入循环或上一次迭代的循环。这种串行顺序实现不适合并行执行由于依赖关系链的顺序执行。另一种方法是采用二叉树的方法。

在基于树的方法,你可以并行执行树的每一个水平超过一定数量的经过。顺序执行,二叉树相比,确实需要更多的内存,因为每一次需要一个临时数组值作为输出。你收到的性能优势远远超过内存使用量增加的成本。GPU编码器产生减少内核通过使用这种基于树的方法在每个线程块降低阵列的一部分。减少并行线程之间的数据交换模块需要部分结果。在旧CUDA®这个数据交换设备,是通过使用共享内存和线程同步。从开普勒架构GPU, CUDA提供shuffle (shfl)指令和快速设备内存原子操作,减少得更快。减少GPU编码器创建使用的内核shfl_down指令减少经纱(32线程)的线程。然后,第一个线程每经使用原子操作指令来更新减少价值。

有关指令的更多信息,请参阅NVIDIA®文档。

矢量和例子

这个例子显示了如何创建CUDA减少内核利用GPU编码器类型。假设您希望创建一个向量v并计算其元素的总和。您可以实现这个例子作为一个MATLAB®函数。

函数s = VecSum (v) s = 0;为i = 1:长度(v) s = s + v(我);结束结束

准备vecSum用于创建内核

GPU编码器不需要特殊的编译指示来推断减少内核。在这个例子中,使用coder.gpu.kernelfun编译指示生成CUDA减少内核。使用修改后的VecSum函数。

请注意

使用coder.gpu.kernel编译指示for循环包含削减不推荐。

函数s = VecSum (v)% # codegens = 0;coder.gpu.kernelfun ();为i = 1:长度(v) s = s + v(我);结束结束

生成的CUDA代码

当你生成CUDA代码通过使用GPU编码器应用程序或从命令行,GPU编码器创建一个内核执行向量和计算。下面是一个片段vecSum_kernel1。

静态__global__ __launch_bounds__(512 1)空白vecSum_kernel1 (const real_T * v, real_T * s) {uint32_T threadId;uint32_T threadStride;uint32_T thdBlkId;uint32_T idx;real_T tmpRed;;;thdBlkId = (threadIdx。z * blockDim。x * blockDim。y + threadIdx。y * blockDim.x) + threadIdx.x; threadId = ((gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y) + blockIdx.x) * (blockDim.x * blockDim.y * blockDim.z) + thdBlkId; threadStride = gridDim.x * blockDim.x * (gridDim.y * blockDim.y) * (gridDim.z * blockDim.z); if (!((int32_T)threadId >= 512)) { tmpRed = 0.0; for (idx = threadId; threadStride < 0U ? idx >= 511U : idx <= 511U; idx += threadStride) { tmpRed += v[idx]; } tmpRed = workGroupReduction1(tmpRed, 0.0); if (thdBlkId == 0U) { atomicOp1(s, tmpRed); } } }

在调用之前VecSum_kernel1,两个cudaMemcpy呼叫转移的向量v和标量年代从主机到设备。内核有一个线程块包含512个线程/块,与输入向量的大小一致。第三个cudaMemcpy调用拷贝的结果计算回主机。以下是主要功能的一个片段。

cudaMemcpy ((void *) gpu_v (void *) v, 4096妳,cudaMemcpyHostToDevice);cudaMemcpy ((void *) gpu_s (void *)谨此告知,8妳,cudaMemcpyHostToDevice);VecSum_kernel1 < < < dim3 (1 u, 1 u, 1 u), dim3 (512 u, 1 u, 1 u) > > > (gpu_v gpu_s);cudaMemcpy(郑清奎gpu_s 8 u, cudaMemcpyDeviceToHost);

请注意

为更好的性能,GPU编码器为主平行内核在减少。如果你减少算法包含在一个并行循环,GPU编码器推断减少作为常规循环并生成内核。

一维GPU上还原操作

您可以使用gpucoder.reduce函数来生成CUDA代码执行效率降低一维GPU的操作。生成的代码使用CUDA洗牌intrinsic执行还原操作。

例如,找到总和和马克斯一个数组元素一个:

函数s = myReduce (A) s = gpucoder。减少(A, {@mysum, @mymax});结束函数c = mysum (a, b) c = a + b;结束函数c = mymax (a, b) c = max (a, b);结束
对于代码生成,gpucoder.reduce这些需求函数:

  • 的输入必须是数字或逻辑数据类型。

  • 函数通过@handle必须是一个二元函数,它接受两个输入,并返回一个输出。输入和输出必须是相同的数据类型。

  • 函数必须交换和关联。

请注意

一些输入的整数数据类型,生成的代码gpucoder.reduce函数可以包含中间计算达到饱和。在这种情况下,生成的代码可能不匹配的结果从MATLAB仿真结果。

另请参阅

|||||

相关的话题