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.")