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):
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="")
