1. 云栖社区>
  2. 技术文集>
  3. 列表>
  4. 正文

OpenCL性能优化实例研究系列2:避免Local Memory Bank Conflicts的两个简单方法

作者:用户 来源:互联网 时间:2018-05-27 09:02:04

opencl

OpenCL性能优化实例研究系列2:避免Local Memory Bank Conflicts的两个简单方法 - 摘要: 本文讲的是OpenCL性能优化实例研究系列2:避免Local Memory Bank Conflicts的两个简单方法, 转自:http://hi.baidu.com/fsword73/item/51df1fafe6083e268919d39e 作者:  fsword73 Bank Conflicts 是存

转自:http://hi.baidu.com/fsword73/item/51df1fafe6083e268919d39e

作者:  fsword73

Bank Conflicts 是存储访问中的常见问题,避免Bank Conflicts有效地提高存储访问速度。下面介绍两个实例, Reduction和Prefix Sum.

1 在Reduction中使用Padding避免Bank Conflicts

    以AMD HD Readon 5870为例,Local Memory 有32Banks, 每个WAVEFronts有64threads, Bank Conflicts的计算公式为

        bank conflicts = STRIDE * 64/ 32 -2 (当STRIDE为偶数个DWORD)   

         bank conflicts = 0                            (当STRIDE为奇数个DWORD)

          STRIDE = 2,   bank conflicts = 2

          STRIDE = 4,   bank conflicts = 6

         STRIDE = 8,   bank conflicts = 14

         STRIDE = 8,   bank conflicts = 14

         STRIDE = 10,   bank conflicts = 18

          STRIDE = 12,   bank conflicts = 22

          Ruduction的代码为UINT4, 因为它的Bank Conflicts = 6  

          原始代码:

__kernel
void 
reduce(__global uint4* input, __global uint4* output, __local uint4* sdata)
{
    // load shared mem
    unsigned int tid = get_local_id(0);
    unsigned int bid = get_group_id(0);
    unsigned int gid = get_global_id(0);

    unsigned int localSize = get_local_size(0);
    sdata[tid] = input[gid];
    barrier(CLK_LOCAL_MEM_FENCE);

    // do reduction in shared mem
    for(unsigned int s = localSize / 2; s > 0; s >>= 1) 
    {
        if(tid < s) 
        {
            sdata[tid] += sdata[tid + s];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    // write result for this block to global mem
    if(tid == 0) output[bid] = sdata[0];
}

优化的代码, 我们必须使用 __attribute__((packed))来定义数据结构来实现5个DWORD的宽度,否则数据的长度是8个DWORDs。

typedef struct __attribute__((packed))
{
uint4 d;
uint r;
}myData;

__kernel
void 
reduce(__global uint4* input, __global uint4* output, __local mydata* sdata)
{
    // load shared mem
    unsigned int tid = get_local_id(0);
    unsigned int bid = get_group_id(0);
    unsigned int gid = get_global_id(0);

    unsigned int localSize = get_local_size(0);
    sdata[tid].d = input[gid];
    barrier(CLK_LOCAL_MEM_FENCE);

    // do reduction in shared mem
    for(unsigned int s = localSize / 2; s > 0; s >>= 1) 
    {
        if(tid < s) 
        {
            sdata[tid].d += sdata[tid + s].d;
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    // write result for this block to global mem
    if(tid == 0) output[bid] = sdata[0].d;
}

2 在Prefix Sum中避免Bank Conflicts

    我们的Intern刘远浩同学(同济大学研究生)使用了一个非常简单的办法来避免Preix Sum的Bank Conflicts.

    #define HD5870_BANKS 32

      #define AVOID_BACNK_CONFLICTS(X) (x + x / HD5870_BANKS)

   original:

      block[2*tid] = input[2*tid];
   Optimized:

     block[AVOID_BACNK_CONFLICTS(2*tid)] = input[2*tid];

    我们来分析Prefix Sum的执行效率

int offset=1;

for(int d = length>>1; d > 0; d >>=1)
{
   barrier(CLK_LOCAL_MEM_FENCE);
  
   if(tid<d)
   {
    int ai = offset*(2*tid + 1) - 1;
    int bi = offset*(2*tid + 2) - 1;
   
    block[bi] += block[ai];
   }
   offset *= 2;
}

如果Length = 512, 静态分析执行效率

    d = 256 : 256 threads = 4 WaveFronts

   d = 128:   128 threads= 2 WaveFronts     

   d = 64 :      64 threads = 1 WaveFronts

   d = 32:       32 threads = 1patial WaveFronts (50% SIMD utlize)

   d = 16:        16 threads = 1patial WaveFronts (25% SIMD utlize)

   d = 8:        8 threads = 1patial WaveFronts (12.5% SIMD utlize)

   d = 4:        4 threads = 1patial WaveFronts (6.25% SIMD utlize)

   d = 2:        2 threads = 1patial WaveFronts (3.125% SIMD utlize)

   d = 1:        1 threads = 1patial WaveFronts (1.5625% SIMD utlize)

   执行效率:

    511 实际计算 Threads / 总共13 实际 WaveFronts = 61.4 %,   所以Prefix最大的Bottleneck是如何提高ALU和Local Memory模块的实际利用效率。

以上是云栖社区小编为您精心准备的的内容,在云栖社区的博客、问答、公众号、人物、课程等栏目也有 的相关内容,欢迎继续使用右上角搜索按钮进行搜索opencl ,以便于您获取更多的相关知识。

opencl-GPU与CPU同时并行问题

...。但感觉还是不够快。最近想把显卡用上。但研究了一下OpenCL和CUDA,现在只能在仅显卡上做。我有两颗E5,如果只用显卡感觉挺浪费。能不能让CPU和GPU同时计算不同的样本。曾经在论坛上搜到在OpenCL的上下文设置一下就成,但找...

linux查看主板内存槽与内存信息的命令dmidecode方法_Linux

...数、那个槽位插了内存,大小是多少 dmidecode|grep -P -A5 "Memory\s+Device"|grep Size|grep -v Range 2、查看最大支持内存数 dmidecode|grep -P 'Maximum\s+Capacity' 3、查看槽位上内存的速率,没插就是unknown。 dmidecode|grep -A16 "Memory Device"|grep 'Spee...

深入研究Servlet多线程安全性问题

摘 要:介绍了Servlet多线程机制,通过一个实例并结合Java 的内存模型说明引起Servlet线程不安全的原因,给出了保证Servlet线程安全的三种解决方案,并说明三种方案在实际开发中的取舍。  关键字:Servlet 线程安全 同步 Java内...

C++ 资源大全

...档的读权限。并发性并发执行和多线程Boost.Compute :用于OpenCL的C++GPU计算库Bolt :针对GPU进行优化的C++模板库C++React :用于C++11的反应性编程库Intel TBB :Intel线程构件块Libclsph:基于OpenCL的GPU加速SPH流体仿真库OpenCL :并行编程的异...

zeppelin入门使用

...S来对即将启动的Spark driver进行配置,例如“-Dspark.executor.memory=6g -Dspark.cores.max=32”。 尊重原创,拒绝转载,http://blog.csdn.net/stark_summer/article/details/48318059

前三篇
后三篇