opencl做图像sobel的问题

[shafa00419]

找到一个例程，写了4了kernel：v1到v4
v1是每个workitem处理一个像素点的sobel。到v2，v3是一个workitem处理16个像素点，v4是一个workitem处理16*16个像素点。处理2048*2048大小的图像，结果看处理时间，v1是最慢的。。。想不通啊，workitem处理的像素点越多应该时间越久啊。
__kernel void Sobel_v1_uchar (__global uchar *pSrcImage, __global uchar *pDstImage)
{
uint dstYStride = get_global_size(0);
uint dstIndex = get_global_id(1) * dstYStride + get_global_id(0);
uint srcYStride = dstYStride + 32;
uint srcIndex = get_global_id(1) * srcYStride + get_global_id(0) + 16;

uint a, b, c;
uint d, /*center*/ f;
uint g, h, i;

// Read data in
a = pSrcImage[srcIndex-1]; b = pSrcImage[srcIndex]; c = pSrcImage[srcIndex+1];
srcIndex += srcYStride;
d = pSrcImage[srcIndex-1]; /*center*/ f = pSrcImage[srcIndex+1];
srcIndex += srcYStride;
g = pSrcImage[srcIndex-1]; h = pSrcImage[srcIndex]; i = pSrcImage[srcIndex+1];

uint xVal = a* 1 + c*-1 +
d* 2 + /*center*/ f*-2 +
g* 1 + i*-1;

uint yVal = a* 1 + b* 2 + c* 1 +
/*center*/
g*-1 + h*-2 + i*-1;

// Write data out
pDstImage[dstIndex] = min((uint)255, (uint)sqrt((float)(xVal*xVal + yVal*yVal)));
}

__kernel void Sobel_v2_uchar16 (__global uchar16* pSrcImage, __global uchar16* pDstImage)
{
uint dstYStride = get_global_size(0);
uint dstIndex = get_global_id(1) * dstYStride + get_global_id(0);
uint srcYStride = dstYStride + 2;
uint srcIndex = get_global_id(1) * srcYStride + get_global_id(0) + 1;

uint a; uint16 b; uint c;
uint d; uint16 e; uint f;
uint g; uint16 h; uint i;

// Read data in
a = ((__global uchar*)(pSrcImage+srcIndex))[-1]; b = convert_uint16(pSrcImage[srcIndex]); c = ((__global uchar*)(pSrcImage+srcIndex))[16];
srcIndex += srcYStride;
d = ((__global uchar*)(pSrcImage+srcIndex))[-1]; e = convert_uint16(pSrcImage[srcIndex]); f = ((__global uchar*)(pSrcImage+srcIndex))[16];
srcIndex += srcYStride;
g = ((__global uchar*)(pSrcImage+srcIndex))[-1]; h = convert_uint16(pSrcImage[srcIndex]); i = ((__global uchar*)(pSrcImage+srcIndex))[16];

uint16 xVal, yVal;

xVal = (uint16)(a, b.s0123, b.s456789ab, b.scde) - (uint16)(b.s123, b.s4567, b.s89abcdef, c) +
2*(uint16)(d, e.s0123, e.s456789ab, e.scde) - 2*(uint16)(e.s123, e.s4567, e.s89abcdef, f) +
(uint16)(g, h.s0123, h.s456789ab, h.scde) - (uint16)(h.s123, h.s4567, h.s89abcdef, i);

yVal = (uint16)(a, b.s0123, b.s456789ab, b.scde) + 2*b + (uint16)(b.s123, b.s4567, b.s89abcdef, c) -
(uint16)(g, h.s0123, h.s456789ab, h.scde) - 2*h - (uint16)(h.s123, h.s4567, h.s89abcdef, i);

// Write data out
pDstImage[dstIndex] = convert_uchar16(min((float16)255.0f, sqrt(convert_float16(xVal*xVal + yVal*yVal))));
}

