OpenCL,C - Pi 的莱布尼茨公式
OpenCL, C - Leibniz Formula for Pi
我正在尝试获得一些使用 OpenCL 的经验,环境已设置,我可以创建和执行内核。我目前正在尝试使用 Leibniz 公式 并行计算圆周率,但收到了一些奇怪的结果。
内核如下:
__kernel void leibniz_cl(__global float *space, __global float *result, int chunk_size)
{
__local float pi[THREADS_PER_WORKGROUP];
pi[get_local_id(0)] = 0.;
for (int i = 0; i < chunk_size; i += THREADS_PER_WORKGROUP) {
// `idx` is the work item's `i` in the grander scheme
int idx = (get_group_id(0) * chunk_size) + get_local_id(0) + i;
float idx_f = 1 / ((2 * (float) idx) + 1);
// Make the fraction negative if needed
if(idx & 1)
idx_f = -idx_f;
pi[get_local_id(0)] += idx_f;
}
// Reduction within workgroups (in `pi[]`)
for(int groupsize = THREADS_PER_WORKGROUP / 2; groupsize > 0; groupsize >>= 1) {
if (get_local_id(0) < groupsize)
pi[get_local_id(0)] += pi[get_local_id(0) + groupsize];
barrier(CLK_LOCAL_MEM_FENCE);
}
如果我在这里结束函数并将 result
设置为 pi[get_local_id(0)]
for !get_global_id(0)
(就像第一组的减少一样),打印 result
打印 -nan
.
内核的剩余部分:
// Reduction amongst workgroups (into `space[]`)
if(!get_local_id(0)) {
space[get_group_id(0)] = pi[get_local_id(0)];
for(int groupsize = get_num_groups(0) / 2; groupsize > 0; groupsize >>= 1) {
if(get_group_id(0) < groupsize)
space[get_group_id(0)] += space[get_group_id(0) + groupsize];
barrier(CLK_LOCAL_MEM_FENCE);
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if(get_global_id(0) == 0)
*result = space[get_group_id(0)] * 4;
}
返回 space[get_group_id(0)] * 4
returns -nan
或一个非常大的数字,显然不是圆周率的近似值。
我无法确定它是我缺少的 OpenCL 概念还是一般的并行执行概念。感谢任何帮助。
链接
缩减模板:OpenCL float sum reduction
莱布尼茨公式:https://www.wikiwand.com/en/Leibniz_formula_for_%CF%80
也许这些不是代码中最关键的问题,但它们可能是问题的根源:
- 您绝对应该在局部缩减之前使用
barrier(CLK_LOCAL_MEM_FENCE);
。如果只有您知道工作组大小等于或小于波前线程数 运行 并行相同指令 - AMD GPU 为 64,NVidia GPU 为 32,则可以避免这种情况。
- 全局缩减必须在多次启动内核时完成,因为
barrier()
仅适用于同一工作组的工作项。将屏障插入内核的清晰且 100% 有效的方法是在需要全局屏障的地方将其一分为二。
我正在尝试获得一些使用 OpenCL 的经验,环境已设置,我可以创建和执行内核。我目前正在尝试使用 Leibniz 公式 并行计算圆周率,但收到了一些奇怪的结果。
内核如下:
__kernel void leibniz_cl(__global float *space, __global float *result, int chunk_size)
{
__local float pi[THREADS_PER_WORKGROUP];
pi[get_local_id(0)] = 0.;
for (int i = 0; i < chunk_size; i += THREADS_PER_WORKGROUP) {
// `idx` is the work item's `i` in the grander scheme
int idx = (get_group_id(0) * chunk_size) + get_local_id(0) + i;
float idx_f = 1 / ((2 * (float) idx) + 1);
// Make the fraction negative if needed
if(idx & 1)
idx_f = -idx_f;
pi[get_local_id(0)] += idx_f;
}
// Reduction within workgroups (in `pi[]`)
for(int groupsize = THREADS_PER_WORKGROUP / 2; groupsize > 0; groupsize >>= 1) {
if (get_local_id(0) < groupsize)
pi[get_local_id(0)] += pi[get_local_id(0) + groupsize];
barrier(CLK_LOCAL_MEM_FENCE);
}
如果我在这里结束函数并将 result
设置为 pi[get_local_id(0)]
for !get_global_id(0)
(就像第一组的减少一样),打印 result
打印 -nan
.
内核的剩余部分:
// Reduction amongst workgroups (into `space[]`)
if(!get_local_id(0)) {
space[get_group_id(0)] = pi[get_local_id(0)];
for(int groupsize = get_num_groups(0) / 2; groupsize > 0; groupsize >>= 1) {
if(get_group_id(0) < groupsize)
space[get_group_id(0)] += space[get_group_id(0) + groupsize];
barrier(CLK_LOCAL_MEM_FENCE);
}
}
barrier(CLK_LOCAL_MEM_FENCE);
if(get_global_id(0) == 0)
*result = space[get_group_id(0)] * 4;
}
返回 space[get_group_id(0)] * 4
returns -nan
或一个非常大的数字,显然不是圆周率的近似值。
我无法确定它是我缺少的 OpenCL 概念还是一般的并行执行概念。感谢任何帮助。
链接
缩减模板:OpenCL float sum reduction
莱布尼茨公式:https://www.wikiwand.com/en/Leibniz_formula_for_%CF%80
也许这些不是代码中最关键的问题,但它们可能是问题的根源:
- 您绝对应该在局部缩减之前使用
barrier(CLK_LOCAL_MEM_FENCE);
。如果只有您知道工作组大小等于或小于波前线程数 运行 并行相同指令 - AMD GPU 为 64,NVidia GPU 为 32,则可以避免这种情况。 - 全局缩减必须在多次启动内核时完成,因为
barrier()
仅适用于同一工作组的工作项。将屏障插入内核的清晰且 100% 有效的方法是在需要全局屏障的地方将其一分为二。