File: low_level.py

package info (click to toggle)
compyle 0.8.1-11
  • links: PTS, VCS
  • area: main
  • in suites: forky, sid, trixie
  • size: 1,100 kB
  • sloc: python: 12,337; makefile: 21
file content (376 lines) | stat: -rw-r--r-- 11,419 bytes parent folder | download | duplicates (2)
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
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
"""Low level utility code. The intention is for users to use these but with the
knowledge that these are not general cross-backend tools but rather specific
tools.

"""

import re
import inspect

import numpy as np

from .config import get_config
from .array import Array, get_backend
from .transpiler import Transpiler
from .types import KnownType, ctype_to_dtype
from .extern import Extern
from .profile import profile


LID_0 = LDIM_0 = GDIM_0 = GID_0 = 0


def local_barrier():
    """Dummy method to keep Python happy.

    This is a valid function in OpenCL but has no meaning in Python for now.
    """
    pass


class LocalMem(object):
    '''A local memory specification for a GPU kernel.

    An example illustrates this best::

       >>> l = LocalMem(2)
       >>> m = l.get('double', 128)
       >>> m.size
       2048

    Note that this is basically ``sizeof(double) * 128 * 2``
    '''

    def __init__(self, size, backend=None):
        '''
        Constructor

        Parameters
        ----------

        size: int: a multiple of the current work group size.
        baackend: str: one of 'opencl', 'cuda'
        '''
        self.backend = get_backend(backend)
        if backend == 'cython':
            raise NotImplementedError(
                'LocalMem is only meaningful for the opencl/cuda backends.'
            )
        self.size = size
        self._cache = {}

    def get(self, c_type, workgroup_size):
        """Return the local memory required given the type and work group size.
        """
        key = (c_type, workgroup_size)
        if key in self._cache:
            return self._cache[key]
        elif self.backend == 'opencl':
            import pyopencl as cl
            dtype = ctype_to_dtype(c_type)
            sz = dtype.itemsize
            mem = cl.LocalMemory(sz * self.size * workgroup_size)
            self._cache[key] = mem
            return mem
        else:
            raise NotImplementedError(
                'Backend %s not implemented' % self.backend
            )


