Merge pull request #7858 from isVoid/deprecate/cuda/ptx

CUDA: Deprecate `ptx` Attribute and Update Tests

numba/cuda/dispatcher.py

@@ -3,11 +3,12 @@
 import sys
 import ctypes
 import functools
+import warnings
 
 from numba.core import config, serialize, sigutils, types, typing, utils
 from numba.core.compiler_lock import global_compiler_lock
 from numba.core.dispatcher import Dispatcher
-from numba.core.errors import NumbaPerformanceWarning
+from numba.core.errors import NumbaPerformanceWarning, NumbaDeprecationWarning
 from numba.core.typing.typeof import Purpose, typeof
 
 from numba.cuda.api import get_current_device
@@ -154,6 +155,13 @@ def ptx(self):
         '''
         PTX code for this kernel.
         '''
+        warnings.warn(
+            "Attribute `ptx` is deprecated and will be removed in the future. "
+            "To retrieve the compiled machine code of the CUDA function for a "
+            "given CUDA compute compatibility `cc`, use the `inspect_asm(cc)` "
+            "method."
+            , NumbaDeprecationWarning
+        )
         return self._codelibrary.get_asm_str()
 
     @property
@@ -809,6 +817,12 @@ def inspect_types(self, file=None):
 
     @property
     def ptx(self):
+        warnings.warn(
+            "Attribute `ptx` is deprecated and will be removed in the future. "
+            "To retrieve the compiled machine code of the CUDA function for a "
+            "given signature `sig`, use the method `inspect_asm(sig)`."
+            , NumbaDeprecationWarning
+        )
         return {sig: overload.ptx for sig, overload in self.overloads.items()}
 
     def bind(self):

numba/cuda/tests/cudapy/test_constmem.py

@@ -100,7 +100,7 @@ def test_const_array(self):
         if not ENABLE_CUDASIM:
             self.assertIn(
                 'ld.const.f64',
-                jcuconst.ptx[sig],
+                jcuconst.inspect_asm(sig),
                 "as we're adding to it, load as a double")
 
     def test_const_empty(self):
@@ -125,7 +125,7 @@ def test_const_array_2d(self):
         if not ENABLE_CUDASIM:
             self.assertIn(
                 'ld.const.u32',
-                jcuconst2d.ptx[sig],
+                jcuconst2d.inspect_asm(sig),
                 "load the ints as ints")
 
     def test_const_array_3d(self):
@@ -146,7 +146,9 @@ def test_const_array_3d(self):
                 complex_load = 'ld.const.f32'
                 description = 'load each half of the complex as f32'
 
-            self.assertIn(complex_load, jcuconst3d.ptx[sig], description)
+            self.assertIn(
+                complex_load, jcuconst3d.inspect_asm(sig), description
+            )
 
     def test_const_record_empty(self):
         jcuconstRecEmpty = cuda.jit('void(int64[:])')(cuconstRecEmpty)

numba/cuda/tests/cudapy/test_dispatcher.py

@@ -1,8 +1,12 @@
 import numpy as np
 import threading
+import warnings
+from contextlib import contextmanager
 
 from numba import cuda, float32, float64, int32, int64, void
+from numba.core.errors import NumbaDeprecationWarning
 from numba.cuda.testing import skip_on_cudasim, unittest, CUDATestCase
+from numba.tests.support import ignore_internal_warnings
 import math
 
 
@@ -288,5 +292,50 @@ def add_device(a, b):
         self.assertEqual("Add two integers, device version", add_device.__doc__)
 
 
+@contextmanager
+def deprecation_warning_catcher():
+    with warnings.catch_warnings(record=True) as w:
+        ignore_internal_warnings()
+        warnings.simplefilter("always", category=NumbaDeprecationWarning)
+        yield w
+
+
+@skip_on_cudasim('Dispatcher objects not used in the simulator')
+class TestDispatcherDeprecation(CUDATestCase):
+    def check_warning(self, warnings, expected_str, category):
+        self.assertEqual(len(warnings), 1)
+        for w in warnings:
+            self.assertEqual(w.category, category)
+            self.assertIn(expected_str, str(w.message))
+
+    def test_ptx_deprecation(self):
+        def f(x):
+            x[0] = 42
+
+        msg = "Attribute `ptx` is deprecated and will be removed in the "
+        "future. "
+
+        # Kernel code path
+        with self.subTest(
+            name="DispatcherOnCudaKernel"
+        ), deprecation_warning_catcher() as w:
+            dispatcher = cuda.jit(f)
+            dispatcher.ptx
+            self.check_warning(w, msg, NumbaDeprecationWarning)
+        with self.subTest(
+            name="KernelObjectOnCudaKernel"
+        ), deprecation_warning_catcher() as w:
+            kernel = dispatcher.compile((float32[::1],))
+            kernel.ptx
+            self.check_warning(w, msg, NumbaDeprecationWarning)
+        # Device function code path
+        with self.subTest(
+            name="DispatcherOnDeviceFunction"
+        ), deprecation_warning_catcher() as w:
+            device_f_dispatcher = cuda.jit(f, device=True)
+            device_f_dispatcher.ptx
+            self.check_warning(w, msg, NumbaDeprecationWarning)
+
+
 if __name__ == '__main__':
     unittest.main()

numba/cuda/tests/cudapy/test_localmem.py

@@ -35,7 +35,7 @@ class TestCudaLocalMem(CUDATestCase):
     def test_local_array(self):
         sig = (int32[:], int32[:])
         jculocal = cuda.jit(sig)(culocal)
-        self.assertTrue('.local' in jculocal.ptx[sig])
+        self.assertTrue('.local' in jculocal.inspect_asm(sig))
         A = np.arange(1000, dtype='int32')
         B = np.zeros_like(A)
         jculocal[1, 1](A, B)

numba/cuda/tests/cudapy/test_sync.py

@@ -186,7 +186,7 @@ def test_threadfence_codegen(self):
         compiled[1, 1](ary)
         self.assertEqual(123 + 321, ary[0])
         if not ENABLE_CUDASIM:
-            self.assertIn("membar.gl;", compiled.ptx[sig])
+            self.assertIn("membar.gl;", compiled.inspect_asm(sig))
 
     def test_threadfence_block_codegen(self):
         # Does not test runtime behavior, just the code generation.
@@ -196,7 +196,7 @@ def test_threadfence_block_codegen(self):
         compiled[1, 1](ary)
         self.assertEqual(123 + 321, ary[0])
         if not ENABLE_CUDASIM:
-            self.assertIn("membar.cta;", compiled.ptx[sig])
+            self.assertIn("membar.cta;", compiled.inspect_asm(sig))
 
     def test_threadfence_system_codegen(self):
         # Does not test runtime behavior, just the code generation.
@@ -206,7 +206,7 @@ def test_threadfence_system_codegen(self):
         compiled[1, 1](ary)
         self.assertEqual(123 + 321, ary[0])
         if not ENABLE_CUDASIM:
-            self.assertIn("membar.sys;", compiled.ptx[sig])
+            self.assertIn("membar.sys;", compiled.inspect_asm(sig))
 
     def test_syncthreads_count(self):
         compiled = cuda.jit("void(int32[:], int32[:])")(use_syncthreads_count)