Add pyopencl.image_from_array.

pyopencl/__init__.py

@@ -642,6 +642,72 @@ def enqueue_copy(queue, dest, src, **kwargs):
 
 # }}}
 
+# {{{ image creation
+
+DTYPE_TO_CHANNEL_TYPE = {
+    np.dtype(np.float32): channel_type.FLOAT,
+    np.dtype(np.int16): channel_type.SIGNED_INT16,
+    np.dtype(np.int32): channel_type.SIGNED_INT32,
+    np.dtype(np.int8): channel_type.SIGNED_INT8,
+    np.dtype(np.uint16): channel_type.UNSIGNED_INT16,
+    np.dtype(np.uint32): channel_type.UNSIGNED_INT32,
+    np.dtype(np.uint8): channel_type.UNSIGNED_INT8,
+    }
+try:
+    np.float16
+except:
+    pass
+else:
+    DTYPE_TO_CHANNEL_TYPE[np.dtype(np.float16)] = channel_type.HALF_FLOAT,
+
+DTYPE_TO_CHANNEL_TYPE_NORM = {
+    np.dtype(np.int16): channel_type.SNORM_INT16,
+    np.dtype(np.int8): channel_type.SNORM_INT8,
+    np.dtype(np.int16): channel_type.UNORM_INT16,
+    np.dtype(np.uint8): channel_type.UNORM_INT8,
+    }
+
+def image_from_array(ctx, ary, num_channels, mode="r", norm_int=False):
+    # FIXME what about vector types?
+
+    if num_channels == 1:
+        shape = ary.shape
+        strides = ary.strides
+    else:
+        if ary.shape[-1] != num_channels:
+            raise RuntimeError("last dimension must be equal to number of channels")
+
+        shape = ary.shape[:-1]
+        strides = ary.strides[:-1]
+
+    if mode == "r":
+        mode_flags = mem_flags.READ_ONLY
+    elif mode == "r":
+        mode_flags = mem_flags.WRITE_ONLY
+    else:
+        raise ValueError("invalid value '%s' for 'mode'" % mode)
+
+    img_format = {
+            1: channel_order.R,
+            2: channel_order.RG,
+            3: channel_order.RGB,
+            4: channel_order.RGBA,
+            }[num_channels]
+
+    assert ary.strides[-1] == ary.dtype.itemsize
+
+    if norm_int:
+        channel_type = DTYPE_TO_CHANNEL_TYPE_NORM[ary.dtype]
+    else:
+        channel_type = DTYPE_TO_CHANNEL_TYPE[ary.dtype]
+
+    return Image(ctx, mode_flags | mem_flags.COPY_HOST_PTR,
+            ImageFormat(img_format, channel_type),
+            shape=shape, pitches=strides[:-1],
+            hostbuf=ary)
+
+# }}}
+
 
 
 

pyopencl/array.py

@@ -1093,7 +1093,4 @@ subset_max = _make_subset_minmax_kernel("max")
 
 # }}}
 
-
-
-
 # vim: foldmethod=marker

test/test_wrapper.py

@@ -220,16 +220,18 @@ class TestCL:
         cl.enqueue_read_buffer(queue, a_buf, a_result).wait()
 
     @pytools.test.mark_test.opencl
-    def test_image_2d(self, device, ctx_factory):
+    def test_image_2d(self, ctx_factory):
         context = ctx_factory()
 
+        device, = context.devices
+
         if not device.image_support:
             from py.test import skip
             skip("images not supported on %s" % device)
 
         prg = cl.Program(context, """
             __kernel void copy_image(
-              __global float4 *dest,
+              __global float *dest,
               __read_only image2d_t src,
               sampler_t samp,
               int width)
@@ -242,18 +244,19 @@ class TestCL:
                 | CLK_ADDRESS_CLAMP
                 | CLK_FILTER_NEAREST;
                 */
-              dest[x + width*y] = read_imagef(src, samp, (float2)(x, y));
+              dest[x + width*y] = read_imagef(src, samp, (float2)(x, y)).x;
               // dest[x + width*y] = get_image_height(src);
             }
             """).build()
 
-        a = np.random.rand(1024, 1024, 4).astype(np.float32)
+        num_channels = 1
+        a = np.random.rand(1024, 1024, num_channels).astype(np.float32)
+        if num_channels == 1:
+            a = a[:,:,0]
+
         queue = cl.CommandQueue(context)
-        mf = cl.mem_flags
-        a_img = cl.Image(context, mf.READ_ONLY | mf.COPY_HOST_PTR,
-                cl.ImageFormat(cl.channel_order.RGBA, cl.channel_type.FLOAT),
-                shape=a.shape[:2], hostbuf=a)
-        a_dest = cl.Buffer(context, mf.READ_WRITE, a.nbytes)
+        a_img = cl.image_from_array(context, a, num_channels)
+        a_dest = cl.Buffer(context, cl.mem_flags.READ_WRITE, a.nbytes)
 
         samp = cl.Sampler(context, False,
                 cl.addressing_mode.CLAMP,
@@ -261,8 +264,7 @@ class TestCL:
         prg.copy_image(queue, a.shape, None, a_dest, a_img, samp, np.int32(a.shape[0]))
 
         a_result = np.empty_like(a)
-        cl.enqueue_read_buffer(queue, a_dest, a_result, is_blocking=True)
-        print(a_result.dtype)
+        cl.enqueue_copy(queue, a_result, a_dest)
 
         assert la.norm(a_result - a) == 0