Hi,
I am one of the developers of darktable an open source DAM software and raw developer for Linux, BSD and MacOS (http://www.darktable.org).
We are using OpenCL quite extensively for our pixelpipe processing. Recently we get reports from several of our users that a certain OpenCL code can not be compiled with NVIDIAs driver if user is working on a GT640. This was reported for drivers 304.51, 304.60 and 310.14. On other hardware (4xx and 5xx series) everything is fine.
The offending kernel is given below. Compilation fails with error code -30. Build log is empty, so no error messages.
__kernel void
blendop_mask_RAW (__read_only image2d_t in_a, __read_only image2d_t in_b, __write_only image2d_t mask, const int width, const int height,
const float gopacity, const int blendif, global const float *blendif_parameters)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
if(x >= width || y >= height) return;
write_imagef(mask, (int2)(x, y), gopacity);
}
This kernel is one in a series all with the same calling conventions. The other kernels (which get successfully compiled) read a pixel from in_a and in_b and then do some data processing before mask is written. Here is an example of one other kernel that works:
__kernel void
blendop_mask_Lab (__read_only image2d_t in_a, __read_only image2d_t in_b, __write_only image2d_t mask, const int width, const int height,
const float gopacity, const int blendif, global const float *blendif_parameters)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
if(x >= width || y >= height) return;
float4 a = read_imagef(in_a, sampleri, (int2)(x, y));
float4 b = read_imagef(in_b, sampleri, (int2)(x, y));
float opacity = gopacity * blendif_factor_Lab(a, b, blendif, blendif_parameters);
write_imagef(mask, (int2)(x, y), opacity);
}
Any ideas what is behind?
Ulrich