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 工作得更好。在本地内存上，随机访问或工作项之间存在间隙都不是问题。这就是为什么它建议在全局使用统一 save/load 而在本地使用 random/scatter/gather。

内核中的第二个循环

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

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

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

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

OpenCl简单优化失败

OpenCl simple unsuccessful optimization

python

opencl

pyopencl