我实现了一个简单的卷积核。我在NVIDIA GT 240 上进行了测试，CUDA写的时候需要70ms，OpenCL写的时候需要100ms。我想，可能是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”，则击中相应模板。
请不要将我的帖子视为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) 本地工作大小 = (32, 8) 对于 AMD 和 (32, 16) 对于 NVIDIA。
执行时间: HD 5850 - 69 毫秒, HD 6670 - 200 毫秒, GT 240 - 100 毫秒。
任何关于我的代码的评论也非常感谢。