We are currently solving the following issue within our repository. Here is the issue text:
--- BEGIN ISSUE ---
python primitive scalar float fails as jit.rawkernel argument
### Description

Passing python primitive float as a jit.rawkernel argument seems not to work.  Numpy.float32 works, python float does not.

### To Reproduce

```py
@jit.rawkernel()
def scalar_multiply(a, m, size):
    tid = jit.blockIdx.x * jit.blockDim.x + jit.threadIdx.x
    ntid = jit.gridDim.x * jit.blockDim.x
    for i in range(tid, size, ntid):
        a[i] = a[i] * m

size = 5
a = cp.random.random(size, dtype=np.float32)

print(a)
scalar_multiply((128,),(1024,),(a, np.float32(5.0), size))
print(a)
scalar_multiply((128,),(1024,),(a, 5.0, size))
print(a)
```
output:

[0.17959814 0.42873758 0.77541053 0.8213136  0.8284943 ]
[0.8979907 2.143688  3.8770528 4.1065683 4.1424713]
[0. 0. 0. 0. 0.]

### Installation

Wheel

### Environment

```
OS                           : Linux-5.15.0-41-generic-x86_64-with-glibc2.29
Python Version               : 3.8.10
CuPy Version                 : 11.0.0
CuPy Platform                : NVIDIA CUDA
NumPy Version                : 1.22.0
SciPy Version                : 1.4.1
Cython Build Version         : 0.29.24
Cython Runtime Version       : 0.29.28
CUDA Root                    : /usr/local/cuda
nvcc PATH                    : /usr/local/cuda/bin/nvcc
CUDA Build Version           : 11070
CUDA Driver Version          : 11070
CUDA Runtime Version         : 11070
cuBLAS Version               : (available)
cuFFT Version                : 10702
cuRAND Version               : 10210
cuSOLVER Version             : (11, 3, 5)
cuSPARSE Version             : (available)
NVRTC Version                : (11, 7)
Thrust Version               : 101500
CUB Build Version            : 101500
Jitify Build Version         : 4a37de0
cuDNN Build Version          : 8400
cuDNN Version                : 8401
NCCL Build Version           : None
NCCL Runtime Version         : None
cuTENSOR Version             : None
cuSPARSELt Build Version     : None
Device 0 Name                : NVIDIA GeForce GTX 950
Device 0 Compute Capability  : 52
Device 0 PCI Bus ID          : 0000:01:00.0
```


### Additional Information

Thanks!

--- END ISSUE ---

