OpenCL select/delete 来自大数组的点
OpenCL select/delete points from large array
我有一组 2M+ 点(计划在适当的时候增加到 20M),我正在通过 OpenCL 运行 进行计算。我想删除落在随机三角形几何图形中的所有点。
如何在 OpenCL 内核进程中执行此操作?
我已经可以了:
识别那些落在三角形外的点(内核中poly算法中的简单点)
将它们的坐标传递到全局输出数组。
但是:
一个 openCL 全局输出数组不能是可变的,所以我初始化它以在大小方面匹配点的输入数组
因此,当一个点落在三角形内时,最终输出中会出现 0,0 个点
因此输出数组本身不会导致任何减少。
可以在openCL上下文中删除0,0点吗?
n.b。我在 OpenFrameworks 中编码,因此 C++ 实现链接到 .cl 文件
如果我理解你的问题,你可以这样做:
--> 在你的内核中,你可以识别三角形中的点并且:
if(element[idx]!=(0,0))
output_array[atomic_inc(number_of_elems)] = element[idx];
最后,在主机中 output_array 的第一个 number_of_elems 中,您将拥有
你的内在点。
希望对你有帮助,
最佳
有替代方案,它们的效果更好或更差,具体取决于数据的外观。我放一个在下面。
删除已识别的点也可以通过在每个工作组的单独数组中注册它们来完成 - 您需要使用与 Moises 的回答相同的 atomic_inc(请参阅我关于在工作组级别执行此操作的评论) !!).最终结果是不需要删除的部分的起点和终点列表。然后,您可以通过不同的线程复制数组的一部分。如果您有需要删除的点簇
,则此方法效果较差
只是大多数点落在原子条件内的情况的替代方案:
可以有一个本地计数器和本地原子。然后将该原子合并到全局值,可以使用 atomic_add()
。 将 return "previous" 全局值。因此,您只需将索引复制到该地址及以上地址即可。
这应该是一个明显的加速,因为线程将在本地同步并且在全局只同步一次。全局副本可以并行,因为地址永远不会重叠。
例如:
__kernel mykernel(__global MyType * global_out, __global int * global_count, _global MyType * global_in){
int lid = get_local_id(0);
int lws = get_local_size(0);
int idx = get_global_id(0);
__local int local_count;
__local int global_val;
//I am using a local container, but a local array of pointers to global is possible as well
__local MyType local_out[WG_SIZE]; //Ensure this is higher than your work_group size
if(lid==0){
local_count = 0; global_val = -1;
}
barrier(CLK_LOCAL_MEM_FENCE);
//Classify them
if(global_in[idx] == ....)
local_out[atomic_inc(local_count)] = global_in[idx];
barrier(CLK_LOCAL_MEM_FENCE);
//If not, we are done
if(local_count > 0){
//Only the first local ID does the atomic to global
if(lid == 0)
global_val = atomic_add(global_count,local_count);
//Resync all the local workers here
barrier(CLK_LOCAL_MEM_FENCE);
//Copy all the data
for(int i=0; i<local_count; i+=lws)
global_out[global_val+i] = local_out[i];
}
}
注意:我没有编译它,但或多或少应该可以工作。
我有一组 2M+ 点(计划在适当的时候增加到 20M),我正在通过 OpenCL 运行 进行计算。我想删除落在随机三角形几何图形中的所有点。
如何在 OpenCL 内核进程中执行此操作?
我已经可以了:
识别那些落在三角形外的点(内核中poly算法中的简单点)
将它们的坐标传递到全局输出数组。
但是:
一个 openCL 全局输出数组不能是可变的,所以我初始化它以在大小方面匹配点的输入数组
因此,当一个点落在三角形内时,最终输出中会出现 0,0 个点
因此输出数组本身不会导致任何减少。
可以在openCL上下文中删除0,0点吗?
n.b。我在 OpenFrameworks 中编码,因此 C++ 实现链接到 .cl 文件
如果我理解你的问题,你可以这样做:
--> 在你的内核中,你可以识别三角形中的点并且:
if(element[idx]!=(0,0))
output_array[atomic_inc(number_of_elems)] = element[idx];
最后,在主机中 output_array 的第一个 number_of_elems 中,您将拥有 你的内在点。
希望对你有帮助, 最佳
有替代方案,它们的效果更好或更差,具体取决于数据的外观。我放一个在下面。
删除已识别的点也可以通过在每个工作组的单独数组中注册它们来完成 - 您需要使用与 Moises 的回答相同的 atomic_inc(请参阅我关于在工作组级别执行此操作的评论) !!).最终结果是不需要删除的部分的起点和终点列表。然后,您可以通过不同的线程复制数组的一部分。如果您有需要删除的点簇
,则此方法效果较差只是大多数点落在原子条件内的情况的替代方案:
可以有一个本地计数器和本地原子。然后将该原子合并到全局值,可以使用 atomic_add()
。 将 return "previous" 全局值。因此,您只需将索引复制到该地址及以上地址即可。
这应该是一个明显的加速,因为线程将在本地同步并且在全局只同步一次。全局副本可以并行,因为地址永远不会重叠。
例如:
__kernel mykernel(__global MyType * global_out, __global int * global_count, _global MyType * global_in){
int lid = get_local_id(0);
int lws = get_local_size(0);
int idx = get_global_id(0);
__local int local_count;
__local int global_val;
//I am using a local container, but a local array of pointers to global is possible as well
__local MyType local_out[WG_SIZE]; //Ensure this is higher than your work_group size
if(lid==0){
local_count = 0; global_val = -1;
}
barrier(CLK_LOCAL_MEM_FENCE);
//Classify them
if(global_in[idx] == ....)
local_out[atomic_inc(local_count)] = global_in[idx];
barrier(CLK_LOCAL_MEM_FENCE);
//If not, we are done
if(local_count > 0){
//Only the first local ID does the atomic to global
if(lid == 0)
global_val = atomic_add(global_count,local_count);
//Resync all the local workers here
barrier(CLK_LOCAL_MEM_FENCE);
//Copy all the data
for(int i=0; i<local_count; i+=lws)
global_out[global_val+i] = local_out[i];
}
}
注意:我没有编译它,但或多或少应该可以工作。