stubs.py 2.76 KB
Newer Older
dugupeiwen's avatar
dugupeiwen committed
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
from numba.core import types, typing, ir

_stub_error = NotImplementedError("This is a stub.")


def get_global_id(*args, **kargs):
    """
    OpenCL get_global_id()
    """
    raise _stub_error


def get_local_id(*args, **kargs):
    """
    OpenCL get_local_id()
    """
    raise _stub_error


def get_global_size(*args, **kargs):
    """
    OpenCL get_global_size()
    """
    raise _stub_error


def get_local_size(*args, **kargs):
    """
    OpenCL get_local_size()
    """
    raise _stub_error


def get_group_id(*args, **kargs):
    """
    OpenCL get_group_id()
    """
    raise _stub_error


def get_num_groups(*args, **kargs):
    """
    OpenCL get_num_groups()
    """
    raise _stub_error


def get_work_dim(*args, **kargs):
    """
    OpenCL get_work_dim()
    """
    raise _stub_error


def barrier(*args, **kargs):
    """
    OpenCL barrier()

    Example:

        # workgroup barrier + local memory fence
        hsa.barrier(hsa.CLK_LOCAL_MEM_FENCE)
        # workgroup barrier + global memory fence
        hsa.barrier(hsa.CLK_GLOBAL_MEM_FENCE)
        # workgroup barrier + global memory fence
        hsa.barrier()

    """
    raise _stub_error


def mem_fence(*args, **kargs):
    """
    OpenCL mem_fence()

    Example:

        # local memory fence
        hsa.mem_fence(hsa.CLK_LOCAL_MEM_FENCE)
        # global memory fence
        hsa.mem_fence(hsa.CLK_GLOBAL_MEM_FENCE)
    """
    raise _stub_error


def wavebarrier():
    """
    HSAIL wavebarrier
    """
    raise _stub_error


def activelanepermute_wavewidth(src, laneid, identity, useidentity):
    """
    HSAIL activelanepermute_wavewidth_*
    """
    raise _stub_error


def ds_permute(src_lane, dest_lane):
    """
    AMDGCN Data Share intrinsic forwards permute (push semantics)
    """
    raise _stub_error


def ds_bpermute(src_lane, dest_lane):
    """
    AMDGCN Data Share intrinsic backwards permute (pull semantics)
    """
    raise _stub_error


class Stub(object):
    """A stub object to represent special objects which is meaningless
    outside the context of HSA-python.
    """
    _description_ = '<ptx special value>'
    __slots__ = ()  # don't allocate __dict__

    def __new__(cls):
        raise NotImplementedError("%s is not instantiable" % cls)

    def __repr__(self):
        return self._description_


class shared(Stub):
    """shared namespace
    """
    _description_ = '<shared>'

    def array(shape, dtype):
        """shared.array(shape, dtype)

        Allocate a shared memory array.
        """


#-------------------------------------------------------------------------------
# atomic

class atomic(Stub):
    """atomic namespace
    """
    _description_ = '<atomic>'

    class add(Stub):
        """add(ary, idx, val)

        Perform atomic ary[idx] += val
        """