Below are some code segments, each from a relevant file. One or more of these files may contain bugs.
--- BEGIN FILES ---
Path: `cupyx/jit/_interface.py`
Content:
```
1 import functools
2 import warnings
3 
4 import numpy
5 
6 from cupy_backends.cuda.api import runtime
7 import cupy
8 from cupy._core import core
9 from cupyx.jit import _compile
10 from cupyx.jit import _cuda_typerules
11 from cupyx.jit import _cuda_types
12 from cupyx.jit import _internal_types
13 
14 
15 class _CudaFunction:
16     """JIT cupy function object
17     """
18 
19     def __init__(self, func, mode, device=False, inline=False):
20         self.attributes = []
21 
22         if device:
23             self.attributes.append('__device__')
24         else:
25             self.attributes.append('__global__')
26 
27         if inline:
28             self.attributes.append('inline')
29 
30         self.name = getattr(func, 'name', func.__name__)
31         self.func = func
32         self.mode = mode
33 
34     def __call__(self, *args, **kwargs):
35         raise NotImplementedError
36 
37     def _emit_code_from_types(self, in_types, ret_type=None):
38         return _compile.transpile(
39             self.func, self.attributes, self.mode, in_types, ret_type)
40 
41 
42 class _JitRawKernel:
43     """JIT CUDA kernel object.
44 
45     The decorator :func:``cupyx.jit.rawkernel`` converts the target function
46     to an object of this class. This class is not inteded to be instantiated
47     by users.
48     """
49 
50     def __init__(self, func, mode, device):
51         self._func = func
52         self._mode = mode
53         self._device = device
54         self._cache = {}
55         self._cached_codes = {}
56 
57     def __call__(
58             self, grid, block, args, shared_mem=0, stream=None):
59         """Calls the CUDA kernel.
60 
61         The compilation will be deferred until the first function call.
62         CuPy's JIT compiler infers the types of arguments at the call
63         time, and will cache the compiled kernels for speeding up any
64         subsequent calls.
65 
66         Args:
67             grid (tuple of int): Size of grid in blocks.
68             block (tuple of int): Dimensions of each thread block.
69             args (tuple):
70                 Arguments of the kernel. The type of all elements must be
71                 ``bool``, ``int``, ``float``, ``complex``, NumPy scalar or
72                 ``cupy.ndarray``.
73             shared_mem (int):
74                 Dynamic shared-memory size per thread block in bytes.
75             stream (cupy.cuda.Stream): CUDA stream.
76 
77         .. seealso:: :ref:`jit_kernel_definition`
78         """
79         in_types = []
80         for x in args:
81             if isinstance(x, cupy.ndarray):
82                 t = _cuda_types.CArray.from_ndarray(x)
83             elif numpy.isscalar(x):
84                 t = _cuda_typerules.get_ctype_from_scalar(self._mode, x)
85             else:
86                 raise TypeError(f'{type(x)} is not supported for RawKernel')
87             in_types.append(t)
88         in_types = tuple(in_types)
89         device_id = cupy.cuda.get_device_id()
90 
91         kern, enable_cg = self._cache.get((in_types, device_id), (None, None))
92         if kern is None:
93             result = self._cached_codes.get(in_types)
94             if result is None:
95                 result = _compile.transpile(
96                     self._func,
97                     ['extern "C"', '__global__'],
98                     self._mode,
99                     in_types,
100                     _cuda_types.void,
101                 )
102                 self._cached_codes[in_types] = result
103 
104             fname = result.func_name
105             enable_cg = result.enable_cooperative_groups
106             # workaround for hipRTC: as of ROCm 4.1.0 hipRTC still does not
107             # recognize "-D", so we have to compile using hipcc...
108             backend = 'nvcc' if runtime.is_hip else 'nvrtc'
109             module = core.compile_with_cache(
110                 source=result.code,
111                 options=('-DCUPY_JIT_MODE', '--std=c++14'),
112                 backend=backend)
113             kern = module.get_function(fname)
114             self._cache[(in_types, device_id)] = (kern, enable_cg)
115 
116         kern(grid, block, args, shared_mem, stream, enable_cg)
117 
118     def __getitem__(self, grid_and_block):
119         """Numba-style kernel call.
120 
121         .. seealso:: :ref:`jit_kernel_definition`
122         """
123         grid, block = grid_and_block
124         if not isinstance(grid, tuple):
125             grid = (grid, 1, 1)
126         if not isinstance(block, tuple):
127             block = (block, 1, 1)
128         return lambda *args, **kwargs: self(grid, block, args, **kwargs)
129 
130     @property
131     def cached_codes(self):
132         """Returns a dict that has input types as keys and codes values.
133 
134         This proprety method is for debugging purpose.
135         The return value is not guaranteed to keep backward compatibility.
136         """
137         if len(self._cached_codes) == 0:
138             warnings.warn(
139                 'No codes are cached because compilation is deferred until '
140                 'the first function call.')
141         return dict([(k, v.code) for k, v in self._cached_codes.items()])
142 
143     @property
144     def cached_code(self):
145         """Returns `next(iter(self.cached_codes.values()))`.
146 
147         This proprety method is for debugging purpose.
148         The return value is not guaranteed to keep backward compatibility.
149         """
150         codes = self.cached_codes
151         if len(codes) > 1:
152             warnings.warn(
153                 'The input types of the kernel could not be inferred. '
154                 'Please use `.cached_codes` instead.')
155         return next(iter(codes.values()))
156 
157 
158 def rawkernel(*, mode='cuda', device=False):
159     """A decorator compiles a Python function into CUDA kernel.
160     """
161     cupy._util.experimental('cupyx.jit.rawkernel')
162 
163     def wrapper(func):
164         return functools.update_wrapper(
165             _JitRawKernel(func, mode, device), func)
166     return wrapper
167 
168 
169 threadIdx = _internal_types.Data('threadIdx', _cuda_types.dim3)
170 blockDim = _internal_types.Data('blockDim', _cuda_types.dim3)
171 blockIdx = _internal_types.Data('blockIdx', _cuda_types.dim3)
172 gridDim = _internal_types.Data('gridDim', _cuda_types.dim3)
173 
174 warpsize = _internal_types.Data(
175     '64' if runtime.is_hip else '32', _cuda_types.uint32)
176 warpsize.__doc__ = r"""Returns the number of threads in a warp.
177 
178 In CUDA this is always 32, and in ROCm/HIP always 64.
179 
180 .. seealso:: :obj:`numba.cuda.warpsize`
181 """
182
```
--- END FILES ---

