在 OpenCL 内核中存储小型常值数组的最佳实践？

Question

我正在编写一个 OpenCL 内核，它使用 5x5 高斯滤波器对图像进行卷积，并且想知道存储滤波器常量的最佳做法是什么。在内核中，32x32 工作组中的每个线程都执行以下操作：

将像素加载到 __local 内存缓冲区，
通过 barrier(CLK_LOCAL_MEM_FENCE)、
然后对其对应的像素进行卷积。

这里是本地图像数据和过滤器的缓冲区：

 __local float4 localRegion[32][32]; // image region w 2 pixel apron
 .... 
 static const float filter[5][5] = { // __constant vs __private ??
    {1/256.0,  4/256.0,  6/256.0,  4/256.0, 1/256.0},
    {4/256.0, 16/256.0, 24/256.0, 16/256.0, 4/256.0},
    {6/256.0, 24/256.0, 36/256.0, 24/256.0, 6/256.0},
    {4/256.0, 16/256.0, 24/256.0, 16/256.0, 4/256.0},
    {1/256.0,  4/256.0,  6/256.0,  4/256.0, 1/256.0}
  };

哪些内存区域可以容纳 filter，哪个最好，在每种情况下如何进行初始化？ Optimally __private 最好，但我不确定你是否可以静态初始化私有数组？ __local 没有意义，除非某些线程负责加载 filter 条目（我认为）？另外，根据 khronos docs Sec 6.5，我不确定 static 和 _private 可以一起使用。

根据 and here，filter 可以存储为 __private，但不清楚如何初始化。

Answer 1

but I am not sure you can statically initialize private array

Opencl 规范说“静态存储-class 说明符只能用于非内核函数、在程序范围内声明的全局变量和函数内的变量在全局或常量地址 space." 中声明。除此之外，编译器（至少是 Amd 的）优化了常量数学运算并与简单的（常量/指令）内存访问进行交换。即使在此之上，当 space 是不够的，私有寄存器溢出到全局内存并且内核开始访问那里。所以当真实数据有时去了其他地方时，静态不能有有意义的描述。

 float filter[5][5] = {  
    {cos(sin(cos(sin(cos(sin(1/256.0f)))))),  4/256.0f,  6/256.0f,  4/256.0f, 1/256.0f},
    {cos(sin(cos(sin(cos(sin(4/256.0f)))))), 16/256.0f, 24/256.0f, 16/256.0f, 4/256.0f},
    {sin(cos(sin(cos(sin(cos(6/256.0f)))))), 24/256.0f, 36/256.0f, 24/256.0f, 6/256.0f},
    {sin(cos(sin(cos(sin(cos(4/256.0f)))))), 16/256.0f, 24/256.0f, 16/256.0f, 4/256.0f},
    {sin(cos(sin(cos(sin(cos(1/256.0f)))))),  4/256.0f,  6/256.0f,  4/256.0f, 1/256.0f}
  };

与

花费相同的时间（r7_240gpu 为 0.78 毫秒）

float filter[5][5] = { 
    {1/256.0f,  4/256.0f,  6/256.0f,  4/256.0f, 1/256.0f},
    {4/256.0f, 16/256.0f, 24/256.0f, 16/256.0f, 4/256.0f},
    {6/256.0f, 24/256.0f, 36/256.0f, 24/256.0f, 6/256.0f},
    {4/256.0f, 16/256.0f, 24/256.0f, 16/256.0f, 4/256.0f},
    {1/256.0f,  4/256.0f,  6/256.0f,  4/256.0f, 1/256.0f}
  };

探查器的 ISA 输出没有任何正弦或余弦函数。在某些内存位置中只写了一些数字。这是没有启用任何优化的情况。

What memory regions can hold filter, which is best

取决于硬件，但通常有不止一种类型：

// defined before kernel
__constant float filter[5][5] = { 
    {1/256.0f,  4/256.0f,  6/256.0f,  4/256.0f, 1/256.0f},
    {4/256.0f, 16/256.0f, 24/256.0f, 16/256.0f, 4/256.0f},
    {6/256.0f, 24/256.0f, 36/256.0f, 24/256.0f, 6/256.0f},
    {4/256.0f, 16/256.0f, 24/256.0f, 16/256.0f, 4/256.0f},
    {1/256.0f,  4/256.0f,  6/256.0f,  4/256.0f, 1/256.0f}
  };

这对 r7_240 gpu 同时执行。请注意，静态索引对于 __constant 内存访问更好（至少在 amd gpu 中）并且对于相同索引访问也不错（组中的所有线程都访问相同的索引，就像在这个例子中一样（嵌套循环） ).使用这些寻址模式，常量内存比全局内存更快，但是当使用不同的索引时，它与全局内存访问没有什么不同（甚至命中缓存）。 "For globally scoped constant arrays, if the size of an array is below 64 kB, it is placed in hardware constant buffers; otherwise, it uses global memory"。（有 Amd-GCN 架构相关，但 Nvidia 和 Intel 可以预期类似的行为）