__kernel void Sobel_v3_uchar16_to_float16 (__global uchar16* pSrcImage, __global uchar16* pDstImage)
{
uint dstYStride = get_global_size(0);
uint dstIndex = get_global_id(1) * dstYStride + get_global_id(0);
uint srcYStride = dstYStride + 2;
uint srcIndex = get_global_id(1) * srcYStride + get_global_id(0) + 1;

float a; float16 b; float c;
float d; float16 e; float f;
float g; float16 h; float i;

// Read data in
a = convert_float(((__global uchar*)(pSrcImage+srcIndex))[-1]);
b = convert_float16(pSrcImage[srcIndex]);
c = convert_float(((__global uchar*)(pSrcImage+srcIndex))[16]);
srcIndex += srcYStride;
d = convert_float(((__global uchar*)(pSrcImage+srcIndex))[-1]);
e = convert_float16(pSrcImage[srcIndex]);
f = convert_float(((__global uchar*)(pSrcImage+srcIndex))[16]);
srcIndex += srcYStride;
g = convert_float(((__global uchar*)(pSrcImage+srcIndex))[-1]);
h = convert_float16(pSrcImage[srcIndex]);
i = convert_float(((__global uchar*)(pSrcImage+srcIndex))[16]);

float16 xVal, yVal;

xVal = (float16)(a, b.s0123, b.s456789ab, b.scde) - (float16)(b.s123, b.s4567, b.s89abcdef, c) +
2.0f*(float16)(d, e.s0123, e.s456789ab, e.scde) - 2.0f*(float16)(e.s123, e.s4567, e.s89abcdef, f) +
(float16)(g, h.s0123, h.s456789ab, h.scde) - (float16)(h.s123, h.s4567, h.s89abcdef, i);

yVal = (float16)(a, b.s0123, b.s456789ab, b.scde) + 2.0f*b + (float16)(b.s123, b.s4567, b.s89abcdef, c) -
(float16)(g, h.s0123, h.s456789ab, h.scde) - 2.0f*h - (float16)(h.s123, h.s4567, h.s89abcdef, i);

// Write data out
pDstImage[dstIndex] = convert_uchar16(min((float16)255.0f, sqrt(xVal*xVal + yVal*yVal)));
}

__kernel void Sobel_v4_uchar16_to_float16_16(__global uchar16* pSrcImage, __global uchar16* pDstImage)
{
uint dstYStride = get_global_size(0);
uint dstIndex = 16 * get_global_id(1) * dstYStride + get_global_id(0);
uint srcYStride = dstYStride + 2;
uint srcIndex = 16 * get_global_id(1) * srcYStride + get_global_id(0) + 1;

float a; float16 b; float c;
float d; float16 e; float f;
float g; float16 h; float i;

// Read data in
a = convert_float(((__global uchar*)(pSrcImage + srcIndex))[-1]);
b = convert_float16(pSrcImage[srcIndex]);
c = convert_float(((__global uchar*)(pSrcImage + srcIndex))[16]);
srcIndex += srcYStride;
d = convert_float(((__global uchar*)(pSrcImage + srcIndex))[-1]);
e = convert_float16(pSrcImage[srcIndex]);
f = convert_float(((__global uchar*)(pSrcImage + srcIndex))[16]);

for (uint k = 0; k < 16; k++) {
srcIndex += srcYStride;
g = convert_float(((__global uchar*)(pSrcImage + srcIndex))[-1]);
h = convert_float16(pSrcImage[srcIndex]);
i = convert_float(((__global uchar*)(pSrcImage + srcIndex))[16]);

float16 xVal, yVal;

xVal = (float16)(a, b.s0123, b.s456789ab, b.scde) - (float16)(b.s123, b.s4567, b.s89abcdef, c) +
2.0f*(float16)(d, e.s0123, e.s456789ab, e.scde) - 2.0f*(float16)(e.s123, e.s4567, e.s89abcdef, f) +
(float16)(g, h.s0123, h.s456789ab, h.scde) - (float16)(h.s123, h.s4567, h.s89abcdef, i);

yVal = (float16)(a, b.s0123, b.s456789ab, b.scde) + 2.0f*b + (float16)(b.s123, b.s4567, b.s89abcdef, c) -
(float16)(g, h.s0123, h.s456789ab, h.scde) - 2.0f*h - (float16)(h.s123, h.s4567, h.s89abcdef, i);

// Write data out
pDstImage[dstIndex] = convert_uchar16(min((float16)255.0f, sqrt(xVal*xVal + yVal*yVal)));

a = d; b = e; c = f;
d = g; e = h; f = i;
dstIndex += dstYStride;
}
}

