OpenCL:将全局内存复制到每个工作组的本地内存
问题描述:
我正在GPU上使用Open CL实现算法。OpenCL:将全局内存复制到每个工作组的本地内存
目前,我发动内核包含128个工作项目在全局内存 .The数据仅供一个工作组的各项工作,项目被多次使用。要采取共享内存的速度优势,我把它抄了到共享内存使用下面的代码。
__kernel void kernel1(__global float2* input,
__global int* bin,
__global float2* DFT,
__local float2* localInput,
__const int N){
size_t itemId = get_local_id(0);
localInput[itemId] = input[itemId];
barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
........................................................
/*Remaining algo here.*/
........................................................
}
上面的代码工作得很好,如果仅存在一个工作group.But如果有多于一个的工作组,假设有两个工作组以相等的数量的项目在他们每个人以上内核只复制第一个工作组共享内存中的前半部分和后一个中的后半部分。
我也尝试了下面的内核:
__kernel void kernel1(__global float2* input,
__global int* bin,
__global float2* DFT,
__local float2* localInput,
__const int N){
size_t itemId = get_local_id(0);
if(itemId == 0){
for(int index = 0;index<N;index++){
localInput[index] = input[index];
}
}
barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
........................................................
/*Remaining algo here.*/
........................................................
}
但上面的代码,因为这会降低性能的条件语句像发散的问题。
可以对代码进行哪些修改,以便整个数组可以有效地复制到每个工作组的共享内存中?
任何建议,非常感谢。
答
根据您正在运行的设备,您很可能会完全忽略本地内存。如果你使用的是桌面GPU,他们使用几乎没有任何缓存,这使得使用本地内存非常重要,但是现在它们有相当的数量。如果你在gpu上的内存中找到相同的部分,它们全都在缓存中(它与共享内存大小一样),它与本地内存一样快(它们是同一块内存,只是分割) 。手动将其拷贝到本地内存可能会额外强加一个小的性能损失
如果你不是在桌面GPU(ARM /等),或者你的要求使这种不切实际的,async_work_group_copy可能是你正在寻找
在一个不相关的说明中,上面的代码只需要做一个屏障(CLK_LOCAL_MEM_FENCE),因为你大概没有修改你的输入