OpenCl简单不成功优化

Question

我正在学习在python中使用opencl，我想优化一个函数。我了解到，这可以通过将全局内存存储在本地内存中来完成。然而，它不应该像它应该那样工作，持续时间是其两倍。这做得好吗？我可以更优化此代码吗？

__kernel void sumOP( __global float *input, 
        __global float *weights, 
        int layer_size, 
        __global float *partialSums,__local float* cache) 

{ 

    private const int i = get_global_id(0); 
    private const int in_layer_s = layer_size; 
    private const int item_id = get_local_id(0);  
    private const int group_id = get_group_id(0); 
    private const int group_count = get_num_groups(0); 

    const int localsize = get_local_size(0); 

      for (int x = 0; x < in_layer_s; x++) 
      { 
       cache[x] = weights[i*in_layer_s + x]; 
      } 


      float total1 = 0; 

      for (int x = 0; x < in_layer_s; x++) 
      { 
       total1 += cache[x] *input[x]; 
      } 
      partialSums[i] = sigmoid(total1); 

} 

Python的通话

l = opencl.LocalMemory(len(inputs)) 
event = program.sumOP(queue, output.shape, np.random.randn(6,).shape, inputs.data, weights.data,np.int32(len(inputs)),output.data,l) 

感谢一些建议

Answer 1

另外一组的所有工作项做写作，以相同的共享内存地址cache[x]数据写竞争条件（如Dithermaster说）和缺少barrier（）函数，可以在修复之后添加一些优化：

内核中的第一个循环

 for (int x = 0; x < in_layer_s; x++) 
     { 
      cache[x] = weights[i*in_layer_s + x]; 
     } 

为每个工作项扫描一个不同的内存区域，一次扫描一个元素。这在全局内存性能方面可能是错误的，因为每个工作项在自己的循环中都可能使用相同的内存通道或甚至是相同的内存组，因此所有工作项都可以串行访问该通道或库。如果in_layer_s获得更大的值，尤其是如果它的权力为2，则情况会更糟糕。要解决此问题，所有工作项都应该与其邻居访问连续的地址。当工作项目统一访问全局内存时，GPU可以更好地工作。在本地存储器上，随机访问或工作项目之间的差距不太成问题。这就是为什么它建议使用统一的保存/加载全局，同时做本地随机/分散/收集。在内核

 for (int x = 0; x < in_layer_s; x++) 
     { 
      total1 += cache[x] *input[x]; 
     } 

第二环路仅使用单个累加器。这是一个依赖链，需要在完成下一个循环之前完成每个循环。至少使用2个临时“总”变量并展开循环。在这里，如果in_layer_s足够小，可以将input数组移动到本地或常量内存中以更快地访问它（由于所有工作项都访问相同的输入数组，所有工作项都会重复）（可能是一半输入到常量内存，另一半是到本地内存增加总带宽）

---

是weights[i*in_layer_s + x];是一个结构数组？如果是的话，你可以通过使其成为数组的结构并完全摆脱第一个循环的优化来实现加速，主机端的代码膨胀会增加，但如果优先级是速度，那么数组的结构在gpu端更快且可读。这也使得从主机端仅将必要的权重数据（SOA阵列）上传到gpu成为可能，进一步减少了总延迟（上传+计算+下载）。

---

您也可以尝试异步本地< - >全局转移功能，使负载和计算重叠的每一个工作项组，隐藏更延迟作为最后的手段。 https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/async_work_group_copy.html