OpenCL AMD与NVIDIA的表现

Question

我实现了一个简单的内核,它是某种卷积.我在NVIDIA GT 240上进行了测量.在CUDA上写入时需要70 ms,在OpenCL上写入时需要100 ms.好吧,我想,NVIDIA编译器更适合CUDA(或者我做错了).我需要在AMD GPU上运行它,所以我迁移到了AMD APP SDK.完全相同的内核代码.

我做了两次测试,他们的结果对我来说令人沮丧:HD 6670为200 ms,HD 5850为70 ms(与GT 240 + CUDA同时).而且我对这种奇怪行为的原因很感兴趣.

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

请不要将我的帖子视为NVIDIA广告.我很明白HD 5850比GT 240更强大.我唯一想知道的是为什么会出现这种差异以及如何解决问题.

更新.下面是内核代码,它在基础代码中查找6个大小相同的模板图像.基本图像的每个像素都被视为其中一个模板的可能来源,并由单独的线程处理.内核比较基本图像和模板1的每个像素的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).

执行时间:HD 5850 - 69 ms,HD 6670 - 200 ms,GT 240 - 100 ms.

关于我的代码的任何评论也非常感谢.

Answer 1

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

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

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

因此，编写提供良好灵活性的代码总是比编写特定于设备的优化更好。