主要内容

遗留代码集成

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

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

笔记

的循环将导致代码生成失败塞瓦尔编码员调用无法映射到内核。请参阅GPU编码器文档中的疑难解答主题,以检查阻止创建内核的问题及其建议的解决方法。如果®代码段包含不支持的函数,则必须删除金宝app塞瓦尔编码员这些部门的电话。

塞瓦尔编码员对于GPU编码器

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

对于代码生成,您必须在调用之前指定参数的类型、大小和复杂性数据类型塞瓦尔编码员

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

遗留代码示例

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

%%立体视差块匹配的改进算法%在此实现中,索引被映射,而不是查找移位图像%相应的节省内存和一些处理RGBA列主要打包%数据用作与CUDA内部函数兼容的输入。卷积%使用可分离过滤器执行(水平和垂直)作用[out_disp] = stereoDisparity_cuda_sample(img0,img1)“cuda_intrinsic.h”);%gpu代码生成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=个([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));%比较可用的最低成本和存储%的差异值如果最低成本(kk,ll)>临时成本最低成本(kk,ll)=临时成本;外部成本(kk,ll)=资产负债表(d)+8;结束结束结束结束结束

的定义__usad4_wrap写在外部文件中cuda_intrinsic.h.该文件位于与入口点功能相同的文件夹中。

__设备uuuuuuuuuuusad4(无符号inta,无符号intb,无符号intc=0){无符号int结果;#if(uuuu-CUDA_uuu-ARCH uuuuu>=300)//Kepler(SM 3.x)支持4向量SAD SIMD asm(“vabsdiff4.u32.u32.u32.u32.金宝app添加“%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.add”%0、%1.b0、%2.b0、%3;“:”=r”(结果):“r”(A),“r”(B),“r”(B),“r”);asm(“vabsdiff.u32.u32.u32.u32.add”%0、%1.b1、%2.b1、%3;“:”=r(结果):“r”(结果):“r”(A),“r”(B),“r”(结果));asm(“vabsdiff.u32.u32.u32.add”%0、%1.b2、%2.u32.u32.u32.u32.u32.add),%3、%r),%r),“r:”结果),“r),“r:”(B),“r),“r:”)asm(“vabsdiff.u32.u32.u32.u32.add”“%0、%1.b3、%2.b3、%3;”:“=r”(结果):“r”(A),“r”(B),“r”(结果));#endif返回结果;#uu设备uuuuuuuuuuu无符号整数压缩字节(const uint8*inBytes){无符号整数压缩=inBytes[0]|(inBytes[1]<<8)|(inBytes[2]<<16)1243];inBytes[uud4]无符号设备)(const uint8_T*A,const uint8_T*B){无符号整数x=packBytes(A);无符号整数y=packBytes(B);返回{uu usad4(x,y);}

生成库达密码

通过创建代码配置对象生成CUDA代码。通过设置自定义代码属性指定自定义C文件的位置(海关包括)下面是一个示例代码生成脚本,它指向cuda_intrinsic.h文件

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

生成的代码

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_立体视差_cuda_样本_i内核调用__usad4_wrap设备函数。以下是e_立体视差_cuda_样本_i内核代码。

静止的__global____launch_bounds__ (512 1)无效的e_立体视差_cuda_样本_i(常数8\u T*img1、常数8\u T*img0、int32\u T d、int32\u T*diff\u 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]=临时成本;}

另见

||||

相关话题