1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
|
How-tos
=======
How to use struct types with PyOpenCL
-------------------------------------
We import and initialize PyOpenCL as usual:
.. doctest::
:options: +ELLIPSIS
>>> import numpy as np
>>> import pyopencl as cl
>>> import pyopencl.tools
>>> import pyopencl.array
>>> ctx = cl.create_some_context(interactive=False)
>>> queue = cl.CommandQueue(ctx)
Then, suppose we would like to declare a struct consisting of an integer and a
floating point number. We first create a :class:`numpy.dtype` along these
lines:
.. doctest::
>>> my_struct = np.dtype([("field1", np.int32), ("field2", np.float32)])
>>> print(my_struct)
[('field1', '<i4'), ('field2', '<f4')]
.. note::
Not all :mod:`numpy` dtypes are supported yet. For example strings (and
generally things that have a shape of their own) are not supported.
Since OpenCL C may have a different opinion for :mod:`numpy` on how the struct
should be laid out, for example because of `alignment
<https://en.wikipedia.org/wiki/Data_structure_alignment>`_. So as a first step, we
match our dtype against CL's version:
.. doctest::
>>> my_struct, my_struct_c_decl = cl.tools.match_dtype_to_c_struct(
... ctx.devices[0], "my_struct", my_struct)
>>> print(my_struct_c_decl)
typedef struct {
int field1;
float field2;
} my_struct;
<BLANKLINE>
<BLANKLINE>
We then tell PyOpenCL about our new type.
.. doctest::
>>> my_struct = cl.tools.get_or_register_dtype("my_struct", my_struct)
Next, we can create some data of that type on the host and transfer it to
the device:
.. doctest::
>>> ary_host = np.empty(20, my_struct)
>>> ary_host["field1"].fill(217)
>>> ary_host["field2"].fill(1000)
>>> ary_host[13]["field2"] = 12
>>> print(ary_host)
[(217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0)
(217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0)
(217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 12.0) (217, 1000.0)
(217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0) (217, 1000.0)]
>>> ary = cl.array.to_device(queue, ary_host)
We can then operate on the array with our own kernels:
.. doctest::
>>> prg = cl.Program(ctx, my_struct_c_decl + """
... __kernel void set_to_1(__global my_struct *a)
... {
... a[get_global_id(0)].field1 = 1;
... }
... """).build()
>>> evt = prg.set_to_1(queue, ary.shape, None, ary.data)
>>> print(ary)
[(1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0)
(1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0)
(1, 1000.0) (1, 12.0) (1, 1000.0) (1, 1000.0) (1, 1000.0) (1, 1000.0)
(1, 1000.0) (1, 1000.0)]
as well as with PyOpenCL's built-in operations:
.. doctest::
>>> from pyopencl.elementwise import ElementwiseKernel
>>> elwise = ElementwiseKernel(ctx, "my_struct *a", "a[i].field1 = 2;",
... preamble=my_struct_c_decl)
>>> evt = elwise(ary)
>>> print(ary)
[(2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0)
(2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0)
(2, 1000.0) (2, 12.0) (2, 1000.0) (2, 1000.0) (2, 1000.0) (2, 1000.0)
(2, 1000.0) (2, 1000.0)]
|
