Monday, February 14, 2011

OpenCL Driver for ATI Radeon is very bad on Mac OS X

I tested a simple kernel for image copy on both my iMac (ATI Radeon HD 5750) and MacBook Pro (Geforce 9600m GT). The MacBook Pro produces a right results, while iMac produces a weird results with some garbage data. So, This may be a bug in OpenCL driver for ATI Radeon on Mac OS X.


__kernel void image_copy_simple
__global unsigned char *src, 
__global unsigned char *dst,
const int pitch,
const int chans
)
{
  const uint gid_x = get_global_id(0) ;
  const uint gid_y = get_global_id(1) ;
  uint index = gid_x *chans + gid_y * pitch ;

  if(chans == 1) 
    dst[index] = src[index] ;
  else if(chans == 2) {
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
  }
  else if(chans == 3) {
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
  }
  else if(chans == 4) {
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
    index++ ;
    dst[index] = src[index] ;
  }
}


Since Apple does not update its graphics driver frequently, this bug may be exist until Lion (OS X 10.7) comes (?). A workaround for this problem is changing the type of 'dst' from 'unsigned char' to 'float' .

Another problem of the OpenCL driver on Mac OS X is it doesn't support the image object. What's funny is that Apple's OpenCL documentation explicitly mentions the usage of image object. But when I tried to use it on the machine with ATI hardware, it fails. Querying CL_DEVICE_IMAGE_SUPPORT via clGetDeviceInfo function gives me 'NO'.  By the way, NVIDIA driver supports the image object well.

No comments:

Post a Comment