AMD 的 opencl 规范说 "L1 and L2 are enabled for images and same-indexed constants."（对于 HD5800 系列 gpu）所以你也可以使用 image2d_t 输入获得类似的性能。对于 GCN，L1 和 L2 比常量内存更快。

Nvidia 的 opencl 最佳实践说：“读取靠近的纹理地址的 p 将达到最佳效果表现。纹理内存还设计用于以常量进行流式读取潜伏;也就是说，缓存命中会减少 DRAM 带宽需求，但不会减少读取延迟。在某些寻址情况下，通过图像对象读取设备内存可以是从全局或常量读取设备内存的有利替代方案记忆。 ”并且还说“它们被缓存，如果有 2D 局部性，可能会表现出更高的带宽在纹理提取中。 “（再次image2d_t）

如果其他地方需要私有内存，您甚至可以拆分过滤器，例如：

// defined before kernel
__constant float filter2[3][5] = {  
    {1/256.0f,  4/256.0f,  6/256.0f,  4/256.0f, 1/256.0f},
    {4/256.0f, 16/256.0f, 24/256.0f, 16/256.0f, 4/256.0f},
    {6/256.0f, 24/256.0f, 36/256.0f, 24/256.0f, 6/256.0f},
  };

   // no need to write __private, automatically private in function body
   float filter[2][5] = { 
        {4/256.0f, 16/256.0f, 24/256.0f, 16/256.0f, 4/256.0f},
        {1/256.0f,  4/256.0f,  6/256.0f,  4/256.0f, 1/256.0f}
    };

这与上面两个示例的时间相同（至少 r7_240）。所有示例都是运行 512x512 大小的图像，512x512 工作项和 16x16 本地工作项。

__local doesn't make sense unless some of the threads are responsible for loading the filter entries

Amd-GCN 上的本地内存是常量内存（相同索引）访问速度的 8 倍，但整个 GPU 的容量是其 5-20 倍（但单个计算单元的容量可能更少）。 Nvidia 的 opencl 最佳实践也是如此。但是 HD5800 系列 amd gpu 具有比本地内存更恒定的内存带宽。 GCN 较新，因此本地内存似乎更好，除非它没有足够的 space.

GCN 上的私有寄存器比本地内存快 5-6 倍，每个计算单元的容量是本地内存的 8 倍。因此，在 GCN 上的私有内存上拥有一些东西意味着终极性能，除非资源消耗停止足够的波阵面来启动（减少延迟隐藏）。

Nvidia 也说了类似的话：“通常，访问寄存器每条指令消耗零个额外的时钟周期，但是由于寄存器写后读依赖和寄存器内存，可能会发生延迟银行冲突。 read-after-write 依赖的延迟大约是 24 个周期，但这延迟在至少有 192 个活动线程的多处理器上完全隐藏（即 6 个经线）。 “

还有一些幽灵墙加载到本地内存中：

    Test gpu was r7_240 so it can work with only 16x16 local threads
    so 20x20 area is loaded from global memory.
    o: each work item's target pixel
    -: needed ghost wall because of filter going out of bounds
    x: ghost corner handled by single threads (yes,non optimized)

    xx----------------xx  
    xx----------------xx
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    --oooooooooooooooo--
    xx----------------xx  
    xx----------------xx

