OpenCL中本地内存的优势是什么？

Question

我想知道本地内存的优势.由于全局内存可以单独和自由地获取项目.我们不能只使用全局内存吗？

例如,我们有一个1000*1000的图像,我们想要添加每个像素值1.我们可以使用1000*1000的全局内存吗？

如果我们使用本地内存并将1000*1000图像转换为100个100*100个部分,我们会更快吗？

如果你给我一个简单的本地记忆代码,我会非常感谢你.

Answer 1

我们不能只使用全球记忆吗？

当然可以.首先写一个实际的工作代码.然后优化.

由于全局内存可以单独和自由地获取项目

我不确定所有架构是否都具有广播能力.但我确定如果所有线程都随机访问内存,它会变得太慢.光线跟踪就是一个例子.每个像素折射/反射到不同的距离和不同的存储区域.这是一个性能打击.如果每个线程都以统一的方式访问全局内存,那么速度会快得多.

我们可以使用1000*1000的全局内存吗？

有一个最小值的最大缓冲区的大小,它可以是围绕128MB或设备存储器的1/4.所有缓冲区的组合大小可能因平台/设备而异,范围为几GB.

如果我们使用本地内存并将1000*1000图像转换为100个100*100个部分,我们会更快吗？

这取决于数据重用率和访问模式的合并性.对本地内存的随机(非合并)访问比对全局内存的随机(非合并)访问快得多.如果你使用太多的本地内存/私有文件,那么它可能会更慢,因为更多的本地内存消耗会导致更少的占用和更少的内存延迟隐藏以及更多的寄存器溢出到全局内存.尝试使用私有寄存器来平衡它.或者,您可以使用压缩技术将更多数据放入本地内存中.

如果你重复使用每个数据256次,那么本地内存比全局内存访问快10-20倍.

这是一个非常简单的用于力计算的2D nbody代码:

// global memory access is only 257 times per item, 1 for private save
//                                                  256 for global broadcast
//                                                  for global-to-local copy
// unoptimized version accesses 65537 times per item.
__kernel void nBodyF(__global float *x, __global float *y,
                     __global float *vx, __global float *vy,
                     __global float *fx, __global float *fy)
{
     int N=65536; // this is total number of masses for this example
     int LN=256;  // this is length of each chunk in local memory, 
                  // means 256 masses per compute unit
    int i=get_global_id(0);  // global thread id keys 0....65535
    int L=get_local_id(0);   // local thread id keys 0...255 for each group
    float2 Fi=(float2)(0,0); // init
    float xi=x[i]; float yi=y[i]; // re-use for 65536 times
    __local xL[256]; __local yL[256]; //declare local mem array with constant length


    for(int k=0;k<N/LN;k++) // number of chunks to fetch from global to local
    {
        barrier(CLK_LOCAL_MEM_FENCE);  //synchronization
        xL[L]=x[k*LN+L]; yL[L]=y[k*LN+L]; //get 256-element chunks into local mem
        barrier(CLK_LOCAL_MEM_FENCE);  //synchronization
        for(int j=0;j<LN;j++)          //start processing local/private variables
        {
            float2 F=(float2)(0,0);          // private force vector init
            float2 r1=(float2)(xi,yi);       // private vector
            float2 r2=(float2)(xL[j],yL[j]); // use local mem to get r2 vector
            float2 dr=r1-r2;                 // private displacement
            F=dr/(0.01f+dot(dr,dr));         // private force calc.
            Fi.x-=F.x; Fi.y-=F.y;            // private force add to private
        }
     }
     fx[i]=Fi.x; fy[i]=Fi.y; //write result to global mem only once
}

上面的示例在本地存储器重用率方面较差.但是有一半的变量存在于私有内存中,并且重复使用了64k次.

最坏的情况:

  1)Big portion of items cannot fit GPU cache.
  2)Only global memory accesses are done
  3)Data is not re-used
  4)Memory is accessed in a very non-uniform way.
  This will make it very slow.
  When data doesnt fit cache and not re-used, you should use __read_only for
  necessary buffers(__write_only for writing).

如果进行卷积(或一些抗锯齿或边缘检测),数据重用将为4到20,本地内存优化至少提供3-4倍的性能.

如果您的GPU具有300GB/s的全局内存带宽,那么其本地内存带宽将约为3-4 TB/s.您也可以优化私人寄存器!那么它可能是15-20 TB/s.但是这种类型的使用区域较少.

编辑:如果您正在读取单个字节,并且如果这些字节之间只有一个较小的值(例如最大值为16),那么您可以将多个变量打包成单个字节并在本地存储器中解密它们.例:

  Global memory(copied to local mem): 
  Reference_byte   Byte0  byte1        byte2         byte3  
  128              +3,-5  +24,+50      -25,-63      0, +2

  Unpacking in local memory:
  Reference_byte   Byte0  byte1 byte2 byte3 Byte4  byte5 byte6 byte7      
  128              131    126   150   200   175    112   112   114

  Computing results on the array
  Reference_byte   Byte0  byte1 byte2 byte3 Byte4  byte5 byte6 byte7 
  128              120    130   140   150   150    150   100   110

  Packing results in local memory:
  Reference_byte   Byte0  byte1        byte2         byte3  
  128              -8,+10 +10,+10      0,0           -50, +10

  Global memory(copied from local mem): 
  Reference_byte   Byte0  byte1        byte2         byte3  
  128              -8,+10 +10,+10      0,0           -50, +10

  //Maybe a coordinate compression for a voxel rendering.

使用分析器为您提供缓存行使用信息.