Please first localize the bug based on the issue statement, and then generate a patch according to the `git diff` format fenced by three backticks.

Here is an example:

```diff
diff --git a/examples/server_async.py b/examples/server_async.py
--- a/examples/server_async.py
+++ b/examples/server_async.py
@@ -313,4 +313,4 @@
 
 
 if __name__ == "__main__":
-    asyncio.run(run_async_server("."), debug=True)
+    asyncio.run(run_async_server(), debug=True)
diff --git a/examples/server_sync.py b/examples/server_sync.py
--- a/examples/server_sync.py
+++ b/examples/server_sync.py
@@ -313,5 +313,5 @@
 
 
 if __name__ == "__main__":
-    server = run_sync_server(".")
+    server = run_sync_server()
     server.shutdown()

```

diff --git a/cupyx/jit/_interface.py b/cupyx/jit/_interface.py
--- a/cupyx/jit/_interface.py
+++ b/cupyx/jit/_interface.py
@@ -10,6 +10,7 @@
 from cupyx.jit import _cuda_typerules
 from cupyx.jit import _cuda_types
 from cupyx.jit import _internal_types
+from cupyx.jit._cuda_types import Scalar
 
 
 class _CudaFunction:
@@ -113,7 +114,16 @@
             kern = module.get_function(fname)
             self._cache[(in_types, device_id)] = (kern, enable_cg)
 
