from numba.cuda.cudadrv import nvvm
from numba.cuda.testing import (
unittest,
skip_on_cudasim,
SerialMixin,
skip_unless_conda_cudatoolkit,
)
from numba.cuda.cuda_paths import (
_get_libdevice_path_decision,
_get_nvvm_path_decision,
_get_cudalib_dir_path_decision,
get_system_ctk,
)
has_cuda = nvvm.is_available()
has_mp_get_context = hasattr(mp, 'get_context')
class LibraryLookupBase(SerialMixin, unittest.TestCase):
def setUp(self):
ctx = mp.get_context('spawn')
qrecv = ctx.Queue()
qsend = ctx.Queue()
self.qsend = qsend
self.qrecv = qrecv
self.child_process = ctx.Process(
target=check_lib_lookup,
args=(qrecv, qsend),
daemon=True,
from numba.cuda.cudadrv import nvvm
from numba.cuda.testing import skip_on_cudasim, SerialMixin
from numba.core import types, utils
import unittest
original = "call void @llvm.memset.p0i8.i64(" \
"i8* align 4 %arg.x.41, i8 0, i64 %0, i1 false)"
missing_align = "call void @llvm.memset.p0i8.i64(" \
"i8* %arg.x.41, i8 0, i64 %0, i1 false)"
@skip_on_cudasim('libNVVM not supported in simulator')
@unittest.skipIf(utils.MACHINE_BITS == 32, "CUDA not support for 32-bit")
@unittest.skipIf(not nvvm.is_available(), "No libNVVM")
class TestNvvmWithoutCuda(SerialMixin, unittest.TestCase):
def test_nvvm_llvm_to_ptx(self):
"""
A simple test to exercise nvvm.llvm_to_ptx()
to trigger issues with mismatch NVVM API.
"""
def foo(x):
x[0] = 123
cukern = compile_kernel(foo, args=(types.int32[::1], ), link=())
llvmir = cukern._func.ptx.llvmir
ptx = nvvm.llvm_to_ptx(llvmir)
self.assertIn("foo", ptx.decode('ascii'))
def test_nvvm_memset_fixup(self):
from __future__ import absolute_import, print_function, division
from numba.cuda.compiler import compile_kernel
from numba.cuda.cudadrv import nvvm
from numba.cuda.testing import skip_on_cudasim
from numba import unittest_support as unittest
from numba import types, utils
@skip_on_cudasim('libNVVM not supported in simulator')
@unittest.skipIf(utils.MACHINE_BITS == 32, "CUDA not support for 32-bit")
@unittest.skipIf(not nvvm.is_available(), "No libNVVM")
class TestNvvmWithoutCuda(unittest.TestCase):
def test_nvvm_llvm_to_ptx(self):
"""
A simple test to exercise nvvm.llvm_to_ptx()
to trigger issues with mismatch NVVM API.
"""
def foo(x):
x[0] = 123
cukern = compile_kernel(foo, args=(types.int32[::1],), link=())
llvmir = cukern._func.ptx.llvmir
ptx = nvvm.llvm_to_ptx(llvmir)
self.assertIn("foo", ptx.decode('ascii'))
if __name__ == '__main__':
unittest.main()
from numba.config import IS_WIN32, IS_OSX
from numba.cuda.cudadrv import nvvm
from numba.cuda.testing import (
unittest,
skip_on_cudasim,
SerialMixin,
skip_unless_conda_cudatoolkit,
)
from numba.cuda.cuda_paths import (
_get_libdevice_path_decision,
_get_nvvm_path_decision,
_get_cudalib_dir_path_decision,
get_system_ctk,
)
has_cuda = nvvm.is_available()
has_mp_get_context = hasattr(mp, 'get_context')
class LibraryLookupBase(SerialMixin, unittest.TestCase):
def setUp(self):
ctx = mp.get_context('spawn')
qrecv = ctx.Queue()
qsend = ctx.Queue()
self.qsend = qsend
self.qrecv = qrecv
self.child_process = ctx.Process(
target=check_lib_lookup,
args=(qrecv, qsend),
daemon=True,
