def test_fine_grain_svm(ctx_factory): import sys is_pypy = '__pypy__' in sys.builtin_module_names ctx = ctx_factory() queue = cl.CommandQueue(ctx) from pyopencl.characterize import has_fine_grain_buffer_svm from pytest import skip if not has_fine_grain_buffer_svm(queue.device): skip("device does not support fine-grain SVM") n = 3000 ary = cl.fsvm_empty(ctx, n, np.float32, alignment=64) if not is_pypy: # https://bitbucket.org/pypy/numpy/issues/52 assert isinstance(ary.base, cl.SVMAllocation) ary.fill(17) orig_ary = ary.copy() prg = cl.Program( ctx, """ __kernel void twice(__global float *a_g) { a_g[get_global_id(0)] *= 2; } """).build() prg.twice(queue, ary.shape, None, cl.SVM(ary)) queue.finish() print(ary) assert np.array_equal(orig_ary * 2, ary)
svm_ary = cl.SVM(cl.csvm_empty(ctx, 10, np.float32)) assert isinstance(svm_ary.mem, np.ndarray) with svm_ary.map_rw(queue) as ary: ary.fill(17) # use from host orig_ary = ary.copy() prg.twice(queue, svm_ary.mem.shape, None, svm_ary) queue.finish() with svm_ary.map_ro(queue) as ary: assert(np.array_equal(orig_ary*2, ary)) print(" done.") if has_fine_grain_buffer_svm(dev): print("Testing fine-grained buffer SVM...", end="") ary = cl.fsvm_empty(ctx, 10, np.float32) assert isinstance(ary.base, cl.SVMAllocation) ary.fill(17) orig_ary = ary.copy() prg.twice(queue, ary.shape, None, cl.SVM(ary)) queue.finish() assert np.array_equal(orig_ary*2, ary) print(" done.")