做最后减少的策略
我正在尝试实现一个OpenCL版本来减少一个float数组。
为实现这一目标,我在网上找到了以下代码段:
__kernel void sumGPU ( __global const double *input, __global double *partialSums, __local double *localSums) { uint local_id = get_local_id(0); uint group_size = get_local_size(0); // Copy from global memory to local memory localSums[local_id] = input[get_global_id(0)]; // Loop for computing localSums for (uint stride = group_size/2; stride>0; stride /=2) { // Waiting for each 2x2 addition into given workgroup barrier(CLK_LOCAL_MEM_FENCE); // Divide WorkGroup into 2 parts and add elements 2 by 2 // between local_id and local_id + stride if (local_id < stride) localSums[local_id] += localSums[local_id + stride]; } // Write result into partialSums[nWorkGroups] if (local_id == 0) partialSums[get_group_id(0)] = localSums[0]; }
这个内核代码运行良好,但我想通过添加每个工作组的所有部分和来计算最终总和。 目前,我通过简单的循环和迭代nWorkGroups
CPU执行此步骤。
我还看到了另一个带有primefaces函数的解决方案,但似乎是针对int实现的,而不是浮点数。 我认为只有CUDA为float提供primefaces函数。
我还看到我可以执行另一个执行sum操作的内核代码但是我想避免这个解决方案以保持简单的可读源。 也许我离不开这个解决方案……
我必须告诉你,我在Radeon HD 7970 Tahiti 3GB
上使用OpenCL 1.2
(由clinfo
返回)(我认为我的卡不支持OpenCL 2.0)。
更一般地说,我想得到关于使用我的显卡模型和OpenCL 1.2执行最后一次总结的最简单方法的建议。
欢迎任何帮助,谢谢
如果该浮点数的数量级小于exa
标度,那么:
代替
if (local_id == 0) partialSums[get_group_id(0)] = localSums[0];
你可以用
if (local_id == 0) { if(strategy==ATOMIC) { long integer_part=getIntegerPart(localSums[0]); atom_add (&totalSumIntegerPart[0] ,integer_part); long float_part=1000000*getFloatPart(localSums[0]); // 1000000 for saving meaningful 7 digits as integer atom_add (&totalSumFloatPart[0] ,float_part); } }
这将溢出浮动部分所以当你在另一个内核中除以1000000时,它可能有超过1000000的值,所以你得到它的整数部分并将它添加到实际的整数部分:
float value=0; if(strategy==ATOMIC) { float float_part=getFloatPart_(totalSumFloatPart[0]); float integer_part=getIntegerPart_(totalSumFloatPart[0]) + totalSumIntegerPart[0]; value=integer_part+float_part; }
只需几个primefaces操作就不应该在整个内核时间内有效。
其中一些get___part
可以使用地板和类似function轻松编写。 有些人需要除以1M。
对不起以前的代码。 它也有问题。
CLK_GLOBAL_MEM_FENCE仅影响当前工作组。 我很困惑。 = [
如果你想通过GPU减少总和,你应该在clFinish(commandQueue)之后通过NDRangeKernel函数将减少内核排入队列。
Plaese只是采取概念。
__kernel void sumGPU ( __global const double *input, __global double *partialSums, __local double *localSums) { uint local_id = get_local_id(0); uint group_size = get_local_size(0); // Copy from global memory to local memory localSums[local_id] = input[get_global_id(0)]; // Loop for computing localSums for (uint stride = group_size/2; stride>0; stride /=2) { // Waiting for each 2x2 addition into given workgroup barrier(CLK_LOCAL_MEM_FENCE); // Divide WorkGroup into 2 parts and add elements 2 by 2 // between local_id and local_id + stride if (local_id < stride) localSums[local_id] += localSums[local_id + stride]; } // Write result into partialSums[nWorkGroups] if (local_id == 0) partialSums[get_group_id(0)] = localSums[0]; barrier(CLK_GLOBAL_MEM_FENCE); if(get_group_id(0)==0){ if(local_id < get_num_groups(0)){ // 16384 for(int n=0 ; n0;s/=2){ if(local_id < s) localSums[local_id] += localSums[local_id+s]; barrier(CLK_LOCAL_MEM_FENCE); } if(local_id == 0) partialSums[0] = localSums[0]; } } }