From 7866e181847c8354e4905b73bf5adc49a590b934 Mon Sep 17 00:00:00 2001 From: Andreas Kloeckner Date: Sun, 21 Aug 2011 17:05:30 -0500 Subject: [PATCH] Fix 2D image creation. (reported by Sebastian Nowozin) --- pyopencl/__init__.py | 5 ++++- test/test_wrapper.py | 14 +++++++------- 2 files changed, 11 insertions(+), 8 deletions(-) diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py index 72f57e34..c505cf71 100644 --- a/pyopencl/__init__.py +++ b/pyopencl/__init__.py @@ -692,6 +692,9 @@ DTYPE_TO_CHANNEL_TYPE_NORM = { def image_from_array(ctx, ary, num_channels, mode="r", norm_int=False): # FIXME what about vector types? + if not ary.flags.c_contiguous: + raise ValueError("array must be C-contiguous") + if num_channels == 1: shape = ary.shape strides = ary.strides @@ -725,7 +728,7 @@ def image_from_array(ctx, ary, num_channels, mode="r", norm_int=False): return Image(ctx, mode_flags | mem_flags.COPY_HOST_PTR, ImageFormat(img_format, channel_type), - shape=shape, pitches=strides[:-1], + shape=shape[::-1], pitches=strides[::-1][1:], hostbuf=ary) # }}} diff --git a/test/test_wrapper.py b/test/test_wrapper.py index 306db603..32a4d0f4 100644 --- a/test/test_wrapper.py +++ b/test/test_wrapper.py @@ -234,23 +234,22 @@ class TestCL: __global float *dest, __read_only image2d_t src, sampler_t samp, - int width) + int stride0) { - int x = get_global_id(0); - int y = get_global_id(1); + int d0 = get_global_id(0); + int d1 = get_global_id(1); /* const sampler_t samp = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; */ - dest[x + width*y] = read_imagef(src, samp, (float2)(x, y)).x; - // dest[x + width*y] = get_image_height(src); + dest[d0*stride0 + d1] = read_imagef(src, samp, (float2)(d1, d0)).x; } """).build() num_channels = 1 - a = np.random.rand(1024, 1024, num_channels).astype(np.float32) + a = np.random.rand(1024, 512, num_channels).astype(np.float32) if num_channels == 1: a = a[:,:,0] @@ -261,7 +260,8 @@ class TestCL: samp = cl.Sampler(context, False, cl.addressing_mode.CLAMP, cl.filter_mode.NEAREST) - prg.copy_image(queue, a.shape, None, a_dest, a_img, samp, np.int32(a.shape[0])) + prg.copy_image(queue, a.shape, None, a_dest, a_img, samp, + np.int32(a.strides[0]/a.dtype.itemsize)) a_result = np.empty_like(a) cl.enqueue_copy(queue, a_result, a_dest) -- GitLab