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)
def test_fine_grain_svm(ctx_factory): ctx = ctx_factory() # queue = cl.CommandQueue(ctx) if (ctx._get_cl_version() < (2, 0) or cl.get_cl_header_version() < (2, 0)): from pytest import skip skip("SVM only available in OpenCL 2.0 and higher") svm_ary = cl.fsvm_empty(ctx, (100, 100), np.float32, alignment=64) assert isinstance(svm_ary.base, cl.SVMAllocation)
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 pytest import skip if (ctx._get_cl_version() < (2, 0) or cl.get_cl_header_version() < (2, 0)): skip("SVM only available in OpenCL 2.0 and higher") if not (ctx.devices[0].svm_capabilities & cl.device_svm_capabilities.FINE_GRAIN_BUFFER): 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)
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 pytest import skip if (ctx._get_cl_version() < (2, 0) or cl.get_cl_header_version() < (2, 0)): skip("SVM only available in OpenCL 2.0 and higher") if not (ctx.devices[0].svm_capabilities & cl.device_svm_capabilities.FINE_GRAIN_BUFFER): 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)
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.") if has_fine_grain_system_svm(dev): print("Testing fine-grained system SVM...", end="")