从生成的代码调用自定义CUDA设备函数- MATLAB & Simulink - MathWorks España万博1manbetx

从生成的代码调用自定义CUDA设备函数

如果你有高度优化的CUDA^®某些子函数的代码，您想合并到您的生成代码，GPU Coder™扩展coder.ceval帮助您实现这一目标的功能。

外部CUDA函数必须使用__device__在GPU设备上执行函数。这些设备函数与全局函数(内核)不同，因为它们只能从其他设备或全局函数调用。因此,coder.ceval对设备函数的调用必须从映射到内核的循环中进行。有关将CUDA内核与生成的代码集成的信息，请参见从生成的代码调用自定义CUDA内核．

请注意

类型的循环将导致代码生成失败coder.ceval调用不能映射到内核。请参阅GPU Coder文档中的故障排除主题，以检查阻止内核创建的问题及其建议的解决方案。如果你用MATLAB^®代码段包含不受支持的函数，则必须删除万博1manbetxcoder.ceval来自这些部分的调用。

调用`__usad4_wrap`CUDA设备功能

立体视差示例测量立体对的左右图像中两个对应点之间的距离。的stereoDisparity_cuda_sample入口点函数调用__usad4_wrap外接装置功能采用coder.ceval函数。

改进的立体视差块匹配算法在这个实现中，索引被映射，而不是寻找移位的图像%相应地节省内存和一些处理RGBA列主要打包% data被用作与CUDA intrinsic兼容的输入。卷积%使用可分离的过滤器执行(水平和垂直)函数[out_disp] = stereoDisparity_cuda_sample(img0,img1) code .cinclude(“cuda_intrinsic.h”）;% gpu代码生成pragmacoder.gpu.kernelfun;%%立体视差参数WIN_RAD是要操作的窗口的半径，min_悬殊是窗口的半径%的最小视差水平继续搜索，max_悬殊是最大值%的视差水平，继续搜索。Win_rad = 8;min_视差= -16;max_视差= 0;用于循环控制的图像尺寸打包的通道数为4 (RGBA)，因此nChannels为4[imgHeight, imgWidth] =大小(img0);nChannels = 4;imgHeight = imgHeight/nChannels;用于存储原始差异diff_img = 0 ([imgHeight+2*WIN_RAD,imgWidth+2*WIN_RAD]，“int32”）;用于存储最小开销min_cost = 0 ([imgHeight,imgWidth]，“int32”）;Min_cost (:，:) = 99999999;存储最终的视差out_disp = 0 ([imgHeight,imgWidth]，“int16”）;用于聚合差异的过滤器% filter_h是用于可分离卷积的水平滤波器% filter_v是用于可分离卷积的垂直滤波器%操作行卷积的输出Filt_h = ones([1 17]，“int32”）;Filt_v = ones([17 1]，“int32”）;为所有视差级别运行的主循环。这个循环目前是%预计在CPU上运行。为d = min_disparity: max_disparity找出当前视差水平的差值矩阵。预计生成一个内核函数。coder.gpu.kernel;为colIdx = 1: imgWidth + 2 * WIN_RAD coder.gpu.kernel;为rowIdx = 1: imgHeight + 2 * WIN_RAD行索引计算ind_h = rowIdx - WIN_RAD;%左图像的列索引计算ind_w1 = colIdx - WIN_RAD;%右侧图像的行索引计算ind_w2 = colIdx + d - WIN_RAD;%行索引的边框夹持如果Ind_h <= 0 Ind_h = 1;结束如果ind_h > imgHeight ind_h = imgHeight;结束%左图像列索引的边界夹紧如果Ind_w1 <= 0 Ind_w1 = 1;结束如果ind_w1 > imgWidth ind_w1 = imgWidth;结束%右侧图像的列索引的边界夹紧如果Ind_w2 <= 0 Ind_w2 = 1;结束如果ind_w2 > imgWidth ind_w2 = imgWidth;结束在这一步中，执行绝对差的和%跨越四个渠道。这段代码是合适的%替换为SAD intrinsictDiff = int32(0);tDiff = code .ceval(“-gpudevicefcn”，“__usad4_wrap”, coder.rref (img0 ((ind_h-1) * (nChannels) + 1, ind_w1)), coder.rref (img1 ((ind_h-1) * (nChannels) + 1, ind_w2)));将SAD成本存储到一个矩阵中diff_img(rowIdx,colIdx) = tDiff;结束结束使用可分离卷积聚合差异。预计这个使用共享内存生成两个内核。第一个内核是%与水平核和第二核的卷积运算%其输出为列级卷积。Cost_v = conv2(diff_img,filt_h，“有效”）;成本= conv2(cost_v,filt_v，“有效”）;此部分通过比较值更新min_cost矩阵%与当前视差水平。期望为此生成一个Kernel。为我= 1:imgWidth为kk = 1: imgHeight装载成本%Temp_cost = int32(cost(kk,ll));%与可用的最低成本进行比较，并存储视差值%如果Min_cost (kk,ll) > temp_cost (kk,ll) = temp_cost;Out_disp (kk,ll) = abs(d) + 8;结束结束结束结束结束

的定义__usad4_wrap是否写入外部文件cuda_intrinsic.h．该文件与入口点函数位于同一文件夹中。

__device__ unsigned int __usad4(unsigned int A, unsigned int B, unsigned int C=0) {unsigned int result;#if (__CUDA_ARCH__ >= 300) //开普勒(SM 3.x)支持一个4向万博1manbetx量SAD SIMD asm("vabsdiff4.u32.u32.u32. asm。添加“%0，%1，%2，%3”;": "=r"(结果):"r"(A)， "r"(B)， "r"(C));#else // SM 2.0 // Fermi (SM 2.x)只万博1manbetx支持1个SAD SIMD， //所以有4个指令asm("vabsdiff.u32.u32.u32. sm2 .x ")添加" " %0，%1。b0, % 2。b0, % 3;": "=r"(结果):"r"(A)， "r"(B)， "r"(C));asm(“vabsdiff.u32.u32.u32。添加" " %0，%1。b1, % 2。b1, % 3;": "=r"(结果):"r"(A)， "r"(B)， "r"(结果));asm(“vabsdiff.u32.u32.u32。添加" " %0，%1。b2, % 2。b2, % 3;": "=r"(结果):"r"(A)， "r"(B)， "r"(结果));asm(“vabsdiff.u32.u32.u32。添加" " %0，%1。b3, % 2。b3, % 3;": "=r"(result):"r"(A), "r"(B), "r"(result)); #endif return result; } __device__ unsigned int packBytes(const uint8_T *inBytes) { unsigned int packed = inBytes[0] | (inBytes[1] << 8) | (inBytes[2] << 16) | (inBytes[3] << 24); return packed; } __device__ unsigned int __usad4_wrap(const uint8_T *A, const uint8_T *B) { unsigned int x = packBytes(A); unsigned int y = packBytes(B); return __usad4(x, y); }

生成CUDA代码

通过创建代码配置对象生成CUDA代码。通过设置自定义代码属性(CustomInclude)配置对象。的位置，下面是一个示例代码生成脚本cuda_intrinsic.h文件。

cfg = code . gpuconfig (墨西哥人的）;cfg。CustomInclude = pwd;codegen配置cfgarg游戏{imgRGB0, imgRGB1}stereoDisparity_cuda_sample_intrinsic；

生成的代码

GPU Coder创建了四个内核。下面是生成的CUDA代码片段。

e_stereoDisparity_cuda_sample_i<<>> (gpu_img1, gpu_img0, d, gpu_diff_img);*/ /*使用可分离卷积聚合差异。*/ /*期望使用共享内存生成两个内核。* //* The first kernel is the convolution with the horizontal kernel and*/ /* second kernel operates on its output the column wise convolution. */ f_stereoDisparity_cuda_sample_i<<>> (gpu_diff_img, gpu_a); g_stereoDisparity_cuda_sample_i<<>> (gpu_a, gpu_cost_v); h_stereoDisparity_cuda_sample_i<<>> (gpu_a, gpu_cost_v); /* This part updates the min_cost matrix with by comparing the values */ /* with current disparity level. Expect to generate a Kernel for this. */ i_stereoDisparity_cuda_sample_i<<>> (d, gpu_cost, gpu_out_disp, gpu_min_cost);

的e_stereoDisparity_cuda_sample_i内核是调用__usad4_wrap设备的功能。下面是一个片段e_stereoDisparity_cuda_sample_i内核代码。

静态__global____launch_bounds__ (512 1)无效e_stereoDisparity_cuda_sample_i(const uint8_T *img1, const uint8_T *img0, int32_T d, int32_T *diff_img) {.．./*在这一步中，跨四个通道执行绝对差值之和*/ /*。这段代码适合*/ /*为替代悲伤的intrinsic* /temp_cost = __usad4_wrap (&img0 (((ind_h - 1) < < 2) + 2132 * (ind_w1 - 1)), &img1 (((ind_h - 1) < < 2) + 2132 * (temp_cost - 1)));/*存储悲伤的成本成一个矩阵* /diff_img[rowIdx + 549 * colIdx] = temp_cost;}}

另请参阅

功能

codegen|coder.wref|coder.ceval|coder.rref|coder.ref|coder.cinclude

对象

coder.gpuConfig|编码器。代码Config|编码器。EmbeddedCodeConfig|编码器。MexCodeConfig