OpenCL: Странное поведение буфера или изображения с NVidia, но не с Amd

У меня есть большая проблема (в Linux): я создаю буфер с определенными данными, затем ядро ​​OpenCL берет эти данные и помещает их в image2d_t. При работе на AMD C50 (Fusion CPU/GPU) программа работает как надо, но на моей GeForce 9500 GT данное ядро ​​очень редко вычисляет правильный результат. Иногда результат правильный, но очень часто неверный. Иногда это зависит от очень странных изменений, таких как удаление неиспользуемых объявлений переменных или добавление новой строки. Я понял, что отключение оптимизации повысит вероятность сбоя. У меня самый актуальный драйвер дисплея в обеих системах.

Вот мой сокращенный код:

#include <CL/cl.h>
#include <string>
#include <iostream>
#include <sstream>
#include <cmath>

    void checkOpenCLErr(cl_int err, std::string name){
        const char* errorString[] = {
            "CL_SUCCESS",
            "CL_DEVICE_NOT_FOUND",
            "CL_DEVICE_NOT_AVAILABLE",
            "CL_COMPILER_NOT_AVAILABLE",
            "CL_MEM_OBJECT_ALLOCATION_FAILURE",
            "CL_OUT_OF_RESOURCES",
            "CL_OUT_OF_HOST_MEMORY",
            "CL_PROFILING_INFO_NOT_AVAILABLE",
            "CL_MEM_COPY_OVERLAP",
            "CL_IMAGE_FORMAT_MISMATCH",
            "CL_IMAGE_FORMAT_NOT_SUPPORTED",
            "CL_BUILD_PROGRAM_FAILURE",
            "CL_MAP_FAILURE",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "",
            "CL_INVALID_VALUE",
            "CL_INVALID_DEVICE_TYPE",
            "CL_INVALID_PLATFORM",
            "CL_INVALID_DEVICE",
            "CL_INVALID_CONTEXT",
            "CL_INVALID_QUEUE_PROPERTIES",
            "CL_INVALID_COMMAND_QUEUE",
            "CL_INVALID_HOST_PTR",
            "CL_INVALID_MEM_OBJECT",
            "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR",
            "CL_INVALID_IMAGE_SIZE",
            "CL_INVALID_SAMPLER",
            "CL_INVALID_BINARY",
            "CL_INVALID_BUILD_OPTIONS",
            "CL_INVALID_PROGRAM",
            "CL_INVALID_PROGRAM_EXECUTABLE",
            "CL_INVALID_KERNEL_NAME",
            "CL_INVALID_KERNEL_DEFINITION",
            "CL_INVALID_KERNEL",
            "CL_INVALID_ARG_INDEX",
            "CL_INVALID_ARG_VALUE",
            "CL_INVALID_ARG_SIZE",
            "CL_INVALID_KERNEL_ARGS",
            "CL_INVALID_WORK_DIMENSION",
            "CL_INVALID_WORK_GROUP_SIZE",
            "CL_INVALID_WORK_ITEM_SIZE",
            "CL_INVALID_GLOBAL_OFFSET",
            "CL_INVALID_EVENT_WAIT_LIST",
            "CL_INVALID_EVENT",
            "CL_INVALID_OPERATION",
            "CL_INVALID_GL_OBJECT",
            "CL_INVALID_BUFFER_SIZE",
            "CL_INVALID_MIP_LEVEL",
            "CL_INVALID_GLOBAL_WORK_SIZE",
        };
        if (err != CL_SUCCESS) {
            std::stringstream str;
            str << errorString[-err] << " (" << err << ")";
            throw std::string(name)+(str.str());
        }
    }

