Pyopencl write_imagef function does not seem to change an image2d_t object

Question

I've been trying to learn OpenCL and stumbled on a bit of a problem. In the code below I create an empty write_only opencl image object and try to get a simple kernel to turn in black (or at least change it in some fashion), however it just returns an empty image. (I was working on an image convolution exercise, which kept returning an empty image, the code below was just an attempt to isolate the problem.)

I've been messing around with it for over 2 hours now, I'm pretty sure I'm stuck.

import matplotlib.pyplot as plt
import scipy.ndimage as si
import pyopencl as cl
import numpy as np
import os

kernel = """
__kernel void black(__write_only image2d_t dst,
                    int rows,
                    int columns)
{
    const int column = get_global_id(0);
    const int row = get_global_id(1);

        if (column < columns && row < rows)
        {
            write_imagef(dst, (int2)(column, row),
                (float4)(1.0f, 1.0f, 1.0f, 1.0f));
        }   
}
"""

ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)

mf = cl.mem_flags

f = cl.ImageFormat(cl.channel_order.R, cl.channel_type.UNSIGNED_INT8)

dst_image = cl.Image(ctx, mf.WRITE_ONLY , f, shape=(100,100,4))

prg = cl.Program(ctx, kernel).build()

prg.black(queue, (100,100), None, dst_image,
          np.int32(100),
          np.int32(100))

postimage = np.zeros((100,100,4), dtype=np.uint8)
cl.enqueue_copy(queue, postimage, dst_image,
                origin=(0, 0, 0),
                region=(100,100,4))

plt.imshow(postimage)
plt.show()

The spec says "write_imageui can only be used with image objects created with image_channel_data_type set to one of the following values: CL_UNSIGNED_INT8, CL_UNSIGNED_INT16, or CL_UNSIGNED_INT32.". I don't see CL_UNORM_INT8 in the list. Try with write_imagef maybe. — Thomas

jprice jprice · Accepted Answer · 2014-01-05T11:04:05

Your kernel code is actually fine (although the code you've posted sets every pixel to white, so it would be hard to tell if it was working!). The issue is in the way you are creating the image. When you specify shape=(100,100,4) in your cl.Image constructor, you are actually asking for a 3D image.

The simplest way to get your example code to produce something useful is to modify these lines:

f = cl.ImageFormat(cl.channel_order.R, cl.channel_type.UNSIGNED_INT8)
....
dst_image = cl.Image(ctx, mf.WRITE_ONLY , f, shape=(100,100,4))
...
cl.enqueue_copy(queue, postimage, dst_image,
                origin=(0, 0, 0),
                region=(100,100,4))

to

f = cl.ImageFormat(cl.channel_order.RGBA, cl.channel_type.UNSIGNED_INT8)
....
dst_image = cl.Image(ctx, mf.WRITE_ONLY , f, shape=(100,100))
....
cl.enqueue_copy(queue, postimage, dst_image,
                origin=(0, 0),
                region=(100,100))

Here I've changed to image format to RGBA, which matches what the matplotlib image display function appears to expect (I'm not familiar with that library, so it might be that you can get it to display single channel images too). I've also modified the lines that create the image and copy it back to the host to be 2D, rather than 3D.

Now, if you change your kernel to write this data to the image:

write_imagef(dst, (int2)(column, row), (float4)(1.0f, 0.0f, 0.0f, 1.0f));

you should get a nice red output image!

Pyopencl write_imagef function does not seem to change an image2d_t object

1 Answers