[shafa00419]

还是说我的理解错了？

[shafa00419]

问到了。 gpu的操作分访问内存操作和运算操作（两种操作速度有差异），本例中应该是一个workitem处理一个像素和处理16个像素的内存访问时间基本一样（应该是矢量类型的功劳），处理16个像素时数据的重复利用率高，节省了运算时间。

[bluewanderer]

引用 2 楼 shafa00419 的回复:

问到了。 gpu的操作分访问内存操作和运算操作（两种操作速度有差异），本例中应该是一个workitem处理一个像素和处理16个像素的内存访问时间基本一样（应该是矢量类型的功劳），处理16个像素时数据的重复利用率高，节省了运算时间。

我只想说... 这边是可以格式化代码的... 另外，并不是，是这个算法数据源的一个像素会被结果里八个像素用到。第一个算法每个像素要读八次，后面的都减少了读取次数。然而第三个我依然看着头疼。

[shafa00419]

这里也能见到你，，大神无处不在。。。刚发现可以格式化代码。。。
这两又想了想。。画了幅图。。（手残作品，功力有限，，挺难看的。。。）
（图片大，传不上，给链接）http://pan.baidu.com/s/1c2sjgoo
由于是并行计算，v1的计算时间应该是一个工作项的计算时间，v1中一个工作项读取8个像素，你说应该是的一个像素会被八个工作项读取。
v3的处理过程v2完全相同，只是计算时将v2的uchar16数据类型转换成float16，结果速度就比v2的快，是不是说gpu处理浮点型的计算比整型的速度快？

[bluewanderer]

引用 4 楼 shafa00419 的回复:

这里也能见到你，，大神无处不在。。。刚发现可以格式化代码。。。 这两又想了想。。画了幅图。。（手残作品，功力有限，，挺难看的。。。） （图片大，传不上，给链接）http://pan.baidu.com/s/1c2sjgoo 由于是并行计算，v1的计算时间应该是一个工作项的计算时间，v1中一个工作项读取8个像素，你说应该是的一个像素会被八个工作项读取。 v3的处理过程v2完全相同，只是计算时将v2的uchar16数据类型转换成float16，结果速度就比v2的快，是不是说gpu处理浮点型的计算比整型的速度快？

我发现我一直压根没发现真正的第三个，我一直以为第四个是第三个... _(:3」∠)_ 至于为什么第三个比第二个快，不知道。不方便试。现在一般的GPU整数加法和浮点数加法速度应该是一样的。区别也就是浮点数可以一个周期执行完a += b * 2的操作，整数不手动优化可能需要两个周期。但是整数转浮点数也是需要2到4个周期的操作。而且直观上看所有这些操作都不足以掩盖内存操作的延迟... 不负责任猜测有可能是指令的排布的偶然变化导致硬件利用率得到提升。至少AMD显卡做OpenCL经常得看编译出来的到底是什么鬼，编译结果基本上100%会有幺蛾子。

[bluewanderer]

GPU编程讲究内存操作和运算穿插，向量运算和标量运算穿插，一个Wavefront一次使用的内存通道尽量少。于是其实这种读一大坨慢慢用的手法对GPU来说是相当不友好的，怎么优化全看编译器心情。

[bluewanderer]

利用这一轮最后一次说话机会确认一下你的显卡什么型号。如果是老AMD显卡的VLIW结构，那和现在普遍的SIMD结构就相当不一样并且我完全不了解了。_(:3」∠)_

[shafa00419]

显卡是Intel HD Graphics 530 这个程序是intel的一个例程。 https://software.intel.com/en-us/INDE-OpenCL-Sobel 虽然问题没有解决，这两天也收获挺多，多谢你了blue

[bluewanderer]

引用

Convert to float16 for 2X math performance

这种说法的话估计是Intel核显整型运算很慢了，同样的处理恐怕在别的显卡上不适用。另外我有点担心这么算搞不好核显处理这种字节数据的性能拼不过CPU一个核... 毕竟CPU上字节的加减法的速度光算峰值也是单精度的4倍。