Skip to content

Commit 037d494

Browse files
gmarkallgmarkall
committed
CUDASIM: Update to include recently-added APIs
Support for recently-added APIs in the CUDA simulator has fallen behind the hardware target. This commit adds in minimal support for the new APIs such that the test suite can complete successfully, and skips inappropriate tests, such as those checking NRT refcounting, or linking behavior.
1 parent f17c0b1 commit 037d494

19 files changed

Lines changed: 252 additions & 93 deletions

‎numba_cuda/numba/cuda/simulator/__init__.py‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,7 @@
3838
sys.modules["numba.cuda.cudadrv.devicearray"] = cudadrv.devicearray
3939
sys.modules["numba.cuda.cudadrv.devices"] = cudadrv.devices
4040
sys.modules["numba.cuda.cudadrv.driver"] = cudadrv.driver
41+
sys.modules["numba.cuda.cudadrv.linkable_code"] = cudadrv.linkable_code
4142
sys.modules["numba.cuda.cudadrv.runtime"] = cudadrv.runtime
4243
sys.modules["numba.cuda.cudadrv.drvapi"] = cudadrv.drvapi
4344
sys.modules["numba.cuda.cudadrv.error"] = cudadrv.error
@@ -46,3 +47,8 @@
4647
from . import compiler
4748

4849
sys.modules["numba.cuda.compiler"] = compiler
50+
51+
from numba.cuda.simulator import memory_management
52+
53+
sys.modules["numba.cuda.memory_management"] = memory_management
54+
sys.modules["numba.cuda.memory_management.nrt"] = memory_management.nrt

‎numba_cuda/numba/cuda/simulator/api.py‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,15 @@
77
from contextlib import contextmanager
88

99
from .cudadrv.devices import require_context, reset, gpus # noqa: F401
10+
from .cudadrv.linkable_code import (
11+
PTXSource, # noqa: F401
12+
CUSource, # noqa: F401
13+
Cubin, # noqa: F401
14+
Fatbin, # noqa: F401
15+
Archive, # noqa: F401
16+
Object, # noqa: F401
17+
LTOIR, # noqa: F401
18+
) # noqa: F401
1019
from .kernel import FakeCUDAKernel
1120
from numba.core.sigutils import is_signature
1221
from numba.core import config
@@ -72,6 +81,10 @@ def list_devices():
7281
return gpus
7382

7483

84+
def get_current_device():
85+
return gpus[0].device
86+
87+
7588
# Events
7689

7790

‎numba_cuda/numba/cuda/simulator/compiler.py‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,3 +7,4 @@
77
compile_for_current_device = None
88
compile_ptx = None
99
compile_ptx_for_current_device = None
10+
declare_device_function = None

‎numba_cuda/numba/cuda/simulator/cudadrv/driver.py‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,8 @@
33
provided to allow tests to import correctly.
44
"""
55

6+
from numba import config
7+
68

79
def device_memset(dst, val, size, stream=0):
810
dst.view("u1")[:size].fill(bytes([val])[0])
@@ -60,3 +62,8 @@ def launch_kernel(*args, **kwargs):
6062

6163

6264
USE_NV_BINDING = False
65+
66+
PyNvJitLinker = None
67+
68+
if config.ENABLE_CUDASIM:
69+
config.CUDA_ENABLE_PYNVJITLINK = False

‎numba_cuda/numba/cuda/simulator/cudadrv/libs.py‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,2 +1,6 @@
11
def check_static_lib(lib):
22
raise FileNotFoundError("Linking libraries not supported by cudasim")
3+
4+
5+
def get_cuda_include_dir():
6+
raise FileNotFoundError("CUDA includes not supported by cudasim")

‎numba_cuda/numba/cuda/simulator/cudadrv/linkable_code.py‎

Lines changed: 57 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,57 @@
1+
class LinkableCode:
2+
"""An object that holds code to be linked from memory.
3+
4+
:param data: A buffer containing the data to link.
5+
:param name: The name of the file to be referenced in any compilation or
6+
linking errors that may be produced.
7+
"""
8+
9+
def __init__(self, data, name=None):
10+
self.data = data
11+
self._name = name
12+
13+
@property
14+
def name(self):
15+
return self._name or self.default_name
16+
17+
18+
class PTXSource(LinkableCode):
19+
"""PTX source code in memory."""
20+
21+
default_name = "<unnamed-ptx>"
22+
23+
24+
class CUSource(LinkableCode):
25+
"""CUDA C/C++ source code in memory."""
26+
27+
default_name = "<unnamed-cu>"
28+
29+
30+
class Fatbin(LinkableCode):
31+
"""An ELF Fatbin in memory."""
32+
33+
default_name = "<unnamed-fatbin>"
34+
35+
36+
class Cubin(LinkableCode):
37+
"""An ELF Cubin in memory."""
38+
39+
default_name = "<unnamed-cubin>"
40+
41+
42+
class Archive(LinkableCode):
43+
"""An archive of objects in memory."""
44+
45+
default_name = "<unnamed-archive>"
46+
47+
48+
class Object(LinkableCode):
49+
"""An object file in memory."""
50+
51+
default_name = "<unnamed-object>"
52+
53+
54+
class LTOIR(LinkableCode):
55+
"""An LTOIR file in memory."""
56+
57+
default_name = "<unnamed-ltoir>"

‎numba_cuda/numba/cuda/simulator/kernelapi.py‎

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -63,7 +63,10 @@ class FakeCUDALocal(object):
6363
CUDA Local arrays
6464
"""
6565