此内核用于上层分析：

            __constant float filter2[3][5] = {  
                        {1/256.0f,  4/256.0f,  6/256.0f,  4/256.0f, 1/256.0f},
                        {4/256.0f, 16/256.0f, 24/256.0f, 16/256.0f, 4/256.0f},
                        {6/256.0f, 24/256.0f, 36/256.0f, 24/256.0f, 6/256.0f},
                      };


            __kernel void test1(__global uchar4 *b2,__global uchar4 *b, __global int * p)
            {
                    int j = get_local_id(0);
                    int g = get_group_id(0);
                    int gx=g%32;
                    int gy=g/32;
                    int lx=j%16;
                    int ly=j/16;
                    int x=gx*16+lx;
                    int y=gy*16+ly;
                    if(gx<2 || gx>29 || gy <2 || gy >29)
                    {
                        b2[((y * 512) + x)] = b[((y * 512) + x)];
                        return;
                    }

                    __local uchar4 localRegion[22][22]; 
                    localRegion[lx+2][ly+2]=b[((y * 512) + x)]; // interior

                    if(lx==0) // left edges
                    {   
                        localRegion[1][ly+2]=b[(( (y) * 512) + x-1)]; // x-1 edge
                        localRegion[0][ly+2]=b[(( (y) * 512) + x-2)]; // x-2 edge
                    }
                    if(lx==15) // right edges
                    {   
                        localRegion[18][ly+2]=b[(( (y) * 512) + x+1)]; // x+1 edge
                        localRegion[19][ly+2]=b[(( (y) * 512) + x+2)]; // x+2 edge
                    }

                    if(ly==0) // top edges
                    {   
                        localRegion[lx+2][1]=b[(( (y-1) * 512) + x)]; // y-1 edge
                        localRegion[lx+2][0]=b[(( (y-2) * 512) + x)]; // y-2 edge
                    }

                    if(ly==15) // bot edges
                    {   
                        localRegion[lx+2][18]=b[(( (y+1) * 512) + x)]; // y+1 edge
                        localRegion[lx+2][19]=b[(( (y+2) * 512) + x)]; // y+2 edge
                    }

                    if(lx==0 && ly==0) // upper-left square
                    {
                        localRegion[0][0]=b[(( (y-2) * 512) + x-2)];
                        localRegion[0][1]=b[(( (y-2) * 512) + x-1)];
                        localRegion[1][0]=b[(( (y-1) * 512) + x-2)];
                        localRegion[1][1]=b[(( (y-1) * 512) + x-1)];
                    }
                    if(lx==15 && ly==0) // upper-right square
                    {
                        localRegion[18][0]=b[(( (y-2) * 512) + x+1)];
                        localRegion[18][1]=b[(( (y-1) * 512) + x+1)];
                        localRegion[19][0]=b[(( (y-2) * 512) + x+2)];
                        localRegion[19][1]=b[(( (y-1) * 512) + x+2)];
                    }
                    if(lx==15 && ly==15) // lower-right square
                    {
                        localRegion[18][18]=b[(( (y+1) * 512) + x+1)];
                        localRegion[18][19]=b[(( (y+2) * 512) + x+1)];
                        localRegion[19][18]=b[(( (y+1) * 512) + x+2)];
                        localRegion[19][19]=b[(( (y+2) * 512) + x+2)];
                    }
                    if(lx==0 && ly==15) // lower-left square
                    {
                        localRegion[0][18]=b[(( (y+1) * 512) + x-2)];
                        localRegion[0][19]=b[(( (y+2) * 512) + x-2)];
                        localRegion[1][18]=b[(( (y+1) * 512) + x-1)];
                        localRegion[1][19]=b[(( (y+2) * 512) + x-1)];
                    }

                    barrier(CLK_LOCAL_MEM_FENCE);



                   float filter[2][5] = { 
                        {4/256.0f, 16/256.0f, 24/256.0f, 16/256.0f, 4/256.0f},
                        {1/256.0f,  4/256.0f,  6/256.0f,  4/256.0f, 1/256.0f}
                    };


                    float4 acc=0;
                    for(int row=-2;row<=0;row++)
                        for(int col=-2;col<=2;col++)
                    {
                        uchar4 tmp=localRegion[lx+col+2][ly+row+2];
                        float tmp2=filter2[row+2][col+2];
                        acc+=((float4)(tmp2,tmp2,tmp2,tmp2)*(float4)((int)tmp.s0,(int)tmp.s1,(int)tmp.s2,(int)tmp.s3));
                    }
                    for(int row=1;row<=2;row++)
                        for(int col=-2;col<=2;col++)
                    {
                        uchar4 tmp=localRegion[lx+col+2][ly+row+2];
                        float tmp2=filter[row-1][col+2];
                        acc+=((float4)(tmp2,tmp2,tmp2,tmp2)*(float4)((int)tmp.s0,(int)tmp.s1,(int)tmp.s2,(int)tmp.s3));
                    }
                    b2[((y * 512) + x)] = (uchar4)(acc.x,acc.y,acc.z,244);
            }

图像为 512x512 rgba（每个通道 8 位）。

源图像（但在作为子步骤过滤之前调整为 512x512）：

结果图片：

我参考的文档：

http://www.nvidia.com/content/cudazone/cudabrowser/downloads/papers/nvidia_opencl_bestpracticesguide.pdf

http://developer.amd.com/wordpress/media/2013/07/AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guide-rev-2.7.pdf

编辑：如果你真的需要 __private、__local、__constant 或 __image2d_t 内存用于内核中的其他内容，你可以完全展开过滤器循环，删除过滤器数组，自己将这些数组元素放入展开的指令中（我试过，它将 VGPR 使用率降低到 21，SGPR 使用率降低到 16）

作为参考，完全消除过滤器计算可将执行时间平均减少 0.05 毫秒，而所有其他版本所需的时间相同。

在 OpenCL 内核中存储小型常值数组的最佳实践？

Best practices for storing a small constant-valued array in OpenCL kernel?

gpu

gpgpu

opencl