int main(){
    try{
        cl_context m_context;
        cl_platform_id* m_platforms;
        unsigned int m_numPlatforms;
        cl_command_queue m_queue;
        cl_device_id m_device;
        cl_int error = 0;   // Used to handle error codes
        clGetPlatformIDs(0,NULL,&m_numPlatforms);
        m_platforms = new cl_platform_id[m_numPlatforms];
        error = clGetPlatformIDs(m_numPlatforms,m_platforms,&m_numPlatforms);
        checkOpenCLErr(error, "getPlatformIDs");

        // Device
        error = clGetDeviceIDs(m_platforms[0], CL_DEVICE_TYPE_GPU, 1, &m_device, NULL);
        checkOpenCLErr(error, "getDeviceIDs");

        // Context
        cl_context_properties properties[] =
            { CL_CONTEXT_PLATFORM, (cl_context_properties)(m_platforms[0]), 0};
        m_context = clCreateContextFromType(properties, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);
        //  m_private->m_context = clCreateContext(properties, 1, &m_private->m_device, NULL, NULL, &error);
        checkOpenCLErr(error, "Create context");
        // Command-queue
        m_queue = clCreateCommandQueue(m_context, m_device, 0, &error);
        checkOpenCLErr(error, "Create command queue");
        //Build program and kernel
        const char* source = "#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable\n"
            "\n"
            "__kernel void bufToImage(__global unsigned char* in,  __write_only image2d_t out, const unsigned int offset_x, const unsigned int image_width , const unsigned int maxval ){\n"
                "\tint i = get_global_id(0);\n"
                "\tint j = get_global_id(1);\n"
                "\tint width = get_global_size(0);\n"
                "\tint height = get_global_size(1);\n"
                "\n"
                "\tint pos = j*image_width*3+(offset_x+i)*3;\n"
                "\tif( maxval < 256 ){\n"
                    "\t\tfloat4 c = (float4)(in[pos],in[pos+1],in[pos+2],1.0f);\n"
                    "\t\tc.x /= maxval;\n"
                    "\t\tc.y /= maxval;\n"
                "\t\tc.z /= maxval;\n"
                "\t\twrite_imagef(out, (int2)(i,j), c);\n"
            "\t}else{\n"
                "\t\tfloat4 c = (float4)(255.0f*in[2*pos]+in[2*pos+1],255.0f*in[2*pos+2]+in[2*pos+3],255.0f*in[2*pos+4]+in[2*pos+5],1.0f);\n"
                "\t\tc.x /= maxval;\n"
                "\t\tc.y /= maxval;\n"
                "\t\tc.z /= maxval;\n"
                "\t\twrite_imagef(out, (int2)(i,j), c);\n"
            "\t}\n"
        "}\n"
        "\n"
        "__constant sampler_t imageSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n"
        "\n"
        "__kernel void imageToBuf(__read_only image2d_t in, __global unsigned char* out, const unsigned int offset_x, const unsigned int image_width ){\n"
            "\tint i = get_global_id(0);\n"
            "\tint j = get_global_id(1);\n"
            "\tint pos = j*image_width*3+(offset_x+i)*3;\n"
            "\tfloat4 c = read_imagef(in, imageSampler, (int2)(i,j));\n"
            "\tif( c.x <= 1.0f && c.y <= 1.0f && c.z <= 1.0f ){\n"
                "\t\tout[pos] = c.x*255.0f;\n"
                "\t\tout[pos+1] = c.y*255.0f;\n"
                "\t\tout[pos+2] = c.z*255.0f;\n"
            "\t}else{\n"
                "\t\tout[pos] = 200.0f;\n"
                "\t\tout[pos+1] = 0.0f;\n"
                "\t\tout[pos+2] = 255.0f;\n"
            "\t}\n"
        "}\n";
    cl_int err;
    cl_program prog = clCreateProgramWithSource(m_context,1,&source,NULL,&err);
    if( -err != CL_SUCCESS ) throw std::string("clCreateProgramWithSources");
    err = clBuildProgram(prog,0,NULL,"-cl-opt-disable",NULL,NULL);
    if( -err != CL_SUCCESS ) throw std::string("clBuildProgram(fromSources)");
    cl_kernel kernel = clCreateKernel(prog,"bufToImage",&err);
    checkOpenCLErr(err,"CreateKernel");

    cl_uint imageWidth = 80;
    cl_uint imageHeight = 90;
    //Initialize datas
    cl_uint maxVal = 255;
    cl_uint offsetX = 0;
    int size = imageWidth*imageHeight*3;
    int resSize = imageWidth*imageHeight*4;
    cl_uchar* data = new cl_uchar[size];
    cl_float* expectedData = new cl_float[resSize];
    for( int i = 0,j=0; i < size; i++,j++ ){
        data[i] = (cl_uchar)i;
        expectedData[j] = (cl_float)((unsigned char)i)/255.0f;
        if ( i%3 == 2 ){
            j++;
            expectedData[j] = 1.0f;
        }
    }
    cl_mem inBuffer = clCreateBuffer(m_context,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,size*sizeof(cl_uchar),data,&err);
    checkOpenCLErr(err, "clCreateBuffer()");

    clFinish(m_queue);
    cl_image_format imgFormat;
    imgFormat.image_channel_order = CL_RGBA;
    imgFormat.image_channel_data_type = CL_FLOAT;
    cl_mem outImg = clCreateImage2D( m_context, CL_MEM_READ_WRITE, &imgFormat, imageWidth, imageHeight, 0, NULL, &err );
    checkOpenCLErr(err,"get2DImage()");
    clFinish(m_queue);
    size_t kernelRegion[]={imageWidth,imageHeight};
    size_t kernelWorkgroup[]={1,1};
    //Fill kernel with data
    clSetKernelArg(kernel,0,sizeof(cl_mem),&inBuffer);
    clSetKernelArg(kernel,1,sizeof(cl_mem),&outImg);
    clSetKernelArg(kernel,2,sizeof(cl_uint),&offsetX);
    clSetKernelArg(kernel,3,sizeof(cl_uint),&imageWidth);
    clSetKernelArg(kernel,4,sizeof(cl_uint),&maxVal);

    //Run kernel
    err = clEnqueueNDRangeKernel(m_queue,kernel,2,NULL,kernelRegion,kernelWorkgroup,0,NULL,NULL);
    checkOpenCLErr(err,"RunKernel");
    clFinish(m_queue);
    //Check resulting data for validty
    cl_float* computedData = new cl_float[resSize];;
    size_t region[]={imageWidth,imageHeight,1};
    const size_t offset[] = {0,0,0};
    err = clEnqueueReadImage(m_queue,outImg,CL_TRUE,offset,region,0,0,computedData,0,NULL,NULL);
    checkOpenCLErr(err, "readDataFromImage()");
    clFinish(m_queue);


    for( int i = 0; i < resSize; i++ ){
        if( fabs(expectedData[i]-computedData[i])>0.1 ){
            std::cout << "Expected: \n";
            for( int j = 0; j < resSize; j++ ){
                std::cout << expectedData[j] << " ";
            }
            std::cout << "\nComputed: \n";
            std::cout << "\n";
            for( int j = 0; j < resSize; j++ ){
                std::cout << computedData[j] << " ";
            }
            std::cout << "\n";
            throw std::string("Error, computed and expected data are not the same!\n");
            }
        }

    }catch(std::string& e){
        std::cout << "\nCaught an exception: " << e << "\n";
        return 1;
    }
    std::cout << "Works fine\n";
    return 0;
}

Я также загрузил для вас исходный код, чтобы упростить его тестирование: http://www.file-upload.net/download-3524302/strangeOpenCLError.cpp.html

Пожалуйста, не могли бы вы сказать мне, если я сделал что-то не так? Есть ли ошибка в коде или это ошибка моего драйвера?

С уважением, Алекс

Редактировать: немного изменил программу (обе: здесь и связанную), чтобы повысить вероятность несоответствия.