4
4
5
5
6
6
Extensions to Numba for Intel GPUs introduce two new features into Numba:
7
- a. A new backend that has a new decorator called @dppl .kernel that
7
+ a. A new backend that has a new decorator called @dppy .kernel that
8
8
exposes an explicit kernel programming interface similar to the
9
- existing Numba GPU code-generation backends. The @dppl .kernel
9
+ existing Numba GPU code-generation backends. The @dppy .kernel
10
10
decorator currently implements a subset of OpenCL’s API through
11
11
Numba’s intrinsic functions.
12
12
20
20
Explicit Kernel Prgoramming with new Docorators:
21
21
22
22
23
- @dppl .kernel
23
+ @dppy .kernel
24
24
25
- The @dppl .kernel decorator can be used with or without extra arguments.
25
+ The @dppy .kernel decorator can be used with or without extra arguments.
26
26
Optionally, users can pass the signature of the arguments to the
27
27
decorator. When a signature is provided to the DK decorator the version
28
28
of the OpenCL kernel generated gets specialized for that type signature.
29
29
30
30
---------------------------------------------------------------------------
31
- @dppl .kernel
31
+ @dppy .kernel
32
32
def data_parallel_sum(a, b, c):
33
- i = dppl .get_global_id(0)
33
+ i = dppy .get_global_id(0)
34
34
c[i] = a[i] + b[i]
35
35
---------------------------------------------------------------------------
36
36
37
37
To invoke the above function users will need to provide a
38
38
global size (OpenCL) which is the size of a (same as b and c) and a
39
- local size (dppl .DEFAULT_LOCAL_SIZE if user don't want to specify).
39
+ local size (dppy .DEFAULT_LOCAL_SIZE if user don't want to specify).
40
40
Example shown below:
41
41
42
42
---------------------------------------------------------------------------
43
- data_parallel_sum[len(a), dppl .DEFAULT_LOCAL_SIZE](dA, dB, dC)
43
+ data_parallel_sum[len(a), dppy .DEFAULT_LOCAL_SIZE](dA, dB, dC)
44
44
---------------------------------------------------------------------------
45
45
46
46
47
- @dppl .func
47
+ @dppy .func
48
48
49
- The @dppl .func decorator is the other decorator provided in the explicit
49
+ The @dppy .func decorator is the other decorator provided in the explicit
50
50
kernel programming model. This decorator allows users to write “device”
51
51
functions that can be invoked from inside DK functions but cannot be invoked
52
52
from the host. The decorator also supports type specialization as with the
53
- DK decorator. Functions decorated with @dppl .func will also be JIT compiled
54
- and inlined into the OpenCL Program containing the @dppl .kernel function
55
- calling it. A @dppl .func will not be launched as an OpenCL kernel.
53
+ DK decorator. Functions decorated with @dppy .func will also be JIT compiled
54
+ and inlined into the OpenCL Program containing the @dppy .kernel function
55
+ calling it. A @dppy .func will not be launched as an OpenCL kernel.
56
56
57
57
---------------------------------------------------------------------------
58
- @dppl .func
58
+ @dppy .func
59
59
def bar(a):
60
60
return a*a
61
61
62
- @dppl .kernel
62
+ @dppy .kernel
63
63
def foo(in, out):
64
- i = dppl .get_global_id(0)
64
+ i = dppy .get_global_id(0)
65
65
out[i] = bar(in[i])
66
66
---------------------------------------------------------------------------
67
67
@@ -71,13 +71,13 @@ def foo(in, out):
71
71
The following table has the list of intrinsic functions that can be directly
72
72
used inside a DK function. All the functions are equivalent to the similarly
73
73
named OpenCL function. Wherever there is an implementation difference
74
- between the Numba-PyDPPL version and the OpenCL version, the difference is
74
+ between the Numba-DPPY version and the OpenCL version, the difference is
75
75
explained in table. Note that these functions cannot be used anywhere else
76
76
outside of a DK function in a Numba application. Readers are referred to the
77
77
OpenCL API specs to review the functionality of each function.
78
78
79
79
+----------------------+----------------------------+----------------------+
80
- | Numba-DPPL intrinsic | Equivalent OpenCL function | Notes |
80
+ | Numba-DPPY intrinsic | Equivalent OpenCL function | Notes |
81
81
+----------------------+----------------------------+----------------------+
82
82
| get_global_id | get_global_id | |
83
83
+----------------------+----------------------------+----------------------+
@@ -121,7 +121,7 @@ def foo(in, out):
121
121
|print |print(varargs) |The print function is a |
122
122
| | |subset of the OpenCL |
123
123
| | |printf function. The |
124
- | | |Numba-DPPL version of |
124
+ | | |Numba-DPPY version of |
125
125
| | |print supports only int, |
126
126
| | |string, and float |
127
127
| | |arguments. |
@@ -160,16 +160,16 @@ def foo(in, out):
160
160
161
161
162
162
163
- Complete Example using @dppl .kernel:
163
+ Complete Example using @dppy .kernel:
164
164
165
165
---------------------------------------------------------------------------
166
166
import numpy as np
167
- import numba_dppy, numba_dppy as dppl
167
+ import numba_dppy, numba_dppy as dppy
168
168
import dpctl
169
169
170
- @dppl .kernel
170
+ @dppy .kernel
171
171
def data_parallel_sum(a, b, c):
172
- i = dppl .get_global_id(0)
172
+ i = dppy .get_global_id(0)
173
173
c[i] = a[i] + b[i]
174
174
175
175
def driver(device_env, a, b, c, global_size):
@@ -181,7 +181,7 @@ def driver(device_env, a, b, c, global_size):
181
181
print("before : ", dA._ndarray)
182
182
print("before : ", dB._ndarray)
183
183
print("before : ", dC._ndarray)
184
- data_parallel_sum[global_size, dppl .DEFAULT_LOCAL_SIZE](dA, dB, dC)
184
+ data_parallel_sum[global_size, dppy .DEFAULT_LOCAL_SIZE](dA, dB, dC)
185
185
device_env.copy_array_from_device(dC)
186
186
print("after : ", dC._ndarray)
187
187
@@ -509,11 +509,11 @@ def main():
509
509
if dppy_present :
510
510
from .device_init import *
511
511
else :
512
- raise ImportError ("Importing dppl failed" )
512
+ raise ImportError ("Importing numba-dppy failed" )
513
513
514
514
def test (* args , ** kwargs ):
515
515
if not dppy_present and not is_available ():
516
- dppl_error ()
516
+ dppy_error ()
517
517
518
518
return numba .testing .test ("numba_dppy.tests" , * args , ** kwargs )
519
519
0 commit comments