遗留代码集成- MATLAB和Simulink万博1manbetx - 万博1manbetx,s manbetx 845,万博尤文图斯

遗留代码集成

如果你有高度优化的CUDA^®要合并到生成代码中的特定子功能的代码，GPU编码器™ 扩展coder.ceval帮助您实现此目标的功能。

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

笔记

的循环将导致代码生成失败coder.ceval调用无法映射到内核。请参阅GPU编码器文档中的疑难解答主题，以检查阻止创建内核的问题及其建议的解决方法。如果^®代码段包含不支持的函数，则必须删除万博1manbetxcoder.ceval来自这些部门的电话。

`coder.ceval`对于GPU编码器

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

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

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

遗留代码示例

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

改进的立体视差块匹配算法%在此实现中，索引被映射，而不是查找移位图像%相应的节省内存和一些处理RGBA列主要打包%数据被用作与CUDA intrinsics兼容的输入。卷积%使用可分离过滤器执行（水平和垂直）函数[out_disp] = stereoDisparity_cuda_sample(img0,img1)“cuda_intrinsic.h”）;图形处理器代码生成pragmacoder.gpu.kernelfun;%%立体视差参数%WIN_RAD是要操作的窗口的半径，min_Distance是%继续搜索的最小视差级别，max_视差是最大的继续搜索的视差级别。WIN_RAD = 8;min_disparity = -16;max_disparity = 0;%%用于循环控制的图像尺寸%压缩的通道数为4（RGBA），因此通道数为4[imgHeight, imgWidth] =大小(img0);nChannels = 4;imgHeight = imgHeight / nChannels;%%存储原始差异diff_img = 0 ([imgHeight + 2 * WIN_RAD imgWidth + 2 * WIN_RAD],“int32”）;%储存成本最低最小成本=零（[imgHeight，imgWidth]，“int32”)；最低成本（：，：）=9999999；存储最后的视差out_disp=零（[imgHeight，imgWidth]，“int16”）;%%用于聚合差异的过滤器%filter_h是用于可分离卷积的水平滤波器% filter_v是用于可分离卷积的垂直滤波器%作用于行卷积的输出Filt_h = ones([1 17]，“int32”)；filt_v=个（[17 1]，“int32”）;%%运行于所有差异级别的主循环。此循环当前为%预期在CPU上运行。对于d=最小视差：最大视差找到当前视差水平的差值矩阵。预计生成一个内核函数。coder.gpu.kernel;对于colIdx = 1: imgWidth + 2 * WIN_RAD coder.gpu.kernel;对于rowIdx=1:imgHeight+2*WIN\u 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；结束结束%使用可分离卷积聚合差异。预期如下%使用共享内存生成两个内核。第一个内核是与水平核和第二核进行卷积%它的输出是列卷积。成本=conv2（差异、过滤、，“有效”); 成本=conv2（成本、过滤、，“有效”）;%这部分通过比较值来更新min_cost矩阵%与当前视差水平。期望为此生成一个内核。对于我= 1:imgWidth对于kk = 1: imgHeight%加载成本临时成本=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向万博1manbetx量SAD SIMD asm("vabsdiff4.u32.u32.u32. ")添加“”%0，%1，%2，%3;": "=r"(结果):"r"(A)， "r"(B)， "r"(C));#else // sm2.0 // Fermi (sm2 .x)只支万博1manbetx持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代码。通过设置自定义代码属性指定自定义C文件的位置(CustomInclude)在配置对象上。下面是一个示例代码生成脚本，它指向cuda_intrinsic.h文件。

cfg=coder.gpuConfig(“墨西哥”); cfg.CustomInclude=pwd；编码基因-配置cfg-args{imgRGB0, imgRGB1}stereoDisparity_cuda_sample_intrinsic;

生成的代码

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

e\U立体视差\U cuda\U样本\U i<>（gpu img1，gpu img0，d，gpu diff\U img）*//*使用可分离的卷积来聚合差异。*//*预计这将使用共享内存生成两个内核。*//*第一个内核是与水平内核的卷积，*/*第二个内核对其输出进行逐列卷积运算。*/f_stereodisparsion_cuda_sample_i<>（gpu diff\u img，gpu a）；g\u立体视差图cuda\u样本i<>（gpu a，gpu成本v）；h\u立体视差图cuda\u样本i<>（gpu a，gpu成本v）/*此部分通过将值*/*与当前差异级别进行比较来更新最小成本矩阵。希望为此生成一个内核。*/i_stereodisparsion_cuda_sample_i<>（d，gpu成本，gpu输出，gpu最小成本）；

这个e_立体视差_cuda_样本_i内核调用__usad4_wrap设备函数。以下是e_立体视差_cuda_样本_i内核代码。

静止的__global____launch_bounds__ (512 1)无效e_立体视差_cuda_样本_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)));/ *存储悲伤的成本进入A.矩阵*／diff_img[rowIdx+549*colIdx]=临时成本；}