gpt4 book ai didi

opencl - 在OpenCL内核中存储小型恒定值数组的最佳实践?

转载 作者:行者123 更新时间:2023-12-02 09:51:26 25 4
gpt4 key购买 nike

我正在编写一个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最佳,但是我不确定您可以静态初始化私有数组吗?除非某些线程负责加载 __local条目,否则 filter没有任何意义(我认为)?另外,根据 khronos docs Sec 6.5,我不确定 static_private是否可以一起使用。

根据 the answers hereherefilter可以存储为 __private,但不清楚如何初始化。

最佳答案

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


Opencl规范说:“静态存储类说明符只能用于
非内核函数,在程序范围内声明的全局变量和函数内部的变量
最重要的是,编译器(至少是Amd的)优化了常数数学运算,并通过简单的(常数/指令)内存访问进行了交换。即使在此之上,当空间不足时,专用寄存器溢出到全局内存,内核开始访问全局内存,因此,当真实数据有时移到其他地方时,static无法具有有意义的描述。

 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.78ms)

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中)更好,而对于相同索引访问则不错(组中的所有线程都访问同一索引,就像在此示例中一样(int嵌套循环))。在这些寻址方式下,常量内存比全局内存要快,但是使用不同的索引时,常量内存与全局内存访问(甚至命中缓存)没有什么不同。 “对于全局范围的常量数组,如果数组的大小小于64 kB,则将其放置在硬件常量缓冲区中;否则,它将使用全局内存”。 (存在与Amd-GCN架构相关的问题,但Nvidia和Intel可以预期类似的行为)

Amd的opencl规范说:“为图像和相同索引的常量启用了L1和L2。”(对于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没有任何意义


Amd-GCN上的本地内存的访问速度是恒定内存(相同索引)访问速度的8倍,但整个GPU的容量是后者的5-20倍(但单个计算单元可能会更少)。 Nvidia的opencl最佳做法也是如此。但是HD5800系列AMD GPU具有比本地内存更多的恒定内存带宽。 GCN较新,因此除非没有足够的空间,否则本地内存似乎更好。

GCN上的专用寄存器比本地内存快5-6倍,每个计算单元的容量是本地内存的8倍。因此,在GCN的私有内存上拥有某些东西意味着最终的性能,除非资源消耗停止了足以启动的波前(减少延迟隐藏)。

Nvidia也说类似的话:“通常,访问一个寄存器会消耗每条指令零个额外的时钟周期,但是
延迟可能会由于寄存器写后读依赖性和寄存器存储器而发生
银行冲突。
写后读取依赖项的等待时间约为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);
}


使用rgba(每个通道8位)的图像为512x512。

源图像(但在进行子步骤过滤之前将其大小调整为512x512):

enter image description here

结果图片:

enter image description here

我引用的文件:

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内存用于内核中的其他内容,则可以完全展开过滤器循环,删除过滤器数组,自己将这些araray元素放入展开的指令中(我尝试过,它将VGPR用法降低到21,SGPR使用率达到16)

作为参考,完全省去的滤波器计算平均可将执行时间减少0.05毫秒,而所有其他版本的执行时间则相同。

关于opencl - 在OpenCL内核中存储小型恒定值数组的最佳实践?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/42725075/

25 4 0
Copyright 2021 - 2024 cfsdn All Rights Reserved 蜀ICP备2022000587号
广告合作:1813099741@qq.com 6ren.com