主要内容

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

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

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

请注意

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

调用__usad4_wrapCUDA设备功能

立体视差示例测量了一对立体图像中左右两个对应点之间的距离。的stereoDisparity_cuda_sample的入口点函数调用__usad4_wrap外部设备功能通过使用coder.ceval函数。

%%改进的立体视差块匹配算法在这个实现中，不是寻找移位的图像，而是映射索引%因此节省内存和一些处理RGBA列主要打包%数据被用作与CUDA intrinsic兼容的输入。卷积%使用可分离过滤器(水平和垂直)执行函数[out_disp] = stereodisity_cuda_sample (img0,img1) code .cinclude(“cuda_intrinsic.h”）;% gpu代码生成pragmacoder.gpu.kernelfun;%%立体视差参数% WIN_RAD是要操作的窗口的半径，min_difference是%最小视差水平继续搜索，max_悬殊是最大%的差异水平，继续搜索。WIN_RAD = 8;min_disparity = -16;max_disparity = 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 = imgHeight;结束左图像列索引的边界夹紧如果Ind_w1 <= 0;结束如果ind_w1 > imgWidth ind_w1 = imgWidth;结束右图像列索引的边框夹紧如果Ind_w2 <= 0;结束如果ind_w2 = imgWidth;结束在此步骤中，执行绝对差的和%通过四个渠道。这段代码是合适的%用于替换SAD intrinsictDiff = 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)));将SAD成本存储到一个矩阵中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;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向量的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)只支持1个SAD SIMD， //所以有4个指令asm("vabsdiff.u32.u32.u32. SIMD ")添加“%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"(A)， "r"(B)， "r"(结果));asm(“vabsdiff.u32.u32.u32。添加“%0，%1”。b2, % 2。b2, % 3;": "= "(结果):"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 = coder.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_iKernel是调用__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;}}

另请参阅

功能

对象

相关的话题

Baidu
map