def splay_cl(queue, n, kernel_specific_max_wg_size=None):
    dev = queue.device
    max_work_items = min(128, dev.max_work_group_size)

    if kernel_specific_max_wg_size is not None:
        max_work_items = min(max_work_items, kernel_specific_max_wg_size)

    min_work_items = min(64, max_work_items)
    full_groups = dev.max_compute_units * 4 * 8
    # 4 to overfill the device
    # 8 is an Nvidia constant--that's how many
    # groups fit onto one compute device

    if n < min_work_items:
        group_count = 1
        work_items_per_group = min_work_items
    elif n < (full_groups * min_work_items):
        group_count = (n + min_work_items - 1) // min_work_items
        work_items_per_group = min_work_items
    elif n < (full_groups * max_work_items):
        group_count = full_groups
        grp = (n + min_work_items - 1) // min_work_items
        work_items_per_group = (
            (grp + full_groups - 1) // full_groups) * min_work_items
    else:
        group_count = (n + max_work_items - 1) // max_work_items
        work_items_per_group = max_work_items

    return (group_count * work_items_per_group,), (work_items_per_group,)


class Kernel(object):
    """A simple abstraction to create GPU kernels with pure Python.

    This will not work currently with the Cython backend.

    The idea is that one can create a Python function with suitable type
    annotations along with standard names from the CLUDA header (`LDIM_0,
    LID_0, GID_0, local_barrier()`, )etc.) to write kernels in pure Python.

    Note
    ----

    This works best with functions with annotations via the @annotate decorator
    or with function annotation as we need the type information for some simple
    type checking of the passed constants.

    """

    def __init__(self, func, backend='opencl'):
        backend = get_backend(backend)
        if backend == 'cython':
            raise NotImplementedError(
                'Kernels only work with opencl/cuda backends.'
            )
        elif backend == 'opencl':
            from .opencl import get_queue
            self.queue = get_queue()
        elif backend == 'cuda':
            from .cuda import set_context
            set_context()
        self.tp = Transpiler(backend=backend)
        self.backend = backend
        self.name = func.__name__
        self.func = func
        self.source = ''  # The generated source.
        self._config = get_config()
        self._use_double = self._config.use_double
        self._func_info = self._get_func_info()
        self._generate()

    def _to_float(self, s):
        return re.sub(r'\bdouble\b', 'float', s)

    def _get_func_info(self):
        try:
            getfullargspec = inspect.getfullargspec
        except AttributeError:
            # compatibility with Python 2.7
            getfullargspec = inspect.getargspec
        argspec = getfullargspec(self.func)
        annotations = getattr(
            argspec, 'annotations', self.func.__annotations__
        )

        arg_info = []
        local_info = {}
        for arg in argspec.args:
            kt = annotations[arg]
            if not self._use_double:
                kt = KnownType(
                    self._to_float(kt.type), self._to_float(kt.base_type)
                )
            if 'LOCAL_MEM' in kt.type:
                local_info[arg] = kt.base_type
            arg_info.append((arg, kt))
        func_info = {
            'args': arg_info,
            'local_info': local_info,
            'return': annotations.get('return', KnownType('void'))
        }
        return func_info

    def _get_local_size(self, args, workgroup_size):
        local_info = self._func_info['local_info']
        arg_info = self._func_info['args']
        total_size = 0
        for arg, a_info in zip(args, arg_info):
            if isinstance(arg, LocalMem):
                dtype = ctype_to_dtype(local_info[a_info[0]])
                total_size += dtype.itemsize
        return workgroup_size * total_size

    def _generate(self):
        self.tp.add(self.func)
        self._correct_opencl_address_space()

        self.tp.compile()
        self.source = self.tp.source

        if self.backend == 'opencl':
            self.knl = getattr(self.tp.mod, self.name)
            import pyopencl as cl
            self._max_work_group_size = self.knl.get_work_group_info(
                cl.kernel_work_group_info.WORK_GROUP_SIZE,
                self.queue.device
            )
        elif self.backend == 'cuda':
            self.knl = self.tp.mod.get_function(self.name)

    def _correct_opencl_address_space(self):
        code = self.tp.blocks[-1].code.splitlines()
        # To remove WITHIN_KERNEL
        code[0] = 'KERNEL ' + code[0][13:]
        self.tp.blocks[-1].code = '\n'.join(code)

    def _massage_arg(self, x, type_info, workgroup_size):
        if isinstance(x, Array):
            if self.backend == 'opencl':
                return x.dev.data
            elif self.backend == 'cuda':
                return x.dev
        elif isinstance(x, LocalMem):
            if self.backend == 'opencl':
                return x.get(type_info.base_type, workgroup_size)
            elif self.backend == 'cuda':
                return np.array(workgroup_size, dtype=np.int32)
        else:
            dtype = ctype_to_dtype(type_info.type)
            return np.array([x], dtype=dtype)

    def _get_args(self, args, workgroup_size):
        arg_info = self._func_info['args']
        c_args = []
        for arg, a_info in zip(args, arg_info):
            c_args.append(self._massage_arg(arg, a_info[1], workgroup_size))
        return c_args

    def _get_workgroup_size(self, global_size):
        if self.backend == 'opencl':
            gs, ls = splay_cl(self.queue, global_size,
                              self._max_work_group_size)
        elif self.backend == 'cuda':
            from pycuda.gpuarray import splay
            gs, ls = splay(global_size)
        return gs, ls

    @profile
    def __call__(self, *args, **kw):
        size = args[0].data.shape
        gs = kw.pop('global_size', size)
        n = np.prod(gs)
        ls = kw.pop('local_size', None)
        if ls is not None:
            local_size = np.prod(ls)
            global_size = ((n + local_size - 1) // local_size) * local_size
            gs = (global_size, )
        else:
            gs, ls = self._get_workgroup_size(n)
        c_args = self._get_args(args, ls[0])
        if self.backend == 'opencl':
            prepend = [self.queue, gs, ls]
            c_args = prepend + c_args
            self.knl(*c_args)
            self.queue.finish()
        elif self.backend == 'cuda':
            import pycuda.driver as drv
            shared_mem_size = int(self._get_local_size(args, ls[0]))
            num_blocks = int((n + ls[0] - 1) / ls[0])
            num_tpb = int(ls[0])
            event = drv.Event()
            self.knl(*c_args, block=(num_tpb, 1, 1), grid=(num_blocks, 1),
                     shared=shared_mem_size)
            event.record()
            event.synchronize()


class _prange(Extern):
    def code(self, backend):
        if backend != 'cython':
            raise NotImplementedError('prange only available with Cython')
        return 'from cython.parallel import prange'

    def __call__(self, *args, **kw):
        # Ignore the kwargs.
        return range(*args)


class _parallel(Extern):
    def code(self, backend):
        if backend != 'cython':
            raise NotImplementedError('prange only available with Cython')
        return 'from cython.parallel import parallel'

    def __call__(self, *args, **kw):
        pass


class _nogil(Extern):
    def code(self, backend):
        if backend != 'cython':
            raise NotImplementedError('prange only available with Cython')
        return ''

    def __call__(self, *args, **kw):
        pass


class _address(Extern):
    def code(self, backend):
        if backend == 'cython':
            return 'from cython import address'
        else:
            return ''

    def __call__(self, *args, **kw):
        pass


class _atomic_inc(Extern):
    def code(self, backend):
        return ''

    def __call__(self, *args, **kw):
        pass


class _atomic_dec(Extern):
    def code(self, backend):
        return ''

    def __call__(self, *args, **kw):
        pass


class _cast(Extern):
    def code(self, backend):
        return ''

    def __call__(self, x, type_str):
        return eval(type_str)(x)


prange = _prange()
parallel = _parallel()
nogil = _nogil()
address = _address()
atomic_inc = _atomic_inc()
atomic_dec = _atomic_dec()
cast = _cast()


class Cython(object):
    def __init__(self, func):
        self.tp = Transpiler(backend='cython')
        self.tp._cgen.set_make_python_methods(True)
        self.name = func.__name__
        self.func = func
        self.source = ''  # The generated source.
        self._generate()

    def _generate(self):
        self.tp.add(self.func)
        self.tp.compile()
        self.source = self.tp.source
        self.c_func = getattr(self.tp.mod, 'py_' + self.name)

    def _massage_arg(self, x):
        if isinstance(x, Array):
            return x.data
        else:
            return x

    def __call__(self, *args):
        args = [self._massage_arg(x) for x in args]
        return self.c_func(*args)