Exemplo n.º 1
0
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))
Exemplo n.º 2
0
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))
Exemplo n.º 3
0
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
Exemplo n.º 4
0
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
Exemplo n.º 5
0
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()
Exemplo n.º 6
0
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()
Exemplo n.º 7
0
    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)
Exemplo n.º 8
0
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")
Exemplo n.º 9
0
 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)
Exemplo n.º 10
0
 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)
Exemplo n.º 11
0
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))
Exemplo n.º 12
0
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))
Exemplo n.º 13
0
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()
Exemplo n.º 14
0
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()
Exemplo n.º 15
0
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))
Exemplo n.º 16
0
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))
Exemplo n.º 17
0
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)
Exemplo n.º 18
0
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)
Exemplo n.º 19
0
    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)
Exemplo n.º 20
0
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)
Exemplo n.º 21
0
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)