OpenCL AMD 与 NVIDIA 性能对比
Posted
技术标签:
【中文标题】OpenCL AMD 与 NVIDIA 性能对比【英文标题】:OpenCL AMD vs NVIDIA performance 【发布时间】:2012-02-16 18:17:28 【问题描述】:我实现了一个简单的内核,它是某种卷积。我在 NVIDIA GT 240 上进行了测量。在 CUDA 上编写需要 70 毫秒,在 OpenCL 上编写需要 100 毫秒。好的,我想,NVIDIA 编译器更适合 CUDA(或者我做错了什么)。 我需要在 AMD GPU 上运行它,所以我迁移到 AMD APP SDK。完全相同的内核代码。
我进行了两次测试,结果让我感到沮丧:HD 6670 为 200 毫秒,HD 5850 为 70 毫秒(与 GT 240 + CUDA 相同)。我对这种奇怪行为的原因非常感兴趣。
所有项目均使用 NVIDIA 和 AMD 示例项目的设置在 VS2010 上构建。
请不要将我的帖子视为 NVIDIA 广告。我相当了解 HD 5850 比 GT 240 更强大。我唯一想知道的是为什么会出现这种差异以及如何解决问题。
更新。下面是内核代码,它在基础图像中查找 6 个大小相同的模板图像。基本图像的每个像素都被视为模板之一的可能来源,并由单独的线程处理。内核比较基础图像和模板图像的每个像素的 R、G、B 值,如果至少有一个差异超过diff
参数,则将相应像素计为不匹配。如果不匹配的像素数小于maxNonmatchQt
,则命中相应的模板。
__constant int tOffset = 8196; // one template size in memory (in bytes)
__kernel void matchImage6( __global unsigned char* image, // pointer to the base image
int imgWidth, // base image width
int imgHeight, // base image height
int imgPitch, // base image pitch (in bytes)
int imgBpp, // base image bytes (!) per pixel
__constant unsigned char* templates, // pointer to the array of templates
int tWidth, // templates width (the same for all)
int tHeight, // templates height (the same for all)
int tPitch, // templates pitch (in bytes, the same for all)
int tBpp, // templates bytes (!) per pixel (the same for all)
int diff, // max allowed difference of intensity
int maxNonmatchQt, // max number of nonmatched pixels
__global int* result, // results
)
int x0 = (int)get_global_id(0);
int y0 = (int)get_global_id(1);
if( x0 + tWidth > imgWidth || y0 + tHeight > imgHeight)
return;
int nonmatchQt[] = 0, 0, 0, 0, 0, 0;
for( int y = 0; y < tHeight; y++)
int ind = y * tPitch;
int baseImgInd = (y0 + y) * imgPitch + x0 * imgBpp;
for( int x = 0; x < tWidth; x++)
unsigned char c0 = image[baseImgInd];
unsigned char c1 = image[baseImgInd + 1];
unsigned char c2 = image[baseImgInd + 2];
for( int i = 0; i < 6; i++)
if( abs( c0 - templates[i * tOffset + ind]) > diff ||
abs( c1 - templates[i * tOffset + ind + 1]) > diff ||
abs( c2 - templates[i * tOffset + ind + 2]) > diff)
nonmatchQt[i]++;
ind += tBpp;
baseImgInd += imgBpp;
if( nonmatchQt[0] > maxNonmatchQt && nonmatchQt[1] > maxNonmatchQt && nonmatchQt[2] > maxNonmatchQt && nonmatchQt[3] > maxNonmatchQt && nonmatchQt[4] > maxNonmatchQt && nonmatchQt[5] > maxNonmatchQt)
return;
for( int i = 0; i < 6; i++)
if( nonmatchQt[i] < maxNonmatchQt)
unsigned int pos = atom_inc( &result[0]) * 3;
result[pos + 1] = i;
result[pos + 2] = x0;
result[pos + 3] = y0;
内核运行配置: 全局工作大小 = (1900, 1200) 本地工作大小 = AMD 的 (32, 8) 和 NVIDIA 的 (32, 16)。
执行时间: 高清 5850 - 69 毫秒, 高清 6670 - 200 毫秒, GT 240 - 100 毫秒。
对我的代码的任何评论也非常感谢。
【问题讨论】:
这里没有足够的信息来回答这个问题!每个 NVidia 和 AMD 卡在架构上都是棘手的野兽,你看到的性能很大程度上取决于代码。了解两者之间的性能差异更加棘手。你能发布你的内核和驱动程序吗? 您在内核中运行什么样的算法?内存访问模式?波前/翘曲尺寸?需要更多信息才能提供建议。 你启动了多少线程?你在向量化数组吗? 谢谢大家的回答。我明天会发布我的内核代码。 maxNonmatchQt 和模板大小 (tW*tH) 之间的比例应该是多少?此信息可以帮助优化。您是每像素使用 3 个字节,还是每像素使用 3*n 字节?我假设每个 RGB 通道都使用相同的字节数。它是否正确?另外,你能解释一下 imgPitch 和 tPitch 吗?我不能完全从代码中弄清楚。谢谢 【参考方案1】:执行时间的差异是由编译器造成的。 您的代码可以轻松矢量化。将图像和模板视为向量类型 char4 的数组(每个 char4 向量的第四坐标始终为 0)。而不是 3 次内存读取:
unsigned char c0 = image[baseImgInd];
unsigned char c1 = image[baseImgInd + 1];
unsigned char c2 = image[baseImgInd + 2];
只使用一个:
unsigned char4 c = image[baseImgInd];
如果:
if( abs( c0 - templates[i * tOffset + ind]) > diff ||
abs( c1 - templates[i * tOffset + ind + 1]) > diff ||
abs( c2 - templates[i * tOffset + ind + 2]) > diff)
nonmatchQt[i]++;
快速使用:
unsigned char4 t = templates[i * tOffset + ind];
nonmatchQt[i] += any(abs_diff(c,t)>diff);
因此,您可以将代码的性能提高多达 3 倍(如果编译器不自行对代码进行矢量化)。我想 AMD OpenCL 编译器没有这种矢量化和其他优化。根据我的经验,NVIDIA GPU 上的 OpenCL 通常可以比 CUDA 更快,因为它更底层。
【讨论】:
HD5850 和 HD6670 GPU 都具有支持矢量化代码的架构,但 AMD 的编译器不一定足够聪明,无法进行这些转换。 Nvidia GPU 上使用的架构不支持矢量化。 你是对的。 NVIDIA GPU 不需要矢量化,但 AMD GPU 需要。【参考方案2】:对此没有完全完美的答案。 OpenCL 性能取决于许多参数。访问全局内存的次数、代码的效率等。此外,在两个设备之间进行比较非常困难,因为它们可能具有不同的本地、全局、常量内存。核心数量、频率、内存带宽,更重要的是硬件架构等。
每个硬件都提供自己的性能提升,例如来自 NVIDIA 的 native_。因此,您需要更多地探索您正在使用的硬件,这可能确实有效。但我个人建议不要使用这种硬件特定的优化,这可能会影响代码的灵活性。
您还可以找到一些发表的论文表明,在相同的 NVIDIA 硬件上,CUDA 的性能比 OpenCL 的性能要好得多。
因此,编写提供良好灵活性的代码总是比针对特定设备的优化更好。
【讨论】:
以上是关于OpenCL AMD 与 NVIDIA 性能对比的主要内容,如果未能解决你的问题,请参考以下文章
使用 Nvidia 显卡安装 AMD OpenCL CPU 驱动程序
在 64 位 Ubuntu 14.04 中使用 Nvidia *和* AMD GPU 进行 OpenCL 开发
可以仅将 AMD gDEBugger 6.0 用于 OpenCL CPU 内核调试吗? NVIDIA GPU 似乎推迟了这一点