def create_struct_types(word_dtype, key_words, counter_words):
key_dtype = dtypes.align(numpy.dtype([('v', (word_dtype, (key_words,)))]))
key_ctype = dtypes.ctype_module(key_dtype)
counter_dtype = dtypes.align(numpy.dtype([('v', (word_dtype, (counter_words,)))]))
counter_ctype = dtypes.ctype_module(counter_dtype)
return key_dtype, key_ctype, counter_dtype, counter_ctype
def create_struct_types(word_dtype, key_words, counter_words):
key_dtype = dtypes.align(numpy.dtype([('v', (word_dtype, (key_words, )))]))
key_ctype = dtypes.ctype_module(key_dtype)
counter_dtype = dtypes.align(
numpy.dtype([('v', (word_dtype, (counter_words, )))]))
counter_ctype = dtypes.ctype_module(counter_dtype)
return key_dtype, key_ctype, counter_dtype, counter_ctype
def test_structural_typing(some_thr):
"""
Checks that ``ctype_module`` for equal dtype objects result in the same module object
(which means that these two types will actually be rendered as a single type).
"""
dtype = dtypes.align(
numpy.dtype([('val1', numpy.int32), ('val2', numpy.float32)]))
struct1 = dtypes.ctype_module(dtype)
struct2 = dtypes.ctype_module(dtype)
# Check that these reference the same object
assert struct1 is struct2
# Just in case, compile a test kernel
program = some_thr.compile("""
KERNEL void test(GLOBAL_MEM float *dest)
{
const SIZE_T i = get_global_id(0);
${struct1} temp1;
${struct2} temp2;
temp1.val1 = 0;
temp1.val2 = 1;
// If struct1 and struct2 correspond to different types,
// this will give a compilation error,
// because C has a nominative typing system.
temp2 = temp1;
dest[i] = temp2.val1 + temp2.val2;
}
""",
render_kwds=dict(struct1=struct1,
struct2=struct2))
def test_structural_typing(some_thr):
"""
Checks that ``ctype_module`` for equal dtype objects result in the same module object
(which means that these two types will actually be rendered as a single type).
"""
dtype = dtypes.align(numpy.dtype([('val1', numpy.int32), ('val2', numpy.float32)]))
struct1 = dtypes.ctype_module(dtype)
struct2 = dtypes.ctype_module(dtype)
# Check that these reference the same object
assert struct1 is struct2
# Just in case, compile a test kernel
program = some_thr.compile(
"""
KERNEL void test(GLOBAL_MEM float *dest)
{
const SIZE_T i = get_global_id(0);
${struct1} temp1;
${struct2} temp2;
temp1.val1 = 0;
temp1.val2 = 1;
// If struct1 and struct2 correspond to different types,
// this will give a compilation error,
// because C has a nominative typing system.
temp2 = temp1;
dest[i] = temp2.val1 + temp2.val2;
}
""", render_kwds=dict(
struct1=struct1,
struct2=struct2))
def test_structure_type(thr):
shape = (100, 100)
dtype = dtypes.align(
numpy.dtype([('i1', numpy.uint32),
('nested', numpy.dtype([
('v', numpy.uint64),
])), ('i2', numpy.uint32)]))
a = get_test_array(shape, dtype)
a_dev = thr.to_device(a)
# Have to construct the resulting array manually,
# since numpy cannot reduce arrays with struct dtypes.
b_ref = numpy.empty(100, dtype)
b_ref['i1'] = a['i1'].sum(0)
b_ref['nested']['v'] = a['nested']['v'].sum(0)
b_ref['i2'] = a['i2'].sum(0)
predicate = Predicate(
Snippet.create(lambda v1, v2: """
${ctype} result = ${v1};
result.i1 += ${v2}.i1;
result.nested.v += ${v2}.nested.v;
result.i2 += ${v2}.i2;
return result;
""",
render_kwds=dict(ctype=dtypes.ctype_module(dtype))),
numpy.zeros(1, dtype)[0])
rd = Reduce(a_dev, predicate, axes=(0, ))
b_dev = thr.empty_like(rd.parameter.output)
rdc = rd.compile(thr)
rdc(b_dev, a_dev)
b_res = b_dev.get()
# Array.get() runs numpy.lib.stride_tricks.as_strided() on the array,
# which adds dummy variables instead of custom offsets (and our `dtype` has them),
# making the result dtype different, and failing the test.
# For now we will just note the difference and convert the result
# back to the original dtype (they are still compatible).
# When the behavior changes, the test will start to fail and we will notice.
# See inducer/compyte issue #26.
wrong_dtype = b_res.dtype != b_dev.dtype
b_res = b_res.astype(dtype)
assert diff_is_negligible(b_res, b_ref)
if wrong_dtype:
pytest.xfail("as_strided() still corrupts the datatype")
else:
pytest.fail("as_strided() does not corrupt the datatype anymore, "
"we can remove the `astype()` now")
def test_align(thr):
"""
Test the correctness of alignment for field offsets adjusted automatically.
"""
dtype_nested = numpy.dtype([('val1', numpy.int32), ('pad', numpy.int8)])
dtype = numpy.dtype([('val1', numpy.int32), ('val2', numpy.int16),
('nested', dtype_nested)])
dtype = dtypes.align(dtype)
check_struct_fill(thr, dtype)
def test_structural_typing_nested(some_thr):
"""
Check that the structural typing behavior works for nested structures.
In other words, a nested dtype gets represented by the same module as
an equal top-level dtype.
"""
dtype_nested = dtypes.align(
numpy.dtype([('val1', numpy.int32), ('val2', numpy.float32)]))
dtype = dtypes.align(numpy.dtype([
('val1', numpy.int32), ('val2', numpy.float32),
('nested', dtype_nested)]))
struct_nested = dtypes.ctype_module(dtype_nested)
struct = dtypes.ctype_module(dtype)
program = some_thr.compile(
"""
KERNEL void test(GLOBAL_MEM float *dest)
{
const SIZE_T i = get_global_id(0);
${struct_nested} temp_nested;
${struct} temp;
temp_nested.val1 = 0;
temp_nested.val2 = 1;
// If the nested structure has a different type from temp_nested,
// this will give a compilation error.
temp.nested = temp_nested;
temp.val1 = 0;
temp.val2 = 1;
dest[i] = temp.val1 + temp.val2 + temp.nested.val1 + temp.nested.val2;
}
""", render_kwds=dict(
struct=struct,
struct_nested=struct_nested))
def test_structural_typing_nested(some_thr):
"""
Check that the structural typing behavior works for nested structures.
In other words, a nested dtype gets represented by the same module as
an equal top-level dtype.
"""
dtype_nested = dtypes.align(
numpy.dtype([('val1', numpy.int32), ('val2', numpy.float32)]))
dtype = dtypes.align(
numpy.dtype([('val1', numpy.int32), ('val2', numpy.float32),
('nested', dtype_nested)]))
struct_nested = dtypes.ctype_module(dtype_nested)
struct = dtypes.ctype_module(dtype)
program = some_thr.compile("""
KERNEL void test(GLOBAL_MEM float *dest)
{
const SIZE_T i = get_global_id(0);
${struct_nested} temp_nested;
${struct} temp;
temp_nested.val1 = 0;
temp_nested.val2 = 1;
// If the nested structure has a different type from temp_nested,
// this will give a compilation error.
temp.nested = temp_nested;
temp.val1 = 0;
temp.val2 = 1;
dest[i] = temp.val1 + temp.val2 + temp.nested.val1 + temp.nested.val2;
}
""",
render_kwds=dict(struct=struct,
struct_nested=struct_nested))
def test_nested_array(thr):
"""
Check that structures with nested arrays are processed correctly.
"""
dtype_nested = numpy.dtype(
dict(names=['val1', 'pad'], formats=[numpy.int8, numpy.int8]))
dtype = numpy.dtype(
dict(names=['pad', 'struct_arr', 'regular_arr'],
formats=[
numpy.int32,
numpy.dtype((dtype_nested, 2)),
numpy.dtype((numpy.int16, 3))
]))
dtype = dtypes.align(dtype)
struct = dtypes.ctype_module(dtype)
program = thr.compile("""
KERNEL void test(GLOBAL_MEM ${struct} *dest)
{
const SIZE_T i = get_global_id(0);
${struct} res;
res.struct_arr[0].val1 = i + 0;
res.struct_arr[1].val1 = i + 1;
res.regular_arr[0] = i + 2;
res.regular_arr[1] = i + 3;
res.regular_arr[2] = i + 4;
dest[i] = res;
}
""",
render_kwds=dict(struct=struct, dtype=dtype))
test = program.test
a_dev = thr.array(64, dtype)
test(a_dev, global_size=64)
a = a_dev.get()
idxs = numpy.arange(64)
assert (a['struct_arr'][:, 0]['val1'] == idxs + 0).all()
assert (a['struct_arr'][:, 1]['val1'] == idxs + 1).all()
assert (a['regular_arr'][:, 0] == idxs + 2).all()
assert (a['regular_arr'][:, 1] == idxs + 3).all()
assert (a['regular_arr'][:, 2] == idxs + 4).all()
def test_align(thr):
"""
Test the correctness of alignment for field offsets adjusted automatically.
"""
dtype_nested = numpy.dtype([
('val1', numpy.int32), ('pad', numpy.int8)])
dtype = numpy.dtype([
('val1', numpy.int32),
('val2', numpy.int16),
('nested', dtype_nested)])
dtype = dtypes.align(dtype)
check_struct_fill(thr, dtype)
def test_nested_array(thr):
"""
Check that structures with nested arrays are processed correctly.
"""
dtype_nested = numpy.dtype(dict(
names=['val1', 'pad'],
formats=[numpy.int8, numpy.int8]))
dtype = numpy.dtype(dict(
names=['pad', 'struct_arr', 'regular_arr'],
formats=[numpy.int32, numpy.dtype((dtype_nested, 2)), numpy.dtype((numpy.int16, 3))]))
dtype = dtypes.align(dtype)
struct = dtypes.ctype_module(dtype)
program = thr.compile(
"""
KERNEL void test(GLOBAL_MEM ${struct} *dest)
{
const SIZE_T i = get_global_id(0);
${struct} res;
res.struct_arr[0].val1 = i + 0;
res.struct_arr[1].val1 = i + 1;
res.regular_arr[0] = i + 2;
res.regular_arr[1] = i + 3;
res.regular_arr[2] = i + 4;
dest[i] = res;
}
""", render_kwds=dict(
struct=struct,
dtype=dtype))
test = program.test
a_dev = thr.array(64, dtype)
test(a_dev, global_size=64)
a = a_dev.get()
idxs = numpy.arange(64)
assert (a['struct_arr'][:,0]['val1'] == idxs + 0).all()
assert (a['struct_arr'][:,1]['val1'] == idxs + 1).all()
assert (a['regular_arr'][:,0] == idxs + 2).all()
assert (a['regular_arr'][:,1] == idxs + 3).all()
assert (a['regular_arr'][:,2] == idxs + 4).all()
def test_structure_type(thr):
shape = (100, 100)
dtype = dtypes.align(numpy.dtype([
('i1', numpy.uint32),
('nested', numpy.dtype([
('v', numpy.uint64),
])),
('i2', numpy.uint32)
]))
a = get_test_array(shape, dtype)
a_dev = thr.to_device(a)
# Have to construct the resulting array manually,
# since numpy cannot reduce arrays with struct dtypes.
b_ref = numpy.empty(100, dtype)
b_ref['i1'] = a['i1'].sum(0)
b_ref['nested']['v'] = a['nested']['v'].sum(0)
b_ref['i2'] = a['i2'].sum(0)
predicate = Predicate(
Snippet.create(lambda v1, v2: """
${ctype} result = ${v1};
result.i1 += ${v2}.i1;
result.nested.v += ${v2}.nested.v;
result.i2 += ${v2}.i2;
return result;
""",
render_kwds=dict(
ctype=dtypes.ctype_module(dtype))),
numpy.zeros(1, dtype)[0])
rd = Reduce(a_dev, predicate, axes=(0,))
b_dev = thr.empty_like(rd.parameter.output)
rdc = rd.compile(thr)
rdc(b_dev, a_dev)
b_res = b_dev.get()
assert diff_is_negligible(b_res, b_ref)
def test_scan_structure_type(thr, exclusive):
shape = (100, 100)
dtype = dtypes.align(
numpy.dtype([('i1', numpy.uint32),
('nested', numpy.dtype([
('v', numpy.uint64),
])), ('i2', numpy.uint32)]))
a = get_test_array(shape, dtype)
a_dev = thr.to_device(a)
# Have to construct the resulting array manually,
# since numpy cannot scan arrays with struct dtypes.
b_ref = numpy.empty(shape, dtype)
b_ref['i1'] = ref_scan(a['i1'], axes=0, exclusive=exclusive)
b_ref['nested']['v'] = ref_scan(a['nested']['v'],
axes=0,
exclusive=exclusive)
b_ref['i2'] = ref_scan(a['i2'], axes=0, exclusive=exclusive)
predicate = Predicate(
Snippet.create(lambda v1, v2: """
${ctype} result = ${v1};
result.i1 += ${v2}.i1;
result.nested.v += ${v2}.nested.v;
result.i2 += ${v2}.i2;
return result;
""",
render_kwds=dict(ctype=dtypes.ctype_module(dtype))),
numpy.zeros(1, dtype)[0])
scan = Scan(a_dev, predicate, axes=(0, ), exclusive=exclusive)
b_dev = thr.empty_like(scan.parameter.output)
scanc = scan.compile(thr)
scanc(b_dev, a_dev)
b_res = b_dev.get()
assert diff_is_negligible(b_res, b_ref)
def test_align(thr, dtype_to_align):
aligned_dtype = dtypes.align(dtype_to_align)
empyric_dtype = get_offsets_from_device(thr, dtype_to_align)
assert aligned_dtype == empyric_dtype
from reikna.cluda import dtypes, any_api
from reikna.algorithms import Reduce, Predicate
from reikna.cluda import Snippet
from reikna.core import Annotation, Type, Transformation, Parameter
# Pick the first available GPGPU API and make a Thread on it.
api = any_api()
thr = api.Thread.create()
# Minmax data type and the corresponding structure.
# Note that the dtype must be aligned before use on a GPU.
mmc_dtype = dtypes.align(numpy.dtype([
("cur_min", numpy.int32),
("cur_max", numpy.int32),
("pad", numpy.int32),
]))
mmc_c_decl = dtypes.ctype_module(mmc_dtype)
# Create the "empty" element for our minmax monoid, that is
# x `minmax` empty == empty `minmax` x == x.
empty = numpy.empty(1, mmc_dtype)[0]
empty["cur_min"] = 1 << 30
empty["cur_max"] = -(1 << 30)
# Reduction predicate for the minmax.
# v1 and v2 get the names of two variables to be processed.
predicate = Predicate(
def test_align(thr, dtype_to_align):
aligned_dtype = dtypes.align(dtype_to_align)
empyric_dtype = get_offsets_from_device(thr, dtype_to_align)
assert aligned_dtype == empyric_dtype
import numpy
from reikna.cluda import dtypes, any_api
from reikna.algorithms import Reduce, Predicate
from reikna.cluda import Snippet
from reikna.core import Annotation, Type, Transformation, Parameter
# Pick the first available GPGPU API and make a Thread on it.
api = any_api()
thr = api.Thread.create()
# Minmax data type and the corresponding structure.
# Note that the dtype must be aligned before use on a GPU.
mmc_dtype = dtypes.align(
numpy.dtype([
("cur_min", numpy.int32),
("cur_max", numpy.int32),
("pad", numpy.int32),
]))
mmc_c_decl = dtypes.ctype_module(mmc_dtype)
# Create the "empty" element for our minmax monoid, that is
# x `minmax` empty == empty `minmax` x == x.
empty = numpy.empty(1, mmc_dtype)[0]
empty["cur_min"] = 1 << 30
empty["cur_max"] = -(1 << 30)
# Reduction predicate for the minmax.
# v1 and v2 get the names of two variables to be processed.
predicate = Predicate(
Snippet.create(lambda v1, v2: """
${ctype} result = ${v1};
