我正在尝试实现一个 OpenCL 版本来减少浮点数组。
为了实现这一目标,我采用了在网上找到的以下代码片段:
__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];
}
该内核代码运行良好,但我想通过添加每个工作组的所有部分总和来计算最终总和。
目前,我通过 CPU 使用简单的循环和迭代来完成最终求和的这一步nWorkGroups
.
我还看到了另一个带有原子函数的解决方案,但它似乎是为 int 实现的,而不是为 float 实现的。我认为只有CUDA提供了float的原子函数。
我还看到我可以使用另一个内核代码来执行此求和操作,但我想避免使用此解决方案以保持简单的可读源。也许没有这个解决方案我就不行......
我必须告诉你,我使用 OpenCL 1.2(由clinfo
)在 Radeon HD 7970 Tahiti 3GB 上(我认为我的卡不支持 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;
}
仅仅一些原子操作不应该在整个内核时间上有效。
其中一些get___part
可以使用 Floor 和类似的函数轻松编写。有些需要除以1M。
本文内容由网友自发贡献,版权归原作者所有,本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容,请联系:hwhale#tublm.com(使用前将#替换为@)