Updating an image's values inside a kernel

OpenCL
1

So it seems like when passing an image2d_t as an argument to a kernel, it needs to be qualified as either __read_only or __write_only which is pretty limiting. What I want to be able to do is update the value in an image pixel instead of rewriting it completely.

So what I want to be able to do is something similar to this

const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                      CLK_ADDRESS_CLAMP_TO_EDGE   |
                      CLK_FILTER_NEAREST;

__kernel void updateImage(image2d_t image)
{
        size_t X = get_global_id(0);
        size_t Y = get_global_id(1);

        float4 oldVal = read_imagef(image, sampler, (int2)(X, Y));

        float4 newVal = (float)4(1,0,0,1);
        float4 result = oldVal + newVal;

        write_imagef(image, (int2)(X, Y), result);
}

So is there some sort of “update” function that I can use or anything instead of “write” that would allow me to achieve a similar functionality?

2

OpenCL 2.0 adds support for images with the read_write qualifier. It is not possible in OpenCL 1.x, you’ll need to use two different images. Note: It might just be faster that way anyway.

3

That’s good to know that OpenCL 2.0 adds in those features but to my knowledge there aren’t currently any CPU/GPUs on the market with with 2.0 driver support? So since I’m restricted to openCL 1.x for now, is the best way really to just use two separate images?

4

You have old knowledge. Intel and AMD are both shipping OpenCL 2.0 drivers.

Intel: OpenCL™ Runtimes for Intel® Processors (2014 r2 is OpenCL 2.0)

AMD: http://support.amd.com/en-us/kb-articles/Pages/OpenCL2-Driver.aspx

NVIDIA is still at OpenCL 1.1

Also, Intel and AMD drivers only provide OpenCL 2.0 on recent hardware (their older hardware will still be OpenCL 1.1 or 1.2).

For for broadest compatibility use two images. And like I said, it might be faster that way any way since an invariant source image can leverage the texture cache but a read/write image might not (depends on hardware design). You’d have to benchmark it to know for sure.