diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py
index 72f57e34a2768ec187785b8d7d34f905ff6b2ea8..c505cf71b84a89d70e4533f6db7af77e23890aa7 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 306db6032fe48d0c374362c0728975218c537369..32a4d0f4857c452aecade9083958a0a42b798ff1 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)