如何将映像从全局内存复制到openCL中的本地内存
问题描述:
我有两个输入映像,我从主机传递给内核。 我的图像尺寸是370x427。如何将映像从全局内存复制到openCL中的本地内存
我想让自己熟悉本地内存,所以我将本地图像传递给内核,并尝试将全局图像复制到本地。
我送我的形象1D array.when我试图显示它不work.My全球worksize结果是{width*height}
和我通过null
本地大小clEnqueueNDRangeKernel
假设opencl会选择本地存储适当的大小。
下面是我的内核代码。
请如果有人能给我一个提示。
__kernel void computeSSD(__global unsigned char *imgL,__global unsigned char *imgR,__global unsigned char *result,__global unsigned char *diff,int width,int MAX_DISP,int WIN_SIZE,__local unsigned char *localLeft,__local unsigned char *localRight)
{
int xCord=get_global_id(0);
int yCord=get_local_id(0);
// copy both images to local memory
localRight[yCord]= imgR[yCord];
localLeft[yCord] = imgL[yCord];
barrier(CLK_LOCAL_MEM_FENCE);
// do operation on local images
result[xCord]=localRight[yCord];
//
}
答
如果您使用3x3过滤器过滤图像。每个工作组都需要工作组像素+1边距。
所以你的内核可以是这样的:
__kernel filter(...){
int x_start = get_group_id(0)*get_local_size(0)-1;
int y_start = get_group_id(1)*get_local_size(1)-1;
int x_end = (get_group_id(0)+1)*get_local_size(0)+1;
int y_end = (get_group_id(1)+1)*get_local_size(1)+1;
__local mytype l[18][18]; //just an example for work sizes 16x16!
//Fetch
//Normally a direct operation per work item is preferred, since it is simpler and the driver will pack all the memory accesses together.
//So just follow coalesced access
//Using CL async workgroup copy has a complex sintax
for(int j=y_start+get_local_id(1); j<y_end; j+=get_local_size(1)){
for(int i=x_start+get_local_id(0); i<x_end; i+=get_local_size(0)){
l[j-y_start][i-x_start] = global[j][i];
}
}
barrier(CLK_GLOBAL_MEM);
//Use the memory for your filtering! (remember to substract y_start & x_start from your original indexes)
//....
}
本地内存是不是一个“耶让所有复制到本地,然后用它,所以它会更快。”这是关于复制该组需要的区域,并在本地重新使用这些读取,因此避免从全局(即3x3过滤器,从全局8次冗余读取,将内存需求减少到1/9)进行大量读取。
此外,如果您没有真正重复使用工作项之间的读取,本地内存可能会变慢。如果每个项目只读写一个内存位置,本地内存将使内核变得更慢,而不是更快。
+0
没问题。此外,大多数设备没有足够的本地内存来存放图像的小部分(370x427)。 – Dithermaster
我的图像尺寸是370x427 – Loraloper