首页
学习
活动
专区
圈层
工具
发布
首页
学习
活动
专区
圈层
工具
社区首页 >问答首页 >OpenCl简单不成功优化

OpenCl简单不成功优化
EN

Stack Overflow用户
提问于 2017-09-29 05:01:00
回答 1查看 120关注 0票数 0

我正在学习在python中使用opencl,我想优化其中一个函数。我了解到,这可以通过在本地内存中存储全局内存来完成。但是,它不能正常工作,它的持续时间是原来的两倍。干得好吗?我可以对这段代码进行更多的优化吗?

代码语言:javascript
代码运行次数:0
运行
AI代码解释
复制
  __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调用

代码语言:javascript
代码运行次数:0
运行
AI代码解释
复制
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)

谢谢你的建议

EN

回答 1

Stack Overflow用户

回答已采纳

发布于 2017-09-29 05:48:34

除了通过组的所有工作项(如Dithermaster所说)写入相同共享内存地址cache[x]的数据写入争用条件和缺少屏障()函数之外,还可以在修复后添加一些优化:

内核中的第一个循环

代码语言:javascript
代码运行次数:0
运行
AI代码解释
复制
        for ( int x = 0; x < in_layer_s; x++ )
        {
          cache[x] = weights[i*in_layer_s + x];
        }

扫描每个工作项的不同内存区域,每次只扫描一个元素。这在全局内存性能方面可能是错误的,因为每个工作项在自己的循环中都可能使用相同的内存通道,甚至是相同的内存库,因此,所有工作项都依次访问该通道或银行。如果in_layer_s得到更大的值,特别是当它是2的幂时,情况就更糟了。为了解决这个问题,所有的工作项都应该访问相邻的连续地址。当全局内存与工作项一致访问时,GPU工作得更好。在本地内存中,随机访问或在工作项之间存在空白的问题较少。这就是为什么它建议在全局上使用统一的保存/加载,而在本地进行随机/分散/聚集。

内核中的第二循环

代码语言:javascript
代码运行次数:0
运行
AI代码解释
复制
        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端数组的结构是更快和更易读的。这也使得只从主机端向gpu上传必要的权重数据(SOA数组)成为可能,从而进一步减少了总延迟(上载+计算+下载)。

您还可以尝试异步本地<->全局传输函数,使每个工作项组的加载和计算重叠,作为最后手段隐藏更多的延迟。copy.html

票数 0
EN
页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持
原文链接:

https://stackoverflow.com/questions/46489375

复制
相关文章

相似问题

领券
问题归档专栏文章快讯文章归档关键词归档开发者手册归档开发者手册 Section 归档
查看详情【社区公告】 技术创作特训营有奖征文