遗留代码集成- MATLAB和Simulink金宝app - 金宝app,下载188bet金宝搏,金宝搏官方网站

遗留代码集成

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

外部CUDA函数必须使用__device__限定符来在GPU设备上执行该函数。这些设备函数与全局函数(内核)不同，它们只能从其他设备或全局函数调用。因此,coder.ceval对设备函数的调用必须来自一个被映射到内核的循环。

请注意

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

`coder.ceval`为GPU编码器

coder.ceval('-gpudevicefcn', 'devicefun_name',devicefun_arguments)是coder.ceval函数MATLAB编码器™这样你就可以打电话了__device__内核中的函数。“-gpudevicefcn”指示coder.ceval目标函数在GPU设备上。devicefun_name是什么名字__device__功能和devicefun_arguments逗号分隔的输入参数列表是否按照这个顺序devicefun_name需要。

对于代码生成，您必须在调用之前指定参数的类型、大小和复杂性数据类型coder.ceval．

此函数是一个代码生成函数，否则使用会导致错误。

遗留代码示例

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

改进的立体视差块匹配算法在这个实现中，索引被映射，而不是查找移位的图像%相应的节省内存和一些处理RGBA列主要打包%数据被用作与CUDA intrinsics兼容的输入。卷积%使用可分离过滤器执行(水平过滤器和垂直过滤器)函数[out_disp] = stereoDisparity_cuda_sample(img0,img1)“cuda_intrinsic.h”）;图形处理器代码生成pragmacoder.gpu.kernelfun;立体声视差参数% WIN_RAD是要操作的窗口的半径，min_视差是%继续搜索的最小视差级别，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 = 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内在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)));将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 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向金宝app量SAD SIMD asm("vabsdiff4.u32.u32.u32. ")添加“”%0，%1，%2，%3;": "=r"(结果):"r"(A)， "r"(B)， "r"(C));#else // sm2.0 // Fermi (sm2 .x)只支金宝app持1个SAD SIMD， //因此有4个指令asm("vabsdiff.u32.u32.u32. ")添加“”%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代码。通过设置自定义代码属性指定自定义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<<>> (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; / /扣除成本} ｝