核心从分散-聚集类型操作- MATLAB和Simulink - MathWorks德国金宝app - 金宝app,下载188bet金宝搏,金宝搏官方网站

来自分散-聚集类型操作的内核

GPU Coder™还支持约简的金宝app概念——对于循环迭代必须是独立的规则来说，这是一个重要的例外。约简变量累积的值依赖于所有的迭代，但与迭代顺序无关。约简变量出现在赋值语句的两边，比如在求和、点积和排序语句中。下面的示例展示了简化变量的典型用法x：

x =.．.；% x的某个初始化为I = 1:n x = x + d(I);结束

的变量x在每次迭代中，要么在进入循环之前获取其值，要么从循环的前一次迭代中获取其值。由于顺序执行中的依赖链，这种串行顺序类型实现不适合并行执行。另一种方法是采用基于二叉树的方法。

在基于树的方法中，您可以在一定次数的传递中并行执行树的每个水平级别。与顺序执行相比，二叉树确实需要更多内存，因为每次传递都需要一个临时值数组作为输出。您获得的性能好处远远超过增加内存使用的成本。GPU Coder通过使用这种基于树的方法创建缩减内核，其中每个线程块减少数组的一部分。并行约简要求在线程块之间交换部分结果数据。在旧的CUDA中^®在设备上，这种数据交换是通过使用共享内存和线程同步来实现的。从Kepler GPU架构开始，CUDA提供了shuffle (shfl)指令和快速的设备存储器原子操作，使缩减更快。GPU编码器创建的缩减内核使用shfl_down减少经线(32线)的指令。然后，每个经线的第一个线程使用原子操作指令更新减少的值。

有关说明的更多信息，请参阅NVIDIA^®文档。

向量和示例

这个例子展示了如何使用GPU Coder创建CUDA缩减类型的内核。假设你想要创建一个向量v然后计算元素的和。你可以用MATLAB实现这个例子^®函数。

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

为内核创建准备vecSum

GPU Coder不需要特殊的pragma来推断缩减内核。在本例中，使用coder.gpu.kernelfunpragma生成CUDA缩减内核。使用修改后的VecSum函数。

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

生成的CUDA代码

当您通过使用GPU Coder应用程序或从命令行生成CUDA代码时，GPU Coder会创建一个执行向量和计算的单个内核。下面是一个片段vecSum_kernel1．

static __global__ __launch_bounds__(512, 1) void 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) + threadadidx .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, 4096ULL, cudaMemcpyHostToDevice);cudamemcpygpu_s ((void *)gpu_s， (void *)&s, 8ULL, cudaMemcpyHostToDevice);VecSum_kernel1 < < < dim3 (1 u, 1 u, 1 u), dim3 (512 u, 1 u, 1 u) > > > (gpu_v gpu_s);cudaMemcpy(&s, gpu_s, 8U, cudaMemcpyDeviceToHost);

请注意

为了获得更好的性能，GPU Coder优先考虑并行内核。如果你的算法在并行循环中包含约简，GPU Coder会将约简推断为常规循环，并为其生成内核。

图形处理器的一维还原操作

您可以使用gpucoder.reduce函数生成CUDA代码，在GPU上执行高效的1-D约简操作。生成的代码使用CUDA shuffle intrinsic来实现约简操作。

例如，找到总和而且马克斯数组的元素一个：

函数s = myReduce(A) s = gucoder。reduce(A， {@mysum， @mymax});结束函数c = mysum(a, b) c = a+b;结束函数c = mymax(a, b) c = max(a,b);结束

对于代码生成，使用gpucoder.reduce函数有以下要求:

输入必须是数字或逻辑数据类型。
通过@handle传递的函数必须是二进制函数，接受两个输入并返回一个输出。输入和输出必须是相同的数据类型。
函数必须是交换律和结合律。