主要内容

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

如果你有高度优化的CUDA®代码为某些子功能你想融入你的生成的代码,GPU编码器™扩展了coder.ceval功能来帮助你实现这个目标。

外部CUDA函数必须使用__device__限定符GPU设备上执行的函数。这些设备的功能不同于全局函数(内核),他们只能从其它设备或全局函数。因此,coder.ceval调用设备功能必须在一个循环中被映射到内核。为信息集成CUDA内核生成的代码,看看从生成的代码调用自定义CUDA内核。

请注意

如果循环包含代码生成失败coder.ceval电话不能映射到内核。看到GPU编码器故障排除主题文档检查问题防止内核创建和提出解决方法。如果你的MATLAB®代码部分包含不支持的功能,那么你必须删除金宝appcoder.ceval电话等部分。

调用__usad4_wrapCUDA设备功能

立体差异的例子措施两个对应点之间的距离在左和右的立体图像对。的stereoDisparity_cuda_sample入口点函数调用__usad4_wrap通过使用外部设备功能coder.ceval函数。

% %修改立体视差块匹配算法%在这个实现中,而不是寻找改变形象,指数映射%因此节省内存和处理RGBA列主要的包装%数据作为输入用于兼容CUDA intrinsic。卷积%执行使用分离过滤器(水平和垂直)函数[out_disp] = stereoDisparity_cuda_sample img0, img1 coder.cinclude (“cuda_intrinsic.h”);% gpu代码生成编译指示coder.gpu.kernelfun;% %立体声差距参数% WIN_RAD操作窗口的半径,min_disparity是%最小差距水平继续搜索,max_disparity是最大的%的差距继续搜索。WIN_RAD = 8;min_disparity = -16;max_disparity = 0;% %形象维度for循环控制%频道包装的数量是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 = 1 (17 [1],“int32”);filt_v = ((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;结束%在这一步中,绝对执行的差异%在四个频道。这段代码是合适的%与悲伤intrinsic替换tDiff = int32 (0);tDiff = coder.ceval (“-gpudevicefcn”,“__usad4_wrap”,coder.rref (img0 ((ind_h-1) * (nChannels) + 1, ind_w1)), coder.rref (img1 ((ind_h-1) * (nChannels) + 1, ind_w2)));%悲伤成本存储到一个矩阵diff_img (rowIdx colIdx) = tDiff;结束结束使用分离卷积%聚合的差异。预计这个%生成两个内核使用共享内存。第一个内核%卷积的内核和第二个内核操作水平%其输出列明智的卷积。cost_v = conv2 (diff_img filt_h,“有效”);成本= conv2 (cost_v filt_v,“有效”);%更新这一部分min_cost矩阵,通过比较值%与当前差距水平。希望生成一个内核。为我= 1:imgWidth为kk = 1: imgHeight%负载成本temp_cost = int32(成本(kk, ll));%比较和存储可用的最低成本%的差异值如果min_cost (kk, ll) > temp_cost min_cost (kk, ll) = temp_cost;out_disp (kk, ll) = abs (d) + 8;结束结束结束结束结束

的定义__usad4_wrap写在一个外部文件吗cuda_intrinsic.h。该文件位于同一文件夹入口点函数。

__device__ unsigned int __usad4 (unsigned int, unsigned int B, unsigned int C = 0) {unsigned int结果;#如果(__CUDA_ARCH__ > = 300) / /开普勒(SM 3. x)支持一个金宝app悲伤的SIMD asm (“vabsdiff4.u32.u32.u32 4向量。添加“% 0,% 1,% 2,% 3;“:”= r”(结果):“r”(一个),“r”(B),“r”(C));# / / SM 2.0 / /其他费米(SM 2. x)只支持1悲金宝app伤SIMD, / / asm (“vabsdiff.u32.u32.u32有4指令。添加“% 0,% 1。b0, % 2。b0, % 3;“:”= r”(结果):“r”(一个),“r”(B),“r”(C));asm (“vabsdiff.u32.u32.u32。添加“% 0,% 1。b1, % 2。b1, % 3;“:”= r”(结果):“r”(一个),“r”(B),“r”(结果));asm (“vabsdiff.u32.u32.u32。添加“% 0,% 1。b2, % 2。b2, % 3;“:”= r”(结果):“r”(一个),“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通过创建一个配置对象。指定自定义C文件的位置通过设置自定义代码属性(CustomInclude配置对象。下面是一个示例代码生成脚本,它指向的位置cuda_intrinsic.h文件。

cfg = coder.gpuConfig (墨西哥人的);cfg。CustomInclude = pwd;codegen配置cfgarg游戏{imgRGB0, imgRGB1}stereoDisparity_cuda_sample_intrinsic;

生成的代码

GPU编码器创建四个内核。下面是生成的CUDA代码的一个片段。

e_stereoDisparity_cuda_sample_i < < < dim3 (704 u, 1 u, 1 u), dim3 (512 u, 1 u, 1 u) > > > (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;}}

另请参阅

功能

对象

相关的话题