在OpenCL内核中存储小常量数组的最佳做法是什么？

Question

在OpenCL内核中存储小常量数组的最佳做法是什么？

3

我正在编写一个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，哪个最好，并且每种情况下初始化是如何发生的？理想情况下，__private 是最好的选择，但我不确定是否可以静态初始化私有数组？除非某些线程负责加载 filter 条目（我认为），否则 __local 没有意义。此外，根据khronos文档第6.5节，我不确定static和_private是否可以同时使用。

根据这里的答案和这里，filter可以存储为__private，但其初始化方式不清楚。

- wcochran

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- huseyin tugrul buyukisik · Accepted Answer

但我不确定您是否可以静态初始化私有数组。

OpenCL规范指出，“静态存储类说明符只能用于非内核函数、在程序作用域中声明的全局变量和在全局或常数地址空间内声明的函数内部变量。”除此之外，编译器（至少AMD的编译器）会将常量计算优化并与简单的（常量/指令）内存访问交换。即使在这种情况下，当空间不足时，私有寄存器也会溢出到全局内存，并开始从那里访问。因此，当真实数据有时已去其他地方时，静态无法具有有意义的描述。

 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 240 GPU为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输出没有任何正弦或余弦函数。只有一些数字写在一些内存位置中。这是在没有启用任何优化的情况下的条件。

哪些内存区域可以容纳过滤器，哪种最好

这取决于硬件，但通常有多种类型：

// 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上），对于相同索引访问（所有线程在一个组中访问相同的索引，就像在此示例中（在嵌套循环中））也不错。使用这些寻址模式时，常量内存比全局内存更快，但当使用不同的索引时，它与全局内存访问没有区别（甚至会命中缓存）。 "对于全局范围的常量数组，如果数组大小低于64 kB，则放置在硬件常量缓冲区中；否则，它使用全局内存"。（这与Amd-GCN架构相关，但可以从Nvidia和Intel期望类似的行为）

Amd的opencl规范说“L1和L2对于图像和相同索引的常量是启用的。”（适用于HD5800系列gpu），因此您也可以使用image2d_t输入获得类似的性能。对于GCN，L1和L2比常量内存更快。

Nvidia的opencl最佳实践指南说：“读取彼此接近的纹理地址的内核会获得最佳性能。纹理内存也设计用于具有恒定延迟的流式读取；也就是说，缓存命中会减少DRAM带宽需求，但不会减少读取延迟。在某些寻址情况下，通过image对象从设备内存读取可以成为从全局或常量内存读取设备内存的有利替代方法。”并且还说“它们被缓存，如果纹理提取具有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没有意义

在Amd-GCN上，本地内存比常量内存（同索引）访问快8倍，但整个GPU的容量要大5-20倍（但单个计算单元可能较少）。 Nvidia的OpenCL最佳实践也说了同样的事情。但是HD5800系列AMD GPU的常量内存带宽比本地内存要高得多。 GCN更新，因此似乎本地内存更好，除非空间不足。

GCN上的私有寄存器比本地内存快5-6倍，并且每个计算单元的容量是本地内存的8倍。因此，在GCN上使用专用内存意味着最终性能，除非资源消耗足以停止足够的wavefronts启动（降低延迟隐藏）。

Nvidia也表示类似的事情：“通常，访问寄存器每个指令不会消耗额外的时钟周期，但是可能会由于寄存器写入后读取的依赖关系和寄存器内存冲突而导致延迟。读取后写入依赖关系的延迟约为24个时钟周期，但是在至少有192个活动线程（即6个warp）的多处理器上完全隐藏了此延迟。”

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

    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毫秒，而所有其他版本都需要同样更长的时间。