title: Simple OpenCL kernel produces vertical bars instead of solid color in output image. Why?
body: <p>I'm trying to write image processing OpenCL application, but my problem that any attempt to alter input image produces artifacts which look like vertical bars. This does not happen if I copy image pixels without altering them. So for example this line produces artifacts:</p>
  
  <pre><code>pixel = (uint4)(image1_pixel.x,
                  image1_pixel.y,
                  image1_pixel.z,
                  255);
  </code></pre>
  
  <p>...but this one works as expected:</p>
  
  <pre><code>pixel = (uint4)(image1_pixel.x,
                  image1_pixel.y,
                  image1_pixel.z,
                  image1_pixel.w);
  </code></pre>
  
  <p><a href="http://science.su/stuff/so/opencl/2012-10-16/input.png" rel="nofollow">Input</a> is opaque 32-bit PNG image, so I expect both code lines to produce the same result. In reality, however, only second line works as expected. First line gives <a href="http://science.su/stuff/so/opencl/2012-10-16/output.png" rel="nofollow">output with artifacts</a>.</p>
  
  <p>Here is my kernel:</p>
  
  <pre><code>__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                                CLK_ADDRESS_CLAMP |
                                CLK_FILTER_NEAREST;
  
  __kernel void test(__read_only image2d_t image1,
                    __write_only image2d_t out) {
    const int2 pos = (int2)(get_global_id(0), get_global_id(1) );
    uint4 image1_pixel = read_imageui(image1, sampler, pos);
    uint4 pixel = (uint4)(image1_pixel.x,
                          image1_pixel.y,
                          image1_pixel.z,
                          255);
    write_imageui(out, pos, pixel);
  }
  </code></pre>
  
  <p>Here is relevant portion of main.cpp code:</p>
  
  <pre><code>  CImg&lt;unsigned char&gt; image1("../input.png");
    ...
    Image2D clImage1 = Image2D(context,
      CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
      ImageFormat(CL_RGBA, CL_UNSIGNED_INT8),
      image1.width(), image1.height(), 0, image1.data() );
    Image2D clResult = Image2D(context, CL_MEM_WRITE_ONLY,
      ImageFormat(CL_RGBA, CL_UNSIGNED_INT8),
      image1.width(), image1.height(), 0, NULL);
    Kernel test = Kernel(program, "test");
    test.setArg(0, clImage1); test.setArg(1, clResult);
    Event kernel_event, read_event;
    queue.enqueueNDRangeKernel(test, NullRange,
      NDRange(image1.width(), image1.height() ),
      NullRange, NULL, &amp;kernel_event);
    cl::size_t&lt;3&gt; origin;
    origin.push_back(0); origin.push_back(0); origin.push_back(0);
    cl::size_t&lt;3&gt; region;
    region.push_back(image1.width() );
    region.push_back(image1.height() ); region.push_back(1);
    queue.enqueueReadImage(clResult, CL_TRUE,
                           origin, region, 0, 0,
                           image1.data(), NULL, NULL);
    kernel_event.wait();
    image1.save("../output.png");
  </code></pre>
  
  <p><a href="http://science.su/stuff/so/opencl/2012-10-16/test.rar" rel="nofollow">Here</a> can be downloaded full source code for my test application <em>(it contains short main.cpp under 30 lines, CMakeLists.txt, readme.txt explaining how to compile and run it, input image and the kernel)</em>. I use CImg library to load and save images. I double-checked that input opens as 32-bit RGBA image. I tried to run the kernel with AMD or NVidia SDK and got the same result.</p>
  
  <p>Any idea why I get unexpected result?</p>