OpenCL AMD与NVIDIA的性能

我实现了一个简单的内核，它是某种卷积。我在NVIDIA GT 240上测量了它。在CUDA上写需要70毫秒，在OpenCL上写需要100毫秒。好吧，我想，NVIDIA编译器更好地针对CUDA进行了优化（或者我做错了什么）。我需要在AMD GPU上运行它，所以我迁移到AMD APP SDK。完全相同的内核代码。

我做了两次测试，结果让我沮丧：HD 6670时为200毫秒，HD 5850时为70毫秒（与GT 240+CUDA相同）。我对这种奇怪行为的原因很感兴趣。

所有项目都是在VS2010上构建的，分别使用NVIDIA和AMD的示例项目中的设置。

请不要将我的帖子视为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;
    }
}

内核运行配置：全球工作规模=（19001200）对于AMD，本地工作大小=（32，8），对于NVIDIA，则为（32，16）。

执行时间：HD 5850-69毫秒，HD 6670-200毫秒，GT 240-100毫秒

任何关于我的代码的评论也将不胜感激。

执行时间的差异是由编译器引起的。您的代码可以很容易地向量化。将图像和模板视为向量类型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：

    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更快，因为它更低级。

这不可能有确切的完美答案。OpenCL的性能取决于许多参数。访问全局内存的次数、代码的效率等。此外，在两个设备之间进行比较非常困难，因为它们可能具有不同的本地、全局、恒定内存。核心数量、频率、内存带宽，更重要的是硬件架构等。

每个硬件都提供自己的性能提升，例如NVIDIA的native_。因此，你需要探索更多关于你正在工作的硬件的信息，这些硬件可能确实有效。但我个人建议不要使用这种特定于硬件的优化，这可能会影响代码的灵活性。

您还可以找到一些发表的论文显示，在相同的NVIDIA硬件上，CUDA的性能比OpenCL的性能要好得多。

相关内容

最新更新

热门标签：