- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
我正在编写一个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
是否可以一起使用。
filter
可以存储为
__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}
};
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}
};
// 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}
};
__constant
内存访问(至少在amd gpu中)更好,而对于相同索引访问则不错(组中的所有线程都访问同一索引,就像在此示例中一样(int嵌套循环))。在这些寻址方式下,常量内存比全局内存要快,但是使用不同的索引时,常量内存与全局内存访问(甚至命中缓存)没有什么不同。 “对于全局范围的常量数组,如果数组的大小小于64 kB,则将其放置在硬件常量缓冲区中;否则,它将使用全局内存”。 (存在与Amd-GCN架构相关的问题,但Nvidia和Intel可以预期类似的行为)
// 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}
};
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);
}
关于opencl - 在OpenCL内核中存储小型恒定值数组的最佳实践?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/42725075/
Intel、AMD 和 Khronos OpenCL 之间有什么区别。我对 OpenCL 完全陌生,想从它开始。我不知道在我的操作系统上安装哪个更好。 最佳答案 OpenCL 是 C 和 C++ 语言
我在这里的一篇文章中看到,我们可以从 OpenCL 内核调用函数。但是在我的情况下,我还需要并行化该复杂函数(由所有可用线程运行),所以我是否必须将该函数也设为内核并像从主内核中调用函数一样直接调
最近我看到一些开发板支持 OpenCL EP,例如 odroid XU。我知道的一件事是 OpenCL EP 适用于 ARM 处理器,但它与基于主要桌面的 OpenCL 在哪些特性上有所不同。 最佳答
我想知道在 OpenCL 中设置为内核函数的参数数量是否有任何限制。设置参数时出现 INVALID_ARG_INDEX 错误。我在内核函数中设置了 9 个参数。请在这方面帮助我。 最佳答案 您可以尝试
我对零拷贝的工作原理有点困惑。 1-要确认以下内容对应于opencl中的零拷贝。 ....................... . . . .
我是 OpenCL 的初学者,我很难理解某些东西。 我想改进主机和设备之间的图像传输。 我制定了一个计划以更好地了解我。 顶部:我现在拥有的 |底部:我想要的 HtD(主机到设备)和 DtH(设备到主
今天我又加了四个 __local变量到我的内核以转储中间结果。但是只需将另外四个变量添加到内核的签名并添加相应的内核参数就会将内核的所有输出呈现为“0”。没有一个 cl 函数返回错误代码。 我进一步尝
我知道工作项被分组到工作组中,并且您不能在工作组之外进行同步。 这是否意味着工作项是并行执行的? 如果是这样,使用 128 个工作项创建 1 个工作组是否可能/有效? 最佳答案 组内的工作项将一起安排
我相当确定经纱仅在 CUDA 中定义。但也许我错了。就 OpenCL 而言,什么是扭曲? 它与工作组不一样,是吗? 任何相关的反馈都受到高度赞赏。谢谢! 最佳答案 它没有在 OpenCL 标准中定义。
已结束。此问题正在寻求书籍、工具、软件库等的推荐。它不满足Stack Overflow guidelines 。目前不接受答案。 我们不允许提出寻求书籍、工具、软件库等推荐的问题。您可以编辑问题,以便
在OpenCL中,我的理解是可以使用barrier()函数来同步工作组中的线程。我(通常)确实了解它们的用途以及何时使用它们。我还知道工作组中的所有线程都必须遇到障碍,否则会出现问题。然而,到目前为止
我的主板上有 Nvidia 显卡 (GeForce GT 640)。我已经在我的盒子上安装了 OpenCL。当我使用“clGetPlatformInfo(参数)”查询平台时,我看到以下输出:-#可用平
我目前正在构建一个 ray marcher 来查看像 mandelbox 等东西。它工作得很好。但是,在我当前的程序中,它使用每个 worker 作为从眼睛转换的光线。这意味着每个 worker 有大
我编写了两个不同的 openCl 内核,使用 nvidia profiler 获取了有关它们的一些信息,发现两者每个工作项都使用 63 个寄存器。 我尝试了一切我能想到的方法来降低这个数字(用 ush
我的主板上有 Nvidia 显卡 (GeForce GT 640)。我已经在我的盒子上安装了 OpenCL。当我使用“clGetPlatformInfo(参数)”查询平台时,我看到以下输出:-#可用平
我目前正在构建一个 ray marcher 来查看像 mandelbox 等东西。它工作得很好。但是,在我当前的程序中,它使用每个 worker 作为从眼睛转换的光线。这意味着每个 worker 有大
我正在尝试使用 OpenCL 加速一些计算,算法的一部分包括矩阵求逆。是否有任何开源库或免费可用的代码来计算用 OpenCL 或 CUDA 编写的矩阵的 lu 分解(lapack dgetrf 和 d
我正在尝试在 OpenCL 内核中使用递归。编译成功,但运行时出现编译错误,所以我想知道,由于 CUDA 现在支持动态并行,OpenCL 是否支持动态并行? 最佳答案 OpenCL 不支持递归。请参阅
考虑以下代码,它从大小为 size 的 double 组创建缓冲区内存对象: coef_mem = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM
OpenCL 中目标平台的示例是什么?例如,它是 Windows、Android、Mac 等操作系统,还是设备中的实际芯片? 最佳答案 OpenCL 平台本质上是一个 OpenCL 实现。它与操作系统
我是一名优秀的程序员,十分优秀!