我为一堆浮点数调整了 openCL 的并行缩减示例。现在我想扩展代码以包含 cl_float3。所以我想在 cl_float3 数组中找到最小值。我认为这是内核中从 float 到 float3 的直接扩展。但是当我从内核返回时，我收到了垃圾值。下面是内核：
__kernel void pmin3(__global float3 *src,
__global float3 *gmin,
__local float3 *lmin,
__global float *dbg,
uint nitems,
uint dev)
{
uint count = nitems / get_global_size(0);
uint idx = (dev == 0) ? get_global_id(0) * count
: get_global_id(0);
uint stride = (dev == 0) ? 1 : get_global_size(0);
// Private min for the work-item
float3 pmin = (float3)(pow(2.0,32.0)-1,pow(2.0,32.0)-1,pow(2.0,32.0)-1);
for (int n = 0; n < count; n++, idx += stride) {
pmin.x = min(pmin.x,src[idx].x);
pmin.y = min(pmin.y,src[idx].y);
pmin.z = min(pmin.z,src[idx].z);
}
// Reduce values within the work-group into local memory
barrier(CLK_LOCAL_MEM_FENCE);
if (get_local_id(0) == 0)
lmin[0] = (float3)(pow(2.0,32.0)-1,pow(2.0,32.0)-1,pow(2.0,32.0)-1);
for (int n = 0; n < get_local_size(0); n++) {
barrier(CLK_LOCAL_MEM_FENCE);
if (get_local_id(0) == n) {
lmin[0].x = min(lmin[0].x,pmin.x);
lmin[0].y = min(lmin[0].y,pmin.y);
lmin[0].z = min(lmin[0].z,pmin.z);
}
}
barrier(CLK_LOCAL_MEM_FENCE);
// Write to __global gmin which will contain the work-group minima
if (get_local_id(0) == 0)
gmin[get_group_id(0)] = lmin[0];
// Collect debug information
if (get_global_id(0) == 0) {
dbg[0] = get_num_groups(0);
dbg[1] = get_global_size(0);
dbg[2] = count;
dbg[3] = stride;
}
}
__kernel void min_reduce3( __global float3 *gmin)
{
for (int n = 0; n < get_global_size(0); n++) {
barrier(CLK_GLOBAL_MEM_FENCE);
if (get_global_id(0) == n) {
gmin[0].x = min(gmin[0].x,gmin[n].x);
gmin[0].y = min(gmin[0].y,gmin[n].y);
gmin[0].z = min(gmin[0].z,gmin[n].z);
}
}
barrier(CLK_GLOBAL_MEM_FENCE);
}
我认为这是 get_global_id(0) 和 get_global_size() 的问题，它给出了整个大小，而不是唯一给出的行数。有什么建议么？