主要内容

来自Element-Wise循环的内核

CUDA最简单的例子®内核创建是用MATLAB编写的®包含按元素进行缩放的数学运算的函数。当基于元素的操作被封装在for循环体中时，可以调用并发CUDA线程来并行计算每个循环迭代。因为CUDA线程的执行没有特定的顺序，并且彼此独立，所以在你的程序中没有迭代是很重要的为-loop依赖于其他迭代的结果。

基于元素的数学例子

这个例子展示了如何从包含元素数学运算的函数中创建CUDA内核。假设你想对矩阵的每个元素进行平方x乘以1 / (i + j),在那里我,我是行索引和列索引。您可以将此示例实现为MATLAB函数。

函数[y] = myFun(x) y = 0(大小(x));为I = 1:size(x,1)为J = 1:size(x,2) y(i, J)=(x(i, J)²)/(i+ J);结束结束结束

为代码生成准备myFun

第一个表述0(大小(A))在myFun函数是初始化结果向量y0。对于CUDA代码生成，预先分配内存y而不会产生将内存初始化为0的开销。将这一行替换为coder.nullcopy(0(大小(y)))．

为了从循环中创建CUDA内核，GPU Coder™提供了另一个pragmacoder.gpu.kernel．指定此内核pragma将覆盖所有并行循环分析。如果不指定任何参数，GPU Coder将根据循环边界和输入大小确定内核边界。它为您提供了一种指定内核启动参数的方法，例如线程而且块大小。但是，只有在知道循环并行化是安全的情况下才使用它。因为myFun示例很简单，不需要指定内核启动参数，您可以利用coder.gpu.kernelfunpragma生成CUDA内核。

经过这些修改，原来的myFun函数适用于代码生成。

函数[y] = myFun(x)% # codegenY = code .nullcopy(零(大小(x)));coder.gpu.kernelfun ();为I = 1:size(x,1)为J = 1:size(x,2) y(i, J)=(x(i, J)²)/(i+ J);结束结束结束

生成的CUDA代码

当您通过使用GPU Coder应用程序或从命令行生成CUDA代码时，GPU Coder会创建一个执行平方和缩放操作的单个内核。下面是myFun_kernel1内核代码。

static __global__ __launch_bounds__(512, 1) void 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, 2097152ULL, cudaMemcpyHostToDevice);myFun_kernel1 < < < dim3 (512 u, 1 u, 1 u), dim3 (512 u, 1 u, 1 u) > > > (gpu_x gpu_y);cudamemcpyy ((void *)y， (void *)gpu_y, 2097152ULL, cudaMemcpyDeviceToHost);

限制

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

另请参阅

||||

相关的话题