From 59c89983c6594827538a6217165b7339c4f5123e Mon Sep 17 00:00:00 2001
From: Andreas Kloeckner <inform@tiker.net>
Date: Wed, 9 Sep 2020 13:39:26 -0500
Subject: [PATCH] Fix, test usage of read-only SVM

---
 pyopencl/__init__.py |  6 ++----
 test/test_wrapper.py | 20 ++++++++++++++++++++
 2 files changed, 22 insertions(+), 4 deletions(-)

diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py
index 4e2ef296..ef56ad0f 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 45e2d747..193f198f 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
-- 
GitLab