AnsweredAssumed Answered

AMD vs NVIDIA performance

Question asked by kcin on Jan 26, 2012
Latest reply on Feb 21, 2012 by LeeHowes

First, I must say, it isn't a holy war post.

I am a newbie with AMD GPUs (previously worked with NVIDIA) so I need help with understanding the unexpected performance lack.

I need to implement on AMD GPUs a simple kernel which searches template images on the base one. I was able to test the kernel on NVIDIA GT 240, AMD HD 6670 and AMD HD 5850. The results were discouraged for me:

GT 240 - 100 ms,

HD 6670 - 200 ms,

HD 5850 - 69 ms.

It seems strange to me as both AMD GPUs are more powerful than GT 240. It looks like I don't get the best of AMD.

Below is the kernel I'm talking about. It looks for 6 equally sized template images in the base one. Every pixel of the base image is considered as a possible origin of one of the templates and is processed by a separate thread. The kernel compares R, G, B values of each pixel of the base image and of the template image, and if the difference for at least one color exceeds diff parameter, the corresponding pixel is counted nonmatched. If the number of nonmatched pixels is less than maxNonmatchQt the corresponding template is hit.

 

__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;
          }
}

 

Kernel run configuration: Global work size = (1900, 1200) Local work size = (32, 8) for AMD and (32, 16) for NVIDIA.

Any remarks about my code are also highly appreciated.

Outcomes