主要内容

gpucoder.stencilKernel

(不推荐)创建CUDA模板函数的代码

gpucoder.stencilKernel不推荐。使用stencilfun代替。有关更多信息,请参见兼容性的考虑

描述

例子

B = gucoder。stencilKernel(乐趣,N [M],形状,param1, param2…)应用函数有趣的对每一个(M, N)输入的滑动窗口一个.函数有趣的为每一个(M, N)子矩阵的一个并计算输出的一个元素B.元素的索引对应于(M, N)窗口。

有趣的是用户定义函数的句柄,该函数返回与输入相同类型的标量输出。

C= FUN(X,param1,param2,…)

X(M, N)原始输入的子矩阵一个X是否可以在必要时进行零填充,例如在输入边界处一个X窗口也可以是1-D的。

C是标量值的输出有趣的.的中心元素计算的输出(M, N)数组X和赋值给输出数组的对应元素B

param1, param2是可选参数。传递这些参数,如果有趣的除了输入窗口外,还需要任何其他参数。

窗口(M, N)必须小于或等于的大小一个的形状一个

如果一个是一维行向量,那么窗口必须是(1, N)

如果一个为一维列向量,则窗口必须为[N, 1]

形状确定输出数组的大小B.它可以有三种可能的值:

  • “相同”-返回输出B它的大小和一个

  • “全部”- (default)返回完整的输出。的大小B>的大小一个,即如果一个大小为(x,y)。的大小B = [x +楼(M/2), y +楼(N/2)]

  • “有效”-只返回输出中没有填充零的边缘的那些部分一个.的大小B = [x - floor(M/2), y - floor(N/2)]

输入一个必须是向量或矩阵,其数字类型由有趣的.的班级B是和班级一样的吗一个

代码生成只支持固定大小的输出。Shape和window必须是编译时常量,因为它们决定了输出的大小。

例子

全部折叠

的用法gpucoder.stencilKernel并生成CUDA®使用模板操作对图像进行过滤的内核。

这个例子执行二维图像的均值滤波。在一个文件中,写入入口点函数测验它接受一个图像矩阵一个.创建子函数my_mean计算均值3 x3子矩阵。

函数B = meanImgFilt(A)% # codegenB = gucoder。stencilKernel (@my_mean, [3 3],“相同”);函数= my_mean(一)=投(意思是((:)),类(A));结束结束

控件的测试输入图像meanImgFilt函数。

inImage = im2double(imread)“cameraman.tif”));

使用codegen函数生成CUDA的MEX函数。

codegen配置coder.gpuConfig(墨西哥人)arg游戏{inImage}报告meanImgFilt

GPU Coder创建了三个内核:meanImgFilt_kernel1对于初始化内存,meanImgFilt_kernel2用于优化输入内存结构,和meanImgFilt_kernel3用于均值滤波操作。下面是生成的代码片段。

cudaMalloc (&gpu_B, 524288妳);cudaMalloc (&gpu_A, 524288妳);cudaMalloc (&gpu_expanded, 532512妳);meanImgFilt_kernel1<<>>(gpu_expanded);cudaMemcpy((void *)gpu_A, (void *)&A[0], 524288ULL, cudaMemcpyHostToDevice);meanImgFilt_kernel2<<>>(gpu_A, gpu_expanded);meanImgFilt_kernel3<<>>(gpu_expanded, gpu_B);cudaMemcpy((void *)&B[0], (void *)gpu_B, 524288ULL, cudaMemcpyDeviceToHost);

meanImgFilt_kernel3使用共享内存(__shared__限定符)来提高内存带宽和数据局部性。

限制

  • 对于非常大的输入尺寸,gpucoder.stencilKernel函数可能产生CUDA代码,在数值上与MATLAB不匹配®模拟。在这种情况下,考虑减少输入的大小以产生准确的结果。

版本历史

在R2017b中引入

全部展开

Baidu
map