-        kern(grid, block, args, shared_mem, stream, enable_cg)
+        new_args = []
+        for a, t in zip(args, in_types):
+            if isinstance(t, Scalar):
+                if t.dtype.char == 'e':
+                    a = numpy.float32(a)
+                else:
+                    a = t.dtype.type(a)
+            new_args.append(a)
+
+        kern(grid, block, tuple(new_args), shared_mem, stream, enable_cg)
 
     def __getitem__(self, grid_and_block):
         """Numba-style kernel call.

{"golden_diff": "diff --git a/cupyx/jit/_interface.py b/cupyx/jit/_interface.py\n--- a/cupyx/jit/_interface.py\n+++ b/cupyx/jit/_interface.py\n@@ -10,6 +10,7 @@\n from cupyx.jit import _cuda_typerules\n from cupyx.jit import _cuda_types\n from cupyx.jit import _internal_types\n+from cupyx.jit._cuda_types import Scalar\n \n \n class _CudaFunction:\n@@ -113,7 +114,16 @@\n             kern = module.get_function(fname)\n             self._cache[(in_types, device_id)] = (kern, enable_cg)\n \n-        kern(grid, block, args, shared_mem, stream, enable_cg)\n+        new_args = []\n+        for a, t in zip(args, in_types):\n+            if isinstance(t, Scalar):\n+                if t.dtype.char == 'e':\n+                    a = numpy.float32(a)\n+                else:\n+                    a = t.dtype.type(a)\n+            new_args.append(a)\n+\n+        kern(grid, block, tuple(new_args), shared_mem, stream, enable_cg)\n \n     def __getitem__(self, grid_and_block):\n         \"\"\"Numba-style kernel call.\n", "issue": "python primitive scalar float fails as jit.rawkernel argument\n### Description\r\n\r\nPassing python primitive float as a jit.rawkernel argument seems not to work.  Numpy.float32 works, python float does not.\r\n\r\n### To Reproduce\r\n\r\n```py\r\n@jit.rawkernel()\r\ndef scalar_multiply(a, m, size):\r\n    tid = jit.blockIdx.x * jit.blockDim.x + jit.threadIdx.x\r\n    ntid = jit.gridDim.x * jit.blockDim.x\r\n    for i in range(tid, size, ntid):\r\n        a[i] = a[i] * m\r\n\r\nsize = 5\r\na = cp.random.random(size, dtype=np.float32)\r\n\r\nprint(a)\r\nscalar_multiply((128,),(1024,),(a, np.float32(5.0), size))\r\nprint(a)\r\nscalar_multiply((128,),(1024,),(a, 5.0, size))\r\nprint(a)\r\n```\r\noutput:\r\n\r\n[0.17959814 0.42873758 0.77541053 0.8213136  0.8284943 ]\r\n[0.8979907 2.143688  3.8770528 4.1065683 4.1424713]\r\n[0. 0. 0. 0. 0.]\r\n\r\n### Installation\r\n\r\nWheel\r\n\r\n### Environment\r\n\r\n```\r\nOS                           : Linux-5.15.0-41-generic-x86_64-with-glibc2.29\r\nPython Version               : 3.8.10\r\nCuPy Version                 : 11.0.0\r\nCuPy Platform                : NVIDIA CUDA\r\nNumPy Version                : 1.22.0\r\nSciPy Version                : 1.4.1\r\nCython Build Version         : 0.29.24\r\nCython Runtime Version       : 0.29.28\r\nCUDA Root                    : /usr/local/cuda\r\nnvcc PATH                    : /usr/local/cuda/bin/nvcc\r\nCUDA Build Version           : 11070\r\nCUDA Driver Version          : 11070\r\nCUDA Runtime Version         : 11070\r\ncuBLAS Version               : (available)\r\ncuFFT Version                : 10702\r\ncuRAND Version               : 10210\r\ncuSOLVER Version             : (11, 3, 5)\r\ncuSPARSE Version             : (available)\r\nNVRTC Version                : (11, 7)\r\nThrust Version               : 101500\r\nCUB Build Version            : 101500\r\nJitify Build Version         : 4a37de0\r\ncuDNN Build Version          : 8400\r\ncuDNN Version                : 8401\r\nNCCL Build Version           : None\r\nNCCL Runtime Version         : None\r\ncuTENSOR Version             : None\r\ncuSPARSELt Build Version     : None\r\nDevice 0 Name                : NVIDIA GeForce GTX 950\r\nDevice 0 Compute Capability  : 52\r\nDevice 0 PCI Bus ID          : 0000:01:00.0\r\n```\r\n\r\n\r\n### Additional Information\r\n\r\nThanks!\n", "before_files": [{"content": "import functools\nimport warnings\n\nimport numpy\n\nfrom cupy_backends.cuda.api import runtime\nimport cupy\nfrom cupy._core import core\nfrom cupyx.jit import _compile\nfrom cupyx.jit import _cuda_typerules\nfrom cupyx.jit import _cuda_types\nfrom cupyx.jit import _internal_types\n\n\nclass _CudaFunction:\n    \"\"\"JIT cupy function object\n    \"\"\"\n\n    def __init__(self, func, mode, device=False, inline=False):\n        self.attributes = []\n\n        if device:\n            self.attributes.append('__device__')\n        else:\n            self.attributes.append('__global__')\n\n        if inline:\n            self.attributes.append('inline')\n\n        self.name = getattr(func, 'name', func.__name__)\n        self.func = func\n        self.mode = mode\n\n    def __call__(self, *args, **kwargs):\n        raise NotImplementedError\n\n    def _emit_code_from_types(self, in_types, ret_type=None):\n        return _compile.transpile(\n            self.func, self.attributes, self.mode, in_types, ret_type)\n\n\nclass _JitRawKernel:\n    \"\"\"JIT CUDA kernel object.\n\n    The decorator :func:``cupyx.jit.rawkernel`` converts the target function\n    to an object of this class. This class is not inteded to be instantiated\n    by users.\n    \"\"\"\n\n    def __init__(self, func, mode, device):\n        self._func = func\n        self._mode = mode\n        self._device = device\n        self._cache = {}\n        self._cached_codes = {}\n\n    def __call__(\n            self, grid, block, args, shared_mem=0, stream=None):\n        \"\"\"Calls the CUDA kernel.\n\n        The compilation will be deferred until the first function call.\n        CuPy's JIT compiler infers the types of arguments at the call\n        time, and will cache the compiled kernels for speeding up any\n        subsequent calls.\n\n        Args:\n            grid (tuple of int): Size of grid in blocks.\n            block (tuple of int): Dimensions of each thread block.\n            args (tuple):\n                Arguments of the kernel. The type of all elements must be\n                ``bool``, ``int``, ``float``, ``complex``, NumPy scalar or\n                ``cupy.ndarray``.\n            shared_mem (int):\n                Dynamic shared-memory size per thread block in bytes.\n            stream (cupy.cuda.Stream): CUDA stream.\n\n        .. seealso:: :ref:`jit_kernel_definition`\n        \"\"\"\n        in_types = []\n        for x in args:\n            if isinstance(x, cupy.ndarray):\n                t = _cuda_types.CArray.from_ndarray(x)\n            elif numpy.isscalar(x):\n                t = _cuda_typerules.get_ctype_from_scalar(self._mode, x)\n            else:\n                raise TypeError(f'{type(x)} is not supported for RawKernel')\n            in_types.append(t)\n        in_types = tuple(in_types)\n        device_id = cupy.cuda.get_device_id()\n\n        kern, enable_cg = self._cache.get((in_types, device_id), (None, None))\n        if kern is None:\n            result = self._cached_codes.get(in_types)\n            if result is None:\n                result = _compile.transpile(\n                    self._func,\n                    ['extern \"C\"', '__global__'],\n                    self._mode,\n                    in_types,\n                    _cuda_types.void,\n                )\n                self._cached_codes[in_types] = result\n\n            fname = result.func_name\n            enable_cg = result.enable_cooperative_groups\n            # workaround for hipRTC: as of ROCm 4.1.0 hipRTC still does not\n            # recognize \"-D\", so we have to compile using hipcc...\n            backend = 'nvcc' if runtime.is_hip else 'nvrtc'\n            module = core.compile_with_cache(\n                source=result.code,\n                options=('-DCUPY_JIT_MODE', '--std=c++14'),\n                backend=backend)\n            kern = module.get_function(fname)\n            self._cache[(in_types, device_id)] = (kern, enable_cg)\n\n        kern(grid, block, args, shared_mem, stream, enable_cg)\n\n    def __getitem__(self, grid_and_block):\n        \"\"\"Numba-style kernel call.\n\n        .. seealso:: :ref:`jit_kernel_definition`\n        \"\"\"\n        grid, block = grid_and_block\n        if not isinstance(grid, tuple):\n            grid = (grid, 1, 1)\n        if not isinstance(block, tuple):\n            block = (block, 1, 1)\n        return lambda *args, **kwargs: self(grid, block, args, **kwargs)\n\n    @property\n    def cached_codes(self):\n        \"\"\"Returns a dict that has input types as keys and codes values.\n\n        This proprety method is for debugging purpose.\n        The return value is not guaranteed to keep backward compatibility.\n        \"\"\"\n        if len(self._cached_codes) == 0:\n            warnings.warn(\n                'No codes are cached because compilation is deferred until '\n                'the first function call.')\n        return dict([(k, v.code) for k, v in self._cached_codes.items()])\n\n    @property\n    def cached_code(self):\n        \"\"\"Returns `next(iter(self.cached_codes.values()))`.\n\n        This proprety method is for debugging purpose.\n        The return value is not guaranteed to keep backward compatibility.\n        \"\"\"\n        codes = self.cached_codes\n        if len(codes) > 1:\n            warnings.warn(\n                'The input types of the kernel could not be inferred. '\n                'Please use `.cached_codes` instead.')\n        return next(iter(codes.values()))\n\n\ndef rawkernel(*, mode='cuda', device=False):\n    \"\"\"A decorator compiles a Python function into CUDA kernel.\n    \"\"\"\n    cupy._util.experimental('cupyx.jit.rawkernel')\n\n    def wrapper(func):\n        return functools.update_wrapper(\n            _JitRawKernel(func, mode, device), func)\n    return wrapper\n\n\nthreadIdx = _internal_types.Data('threadIdx', _cuda_types.dim3)\nblockDim = _internal_types.Data('blockDim', _cuda_types.dim3)\nblockIdx = _internal_types.Data('blockIdx', _cuda_types.dim3)\ngridDim = _internal_types.Data('gridDim', _cuda_types.dim3)\n\nwarpsize = _internal_types.Data(\n    '64' if runtime.is_hip else '32', _cuda_types.uint32)\nwarpsize.__doc__ = r\"\"\"Returns the number of threads in a warp.\n\nIn CUDA this is always 32, and in ROCm/HIP always 64.\n\n.. seealso:: :obj:`numba.cuda.warpsize`\n\"\"\"\n", "path": "cupyx/jit/_interface.py"}], "after_files": [{"content": "import functools\nimport warnings\n\nimport numpy\n\nfrom cupy_backends.cuda.api import runtime\nimport cupy\nfrom cupy._core import core\nfrom cupyx.jit import _compile\nfrom cupyx.jit import _cuda_typerules\nfrom cupyx.jit import _cuda_types\nfrom cupyx.jit import _internal_types\nfrom cupyx.jit._cuda_types import Scalar\n\n\nclass _CudaFunction:\n    \"\"\"JIT cupy function object\n    \"\"\"\n\n    def __init__(self, func, mode, device=False, inline=False):\n        self.attributes = []\n\n        if device:\n            self.attributes.append('__device__')\n        else:\n            self.attributes.append('__global__')\n\n        if inline:\n            self.attributes.append('inline')\n\n        self.name = getattr(func, 'name', func.__name__)\n        self.func = func\n        self.mode = mode\n\n    def __call__(self, *args, **kwargs):\n        raise NotImplementedError\n\n    def _emit_code_from_types(self, in_types, ret_type=None):\n        return _compile.transpile(\n            self.func, self.attributes, self.mode, in_types, ret_type)\n\n\nclass _JitRawKernel:\n    \"\"\"JIT CUDA kernel object.\n\n    The decorator :func:``cupyx.jit.rawkernel`` converts the target function\n    to an object of this class. This class is not inteded to be instantiated\n    by users.\n    \"\"\"\n\n    def __init__(self, func, mode, device):\n        self._func = func\n        self._mode = mode\n        self._device = device\n        self._cache = {}\n        self._cached_codes = {}\n\n    def __call__(\n            self, grid, block, args, shared_mem=0, stream=None):\n        \"\"\"Calls the CUDA kernel.\n\n        The compilation will be deferred until the first function call.\n        CuPy's JIT compiler infers the types of arguments at the call\n        time, and will cache the compiled kernels for speeding up any\n        subsequent calls.\n\n        Args:\n            grid (tuple of int): Size of grid in blocks.\n            block (tuple of int): Dimensions of each thread block.\n            args (tuple):\n                Arguments of the kernel. The type of all elements must be\n                ``bool``, ``int``, ``float``, ``complex``, NumPy scalar or\n                ``cupy.ndarray``.\n            shared_mem (int):\n                Dynamic shared-memory size per thread block in bytes.\n            stream (cupy.cuda.Stream): CUDA stream.\n\n        .. seealso:: :ref:`jit_kernel_definition`\n        \"\"\"\n        in_types = []\n        for x in args:\n            if isinstance(x, cupy.ndarray):\n                t = _cuda_types.CArray.from_ndarray(x)\n            elif numpy.isscalar(x):\n                t = _cuda_typerules.get_ctype_from_scalar(self._mode, x)\n            else:\n                raise TypeError(f'{type(x)} is not supported for RawKernel')\n            in_types.append(t)\n        in_types = tuple(in_types)\n        device_id = cupy.cuda.get_device_id()\n\n        kern, enable_cg = self._cache.get((in_types, device_id), (None, None))\n        if kern is None:\n            result = self._cached_codes.get(in_types)\n            if result is None:\n                result = _compile.transpile(\n                    self._func,\n                    ['extern \"C\"', '__global__'],\n                    self._mode,\n                    in_types,\n                    _cuda_types.void,\n                )\n                self._cached_codes[in_types] = result\n\n            fname = result.func_name\n            enable_cg = result.enable_cooperative_groups\n            # workaround for hipRTC: as of ROCm 4.1.0 hipRTC still does not\n            # recognize \"-D\", so we have to compile using hipcc...\n            backend = 'nvcc' if runtime.is_hip else 'nvrtc'\n            module = core.compile_with_cache(\n                source=result.code,\n                options=('-DCUPY_JIT_MODE', '--std=c++14'),\n                backend=backend)\n            kern = module.get_function(fname)\n            self._cache[(in_types, device_id)] = (kern, enable_cg)\n\n        new_args = []\n        for a, t in zip(args, in_types):\n            if isinstance(t, Scalar):\n                if t.dtype.char == 'e':\n                    a = numpy.float32(a)\n                else:\n                    a = t.dtype.type(a)\n            new_args.append(a)\n\n        kern(grid, block, tuple(new_args), shared_mem, stream, enable_cg)\n\n    def __getitem__(self, grid_and_block):\n        \"\"\"Numba-style kernel call.\n\n        .. seealso:: :ref:`jit_kernel_definition`\n        \"\"\"\n        grid, block = grid_and_block\n        if not isinstance(grid, tuple):\n            grid = (grid, 1, 1)\n        if not isinstance(block, tuple):\n            block = (block, 1, 1)\n        return lambda *args, **kwargs: self(grid, block, args, **kwargs)\n\n    @property\n    def cached_codes(self):\n        \"\"\"Returns a dict that has input types as keys and codes values.\n\n        This proprety method is for debugging purpose.\n        The return value is not guaranteed to keep backward compatibility.\n        \"\"\"\n        if len(self._cached_codes) == 0:\n            warnings.warn(\n                'No codes are cached because compilation is deferred until '\n                'the first function call.')\n        return dict([(k, v.code) for k, v in self._cached_codes.items()])\n\n    @property\n    def cached_code(self):\n        \"\"\"Returns `next(iter(self.cached_codes.values()))`.\n\n        This proprety method is for debugging purpose.\n        The return value is not guaranteed to keep backward compatibility.\n        \"\"\"\n        codes = self.cached_codes\n        if len(codes) > 1:\n            warnings.warn(\n                'The input types of the kernel could not be inferred. '\n                'Please use `.cached_codes` instead.')\n        return next(iter(codes.values()))\n\n\ndef rawkernel(*, mode='cuda', device=False):\n    \"\"\"A decorator compiles a Python function into CUDA kernel.\n    \"\"\"\n    cupy._util.experimental('cupyx.jit.rawkernel')\n\n    def wrapper(func):\n        return functools.update_wrapper(\n            _JitRawKernel(func, mode, device), func)\n    return wrapper\n\n\nthreadIdx = _internal_types.Data('threadIdx', _cuda_types.dim3)\nblockDim = _internal_types.Data('blockDim', _cuda_types.dim3)\nblockIdx = _internal_types.Data('blockIdx', _cuda_types.dim3)\ngridDim = _internal_types.Data('gridDim', _cuda_types.dim3)\n\nwarpsize = _internal_types.Data(\n    '64' if runtime.is_hip else '32', _cuda_types.uint32)\nwarpsize.__doc__ = r\"\"\"Returns the number of threads in a warp.\n\nIn CUDA this is always 32, and in ROCm/HIP always 64.\n\n.. seealso:: :obj:`numba.cuda.warpsize`\n\"\"\"\n", "path": "cupyx/jit/_interface.py"}]}