diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py
index 4e2ef296eea79b54d3ececefcfda9c7a8e40ad66..ef56ad0fd8782286e6d5f35b331e99b3df1b82f5 100644
--- a/pyopencl/__init__.py
+++ b/pyopencl/__init__.py
@@ -1114,10 +1114,8 @@ def _add_functionality():
         """
         svmallocation_old_init(self, ctx, size, alignment, flags)
 
-        read_write = (
-                flags & mem_flags.WRITE_ONLY != 0
-                or flags & mem_flags.READ_WRITE != 0)
-
+        # mem_flags.READ_ONLY applies to kernels, not the host
+        read_write = True
         _interface["data"] = (
                 int(self._ptr_as_int()), not read_write)
 
diff --git a/test/test_wrapper.py b/test/test_wrapper.py
index 45e2d7476bf2939bcdee66c487c412667183e831..193f198fa8d305780b7bac2dd252c0eeba9b7f91 100644
--- a/test/test_wrapper.py
+++ b/test/test_wrapper.py
@@ -1030,6 +1030,26 @@ def test_coarse_grain_svm(ctx_factory):
         cl.enqueue_copy(queue, new_ary, svm_ary)
         assert np.array_equal(orig_ary*2, new_ary)
 
+    # {{{ https://github.com/inducer/pyopencl/issues/372
+
+    svm_buf_arr = cl.svm_empty(ctx, cl.svm_mem_flags.READ_ONLY, 10, np.int32)
+    svm_out_arr = cl.svm_empty(ctx, cl.svm_mem_flags.READ_WRITE, 10, np.int32)
+
+    with cl.SVM(svm_buf_arr).map_rw(queue) as ary:
+        ary.fill(17)
+
+    prg_ro = cl.Program(ctx, """
+        __kernel void twice_ro(__global int *out_g, __global int *in_g)
+        {
+          out_g[get_global_id(0)] = 2*in_g[get_global_id(0)];
+        }
+        """).build()
+
+    prg_ro.twice_ro(queue, svm_buf_arr.shape, None,
+            cl.SVM(svm_out_arr), cl.SVM(svm_buf_arr))
+
+    # }}}
+
 
 def test_fine_grain_svm(ctx_factory):
     import sys