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

Question

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

例如，我们有一个 1000*1000 的图像，我们想将每个像素值加 1。我们可以使用 1000*1000 的全局内存对吗？

如果我们使用本地内存，将 1000*1000 的图像变成 100 个 100*100 的部分，对我们来说会更快吗？

如果你给我一个简单的本地内存代码，我会很感激你的。

Answer 1

Cann't we just use the global memory?

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

Since the global memory can get the item separately and freely

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

We can use 1000*1000's global memory right?

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

Will it be faster for us if we use local memory and turn the 1000*1000 image into 100 100*100 parts?

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

如果您将每个数据重复使用 256 次，那么本地内存的访问速度将比全局内存访问快 10-20 倍左右。

这是一个非常简单的用于力计算的二维 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
}

上例在本地内存重用率方面较差。但是一半的变量在私有(private)内存中，并且被重复使用了 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 左右。您也可以针对私有(private)寄存器进行优化!那么它可能是 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.

使用为您提供缓存线使用信息的分析器。