Link Error when using max_constant_size attribute on 4850

Discussion created by timchist on May 28, 2010
Latest reply on Jun 24, 2010 by timchist
Link Error when using max_constant_size attribute on ATI 4850

I'm trying to build the following kernel:

__kernel void ConstDemo1(__global float* src, __global float* dst, int width, int height, __constant float *constDemoData __attribute__((max_constant_size(256))))

    int x = get_global_id(0);
    if(x < width)
        int y = get_global_id(1);
        if(y < height)
            int offset = y * width + x;
            float sum = 0;
            for(int i = 0; i < 20; i++)
               sum += constDemoData[(offset + i) * 17 % 1000];
            dst[offset] = sum;


When I call clBuildProgram(), it returns CL_SUCCESS. Then I extract binary using clGetProgramInfo(...CL_PROGRAM_BINARIES...), and create another cl_program object using clCreateProgramWithBinary. Now, when I call clBuildProgram() for this new program, it returns -11. clGetProgramBuildInfo(...CL_PROGRAM_BUILD_LOG...) returns

Internal error: Link failed.
Make sure the system setup is correct.

When I remove the max_constant_size attribute, everything works fine. However, without the attribute it does not look that readings are cached (the performance is the same as if constDemoData was a __global pointer).

I have an iMac 27" with ATI Radeon 4850. Driver version is 10.5 (8.732.0.0). ATI Stream SDK 2.1 64 bit. Windows 7 64 bit, a 32-bit application.