主要内容

内核从Element-Wise循环

CUDA的简单的例子®从MATLAB内核创建®函数包含scalarized element-wise数学操作。当element-wise操作是封闭在一个循环的身体,并发CUDA线程可以调用并行计算每个循环迭代。因为CUDA线程在任何特定的顺序执行,是相互独立的,至关重要的是,在你没有迭代为循环的结果取决于其他迭代。

Element-Wise数学例子

这个例子显示了如何创建CUDA内核函数包含element-wise数学操作。假设您想要平方矩阵的每个元素x和规模的一个因素1 / (i + j),在那里我,我行和列的索引。您可以实现这个例子作为一个MATLAB函数。

函数[y] = myFun (x) y = 0(大小(x));为i = 1:尺寸(x, 1)为j = 1:尺寸(x, 2) y (i, j) = (x (i, j) ^ 2) / (i + j);结束结束结束

准备myFun代码生成

第一个语句0(大小(A))在myFun功能是初始化结果向量y0。CUDA代码生成,预先分配内存y在不引起零初始化内存的开销。把这条线换成coder.nullcopy(0(大小(y)))。

从循环创建CUDA内核,GPU编码器™提供了另一个编译指示coder.gpu.kernel。指定这个内核编译指示覆盖所有并行循环分析。如果你不指定任何参数,GPU编码器确定内核范围基于循环边界和输入的大小。它为你提供了一种方法来指定内核启动参数等线程和块大小。然而,只有当你使用它知道循环并行化是安全的。因为myFun例子很简单,不需要规范的内核启动参数,您可以使用coder.gpu.kernelfun编译指示生成CUDA内核。

与这些修改,原来的myFun适用于代码生成功能。

函数[y] = myFun (x)% # codegeny = coder.nullcopy(0(大小(x)));coder.gpu.kernelfun ();为i = 1:尺寸(x, 1)为j = 1:尺寸(x, 2) y (i, j) = (x (i, j) ^ 2) / (i + j);结束结束结束

生成的CUDA代码

当你生成CUDA代码通过使用GPU编码器应用程序或从命令行,GPU编码器创建一个内核执行的棱角和缩放操作。下面的代码片段myFun_kernel1内核代码。

静态__global__ __launch_bounds__(512 1)空白myFun_kernel1 (const real_T * x, real_T * y) {…threadId = ((((gridDim。x * gridDim。y* blockIdx.z + gridDim.x * blockIdx.y) + blockIdx.x) * (blockDim.x * blockDim.y * blockDim.z) + threadIdx.z * blockDim.x * blockDim.y) + threadIdx.y * blockDim.x) + threadIdx.x; i = (int32_T)(threadId / 512U); j = (int32_T)(threadId - (uint32_T)i * 512U); if ((!(j <= 512)) && (!(i <= 512))) { y[i + (j << 9)] = x[i + (j << 9)] * x[i + (j << 9)] / ((real_T)(i + j) + 2.0); } }

以下是主要的一个片段myFun函数。在调用之前myFun_kernel1,有一个cudaMemcpy称之为转移矩阵x从主机(x)设备(gpu_x)。包含512个线程的内核有512块/块,与输入向量的大小一致。第二个cudaMemcpy调用拷贝的结果计算回主机。

cudaMemcpy ((void *) gpu_x (void *) x, 2097152妳,cudaMemcpyHostToDevice);myFun_kernel1 < < < dim3 (512 u, 1 u, 1 u), dim3 (512 u, 1 u, 1 u) > > > (gpu_x gpu_y);cudaMemcpy ((void *) y (void *) gpu_y, 2097152妳,cudaMemcpyDeviceToHost);

限制

  • 如果循环边界的无符号数据类型、代码生成器可以添加条件检查确定循环范围是有效的。这些条件检查可能会限制优化软件和执行的引入减少内核可以影响性能。

另请参阅

||||

相关的话题