在并行归约示例 opencl 中使用 cl_float3

Question

我将 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() 的问题，它给出了整个大小，而不是仅给出的行数。有什么建议么？

Answer 1

正如其他人提到的，出于大小和对齐的目的，float3（和其他type3类型）的行为与float4（和其他类型）相同。type4使用内置vec_step函数也可以看出这一点，该函数返回输入对象类型中的元素数量，但对于type3对象则返回 4。

如果您的主机代码生成一个打包 float3数组 - 每个对象的大小和对齐方式仅为 3 个浮点数 - 那么从 OpenCL 使用它的正确方法是：

• 使用float*参数代替float3*
• 使用加载数据vload3
• 存储数据使用vstore3