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 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 check_struct_fill(thr, dtype):
"""
Fill every field of the given ``dtype`` with its number and check the results.
This helps to detect issues with offsets in the struct.
"""
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;
%for i, field_info in enumerate(dtypes.flatten_dtype(dtype)):
res.${".".join(field_info[0])} = ${i};
%endfor
dest[i] = res;
}
""",
render_kwds=dict(struct=struct, dtype=dtype))
test = program.test
a_dev = thr.array(128, dtype)
test(a_dev, global_size=128)
a = a_dev.get()
for i, field_info in enumerate(dtypes.flatten_dtype(dtype)):
path, _ = field_info
assert (dtypes.extract_field(a, path) == i).all()
def check_struct_fill(thr, dtype):
"""
Fill every field of the given ``dtype`` with its number and check the results.
This helps to detect issues with offsets in the struct.
"""
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;
%for i, field_info in enumerate(dtypes.flatten_dtype(dtype)):
res.${".".join(field_info[0])} = ${i};
%endfor
dest[i] = res;
}
""", render_kwds=dict(
struct=struct,
dtype=dtype))
test = program.test
a_dev = thr.array(128, dtype)
test(a_dev, global_size=128)
a = a_dev.get()
for i, field_info in enumerate(dtypes.flatten_dtype(dtype)):
path, _ = field_info
assert (dtypes.extract_field(a, path) == i).all()
def __init__(self, dtype, shape=None, strides=None, offset=0, nbytes=None):
self.shape = tuple() if shape is None else wrap_in_tuple(shape)
self.size = product(self.shape)
self.dtype = dtypes.normalize_type(dtype)
self.ctype = dtypes.ctype_module(self.dtype)
default_strides = helpers.default_strides(self.shape,
self.dtype.itemsize)
if strides is None:
strides = default_strides
else:
strides = tuple(strides)
self._default_strides = strides == default_strides
self.strides = strides
default_nbytes = helpers.min_buffer_size(self.shape,
self.dtype.itemsize,
self.strides)
if nbytes is None:
nbytes = default_nbytes
self._default_nbytes = nbytes == default_nbytes
self.nbytes = nbytes
self.offset = offset
self._cast = dtypes.cast(self.dtype)
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 __init__(self, dtype, shape=None, strides=None):
self.shape = tuple() if shape is None else wrap_in_tuple(shape)
self.size = product(self.shape)
self.dtype = dtypes.normalize_type(dtype)
self.ctype = dtypes.ctype_module(self.dtype)
if strides is None:
self.strides = tuple([
self.dtype.itemsize * product(self.shape[i+1:]) for i in range(len(self.shape))])
else:
self.strides = strides
self._cast = dtypes.cast(self.dtype)
def __init__(self, dtype, shape=None, strides=None):
self.shape = tuple() if shape is None else wrap_in_tuple(shape)
self.size = product(self.shape)
self.dtype = dtypes.normalize_type(dtype)
self.ctype = dtypes.ctype_module(self.dtype)
if strides is None:
self.strides = tuple([
self.dtype.itemsize * product(self.shape[i + 1:])
for i in range(len(self.shape))
])
else:
self.strides = strides
self._cast = dtypes.cast(self.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_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 get_offsets_from_device(thr, dtype):
"""
Returns a new data type object with the same fields as in ``dtype``
and the field offsets set by the compiler in ``thr``.
All existing offsets in ``dtype`` are ignored.
"""
if dtype.names is None:
return dtype
struct = dtypes.ctype_module(dtype, ignore_alignment=True)
program = thr.compile(
"""
#define my_offsetof(type, field) ((size_t)(&((type *)0)->field))
KERNEL void test(GLOBAL_MEM int *dest)
{
%for i, name in enumerate(dtype.names):
dest[${i}] = my_offsetof(${struct}, ${name});
%endfor
dest[${len(dtype.names)}] = sizeof(${struct});
}
""", render_kwds=dict(
struct=struct,
dtype=dtype))
offsets = thr.array(len(dtype.names) + 1, numpy.int32)
test = program.test
test(offsets, global_size=1)
offsets = offsets.get()
# Casting to Python ints, becase numpy ints as dtype offsets make it unhashable.
offsets = [int(offset) for offset in offsets]
aligned_dtypes = [
get_offsets_from_device(thr, dtype.fields[name][0])
for name in dtype.names]
return numpy.dtype(dict(
names=dtype.names,
formats=aligned_dtypes,
offsets=offsets[:-1],
itemsize=offsets[-1],
aligned=True))
def get_offsets_from_device(thr, dtype):
"""
Returns a new data type object with the same fields as in ``dtype``
and the field offsets set by the compiler in ``thr``.
All existing offsets in ``dtype`` are ignored.
"""
if dtype.names is None:
return dtype
struct = dtypes.ctype_module(dtype, ignore_alignment=True)
program = thr.compile("""
#define my_offsetof(type, field) ((size_t)(&((type *)0)->field))
KERNEL void test(GLOBAL_MEM int *dest)
{
%for i, name in enumerate(dtype.names):
dest[${i}] = my_offsetof(${struct}, ${name});
%endfor
dest[${len(dtype.names)}] = sizeof(${struct});
}
""",
render_kwds=dict(struct=struct, dtype=dtype))
offsets = thr.array(len(dtype.names) + 1, numpy.int32)
test = program.test
test(offsets, global_size=1)
offsets = offsets.get()
# Casting to Python ints, becase numpy ints as dtype offsets make it unhashable.
offsets = [int(offset) for offset in offsets]
aligned_dtypes = [
get_offsets_from_device(thr, dtype.fields[name][0])
for name in dtype.names
]
return numpy.dtype(
dict(names=dtype.names,
formats=aligned_dtypes,
offsets=offsets[:-1],
itemsize=offsets[-1],
aligned=True))
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 __init__(self, dtype, shape=None, strides=None, offset=0, nbytes=None):
self.shape = tuple() if shape is None else wrap_in_tuple(shape)
self.size = product(self.shape)
self.dtype = dtypes.normalize_type(dtype)
self.ctype = dtypes.ctype_module(self.dtype)
default_strides = helpers.default_strides(self.shape, self.dtype.itemsize)
if strides is None:
strides = default_strides
else:
strides = tuple(strides)
self._default_strides = strides == default_strides
self.strides = strides
default_nbytes = helpers.min_buffer_size(self.shape, self.dtype.itemsize, self.strides)
if nbytes is None:
nbytes = default_nbytes
self._default_nbytes = nbytes == default_nbytes
self.nbytes = nbytes
self.offset = offset
self._cast = dtypes.cast(self.dtype)
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};
if (${v2}.cur_min < result.cur_min)
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};
if (${v2}.cur_min < result.cur_min)
result.cur_min = ${v2}.cur_min;
if (${v2}.cur_max > result.cur_max)