File: C:/Users/fred/anaconda3/Lib/site-packages/numba/cuda/tests/cudapy/test_sm_creation.py
import numpy as np
from numba import cuda, float32, int32, void
from numba.core.errors import TypingError
from numba.cuda.testing import unittest, CUDATestCase
from numba.cuda.testing import skip_on_cudasim
from .extensions_usecases import test_struct_model_type
GLOBAL_CONSTANT = 5
GLOBAL_CONSTANT_2 = 6
GLOBAL_CONSTANT_TUPLE = 5, 6
def udt_global_constants(A):
sa = cuda.shared.array(shape=GLOBAL_CONSTANT, dtype=float32)
i = cuda.grid(1)
A[i] = sa[i]
def udt_global_build_tuple(A):
sa = cuda.shared.array(shape=(GLOBAL_CONSTANT, GLOBAL_CONSTANT_2),
dtype=float32)
i, j = cuda.grid(2)
A[i, j] = sa[i, j]
def udt_global_build_list(A):
sa = cuda.shared.array(shape=[GLOBAL_CONSTANT, GLOBAL_CONSTANT_2],
dtype=float32)
i, j = cuda.grid(2)
A[i, j] = sa[i, j]
def udt_global_constant_tuple(A):
sa = cuda.shared.array(shape=GLOBAL_CONSTANT_TUPLE, dtype=float32)
i, j = cuda.grid(2)
A[i, j] = sa[i, j]
def udt_invalid_1(A):
sa = cuda.shared.array(shape=A[0], dtype=float32)
i = cuda.grid(1)
A[i] = sa[i]
def udt_invalid_2(A):
sa = cuda.shared.array(shape=(1, A[0]), dtype=float32)
i, j = cuda.grid(2)
A[i, j] = sa[i, j]
def udt_invalid_3(A):
sa = cuda.shared.array(shape=(1, A[0]), dtype=float32)
i = cuda.grid(1)
A[i] = sa[i, 0]
class TestSharedMemoryCreation(CUDATestCase):
def getarg(self):
return np.array(100, dtype=np.float32, ndmin=1)
def getarg2(self):
return self.getarg().reshape(1,1)
def test_global_constants(self):
udt = cuda.jit((float32[:],))(udt_global_constants)
udt[1, 1](self.getarg())
def test_global_build_tuple(self):
udt = cuda.jit((float32[:, :],))(udt_global_build_tuple)
udt[1, 1](self.getarg2())
@skip_on_cudasim('Simulator does not prohibit lists for shared array shape')
def test_global_build_list(self):
with self.assertRaises(TypingError) as raises:
cuda.jit((float32[:, :],))(udt_global_build_list)
self.assertIn("No implementation of function "
"Function(<function shared.array",
str(raises.exception))
self.assertIn("found for signature:\n \n "
">>> array(shape=list(int64)<iv=[5, 6]>, "
"dtype=class(float32)",
str(raises.exception))
def test_global_constant_tuple(self):
udt = cuda.jit((float32[:, :],))(udt_global_constant_tuple)
udt[1, 1](self.getarg2())
@skip_on_cudasim("Can't check for constants in simulator")
def test_invalid_1(self):
# Scalar shape cannot be a floating point value
with self.assertRaises(TypingError) as raises:
cuda.jit((float32[:],))(udt_invalid_1)
self.assertIn("No implementation of function "
"Function(<function shared.array",
str(raises.exception))
self.assertIn("found for signature:\n \n "
">>> array(shape=float32, dtype=class(float32))",
str(raises.exception))
@skip_on_cudasim("Can't check for constants in simulator")
def test_invalid_2(self):
# Tuple shape cannot contain a floating point value
with self.assertRaises(TypingError) as raises:
cuda.jit((float32[:, :],))(udt_invalid_2)
self.assertIn("No implementation of function "
"Function(<function shared.array",
str(raises.exception))
self.assertIn("found for signature:\n \n "
">>> array(shape=Tuple(Literal[int](1), "
"array(float32, 1d, A)), dtype=class(float32))",
str(raises.exception))
@skip_on_cudasim("Can't check for constants in simulator")
def test_invalid_3(self):
# Scalar shape must be literal
with self.assertRaises(TypingError) as raises:
cuda.jit((int32[:],))(udt_invalid_1)
self.assertIn("No implementation of function "
"Function(<function shared.array",
str(raises.exception))
self.assertIn("found for signature:\n \n "
">>> array(shape=int32, dtype=class(float32))",
str(raises.exception))
@skip_on_cudasim("Can't check for constants in simulator")
def test_invalid_4(self):
# Tuple shape must contain only literals
with self.assertRaises(TypingError) as raises:
cuda.jit((int32[:],))(udt_invalid_3)
self.assertIn("No implementation of function "
"Function(<function shared.array",
str(raises.exception))
self.assertIn("found for signature:\n \n "
">>> array(shape=Tuple(Literal[int](1), int32), "
"dtype=class(float32))",
str(raises.exception))
def check_dtype(self, f, dtype):
# Find the typing of the dtype argument to cuda.shared.array
annotation = next(iter(f.overloads.values()))._type_annotation
l_dtype = annotation.typemap['s'].dtype
# Ensure that the typing is correct
self.assertEqual(l_dtype, dtype)
@skip_on_cudasim("Can't check typing in simulator")
def test_numba_dtype(self):
# Check that Numba types can be used as the dtype of a shared array
@cuda.jit(void(int32[::1]))
def f(x):
s = cuda.shared.array(10, dtype=int32)
s[0] = x[0]
x[0] = s[0]
self.check_dtype(f, int32)
@skip_on_cudasim("Can't check typing in simulator")
def test_numpy_dtype(self):
# Check that NumPy types can be used as the dtype of a shared array
@cuda.jit(void(int32[::1]))
def f(x):
s = cuda.shared.array(10, dtype=np.int32)
s[0] = x[0]
x[0] = s[0]
self.check_dtype(f, int32)
@skip_on_cudasim("Can't check typing in simulator")
def test_string_dtype(self):
# Check that strings can be used to specify the dtype of a shared array
@cuda.jit(void(int32[::1]))
def f(x):
s = cuda.shared.array(10, dtype='int32')
s[0] = x[0]
x[0] = s[0]
self.check_dtype(f, int32)
@skip_on_cudasim("Can't check typing in simulator")
def test_invalid_string_dtype(self):
# Check that strings of invalid dtypes cause a typing error
re = ".*Invalid NumPy dtype specified: 'int33'.*"
with self.assertRaisesRegex(TypingError, re):
@cuda.jit(void(int32[::1]))
def f(x):
s = cuda.shared.array(10, dtype='int33')
s[0] = x[0]
x[0] = s[0]
@skip_on_cudasim("Can't check typing in simulator")
def test_type_with_struct_data_model(self):
@cuda.jit(void(test_struct_model_type[::1]))
def f(x):
s = cuda.shared.array(10, dtype=test_struct_model_type)
s[0] = x[0]
x[0] = s[0]
self.check_dtype(f, test_struct_model_type)
if __name__ == '__main__':
unittest.main()