nvidia driver differences 275.50 and 280.19 for GTX480 GPU in opencl resampling

Question

the resampling of an image with driver version 275.50 and 280.19 for Nvidia GTX480 GPU gives little changes in the gray values. Maybe it is an interpolation issue. I can not determine what has changed in the new version except for the implementation of OpenCL version 1.1. Using only OpenCL 1.0 gives the same little changes in the gray value.

In following there is the code which gives different results betwenn driver version 275.50 and 280.19:

success oclInitImgData(struct _Image2d *image)
{
cl_image_format volume_format;
volume_format.image_channel_order       = CL_R;
volume_format.image_channel_data_type   = CL_UNORM_INT16;

size = len[0] * sizeof(unsigned short);
img_h = clCreateImage2D(Ocl._GPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    &volume_format, len[0], len[1], size, data, &error);

resampledImg_h = clCreateBuffer(Ocl._GPUContext, CL_MEM_READ_WRITE, size, NULL, &error);
oclResampleImg(Ocl._GPUCommandQueue[posSet], Ocl._pGPUKernels[posSet][K_IMG_RESAMPLE], img_h, resampledImg_h , Size, PixelSize, mm_per_p, vm_h);
}

__kernel void resampleImage( __read_only image2d_t IN image, __global uint OUT *resampledImage)
    {
//get resampled position
int2 posResampledImg = (int2)(get_global_id(0), get_global_id(1));

//get dimension of the image
int2 imageSize = get_image_dim(image);

//calulate image size in mm
float2 imageSizemm = convert_float2(imageSize)*imagePixelSize;
float2 posImg = (posResampledImgmm + (imageSizemm*0.5f));
float2 posImgnorm = posImg/imageSizemm;
int2 posImgpix = convert_int2(posImg/imagePixelSize);

uint sample = 0;
if(read_imageui(image, CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST, posImgpix).x == 1)
{
    //get grey value and store it in the resampled image
    sample = read_imagef(image, CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_LINEAR, posImgnorm).x * 65535.0f;
}
resampledImg[posResampledImg.y*convert_int(imageSizeResampled.x)+posResampledImg.x] = sample;
}

Answer 1

To get a useful result you should:

(1) Understand if the difference in behavior is covered by the OpenCL spec.

(2) If yes, create a testcase that demonstrates it

(3) File a bug report with Nvidia

(4) ???

(5) Profit

I don't see how people here can help you beyond steps 1-5.