66-
def array(self, shape, dtype):
66+
def array(self, shape, dtype, alignment=None):
67+
if alignment is not None:
68+
raise RuntimeError("Array alignment is not supported in cudasim")
69+
6770
if isinstance(dtype, types.Type):
6871
dtype = numpy_support.as_dtype(dtype)
6972
return np.empty(shape, dtype)
@@ -102,7 +105,10 @@ def __init__(self, dynshared_size):
102105
self._dynshared_size = dynshared_size
103106
self._dynshared = np.zeros(dynshared_size, dtype=np.byte)
104107

105-
def array(self, shape, dtype):
108+
def array(self, shape, dtype, alignment=None):
109+
if alignment is not None:
110+
raise RuntimeError("Array alignment is not supported in cudasim")
111+
106112
if isinstance(dtype, types.Type):
107113
dtype = numpy_support.as_dtype(dtype)
108114
# Dynamic shared memory is requested with size 0 - this all shares the

‎numba_cuda/numba/cuda/simulator/memory_management/__init__.py‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
from .nrt import rtsys # noqa: F401

‎numba_cuda/numba/cuda/simulator/memory_management/nrt.py‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,6 @@
1+
from numba import config
2+
3+
rtsys = None
4+
5+
config.CUDA_NRT_STATS = False
6+
config.CUDA_ENABLE_NRT = False

‎numba_cuda/numba/cuda/testing.py‎

Lines changed: 10 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -116,20 +116,26 @@ def skip_on_arm(reason):
116116
def skip_if_cuda_includes_missing(fn):
117117
# Skip when cuda.h is not available - generally this should indicate
118118
# whether the CUDA includes are available or not
119-
cuda_include_path = libs.get_cuda_include_dir()
119+
reason = "CUDA include dir not available on this system"
120+
try:
121+
cuda_include_path = libs.get_cuda_include_dir()
122+
except FileNotFoundError:
123+
return unittest.skip(reason)(fn)
120124
cuda_h = os.path.join(cuda_include_path, "cuda.h")
121125
cuda_h_file = os.path.exists(cuda_h) and os.path.isfile(cuda_h)
122-
reason = "CUDA include dir not available on this system"
123126
return unittest.skipUnless(cuda_h_file, reason)(fn)
124127

125128

126129
def skip_if_curand_kernel_missing(fn):
127-
cuda_include_path = libs.get_cuda_include_dir()
130+
reason = "curand_kernel.h not available on this system"
131+
try:
132+
cuda_include_path = libs.get_cuda_include_dir()
133+
except FileNotFoundError:
134+
return unittest.skip(reason)(fn)
128135
curand_kernel_h = os.path.join(cuda_include_path, "curand_kernel.h")
129136
curand_kernel_h_file = os.path.exists(curand_kernel_h) and os.path.isfile(
130137
curand_kernel_h
131138
)
132-
reason = "curand_kernel.h not available on this system"
133139
return unittest.skipUnless(curand_kernel_h_file, reason)(fn)
134140

135141

0 commit comments

Comments
 (0)