[Refactor][NFC] Vendor-in compiler_lock for future CUDA-specific changes

Author: atmnp

numba_cuda/numba/cuda/compiler.py

@@ -10,7 +10,7 @@
 from numba.cuda.core import bytecode
 from numba.cuda import types
 from numba.cuda.core.options import ParallelOptions
-from numba.core.compiler_lock import global_compiler_lock
+from numba.cuda.core.compiler_lock import global_compiler_lock
 from numba.cuda.core.errors import NumbaWarning, NumbaInvalidConfigWarning
 from numba.cuda.core.interpreter import Interpreter
 

numba_cuda/numba/cuda/core/base.py

@@ -15,7 +15,7 @@
 from numba.cuda.core import imputils, targetconfig, funcdesc
 from numba.cuda import cgutils, debuginfo, types, utils, datamodel, config
 from numba.core import errors
-from numba.core.compiler_lock import global_compiler_lock
+from numba.cuda.core.compiler_lock import global_compiler_lock
 from numba.cuda.core.pythonapi import PythonAPI
 from numba.cuda.core.imputils import (
     user_function,

numba_cuda/numba/cuda/core/compiler_lock.py

@@ -0,0 +1,59 @@
+# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: BSD-2-Clause
+
+import threading
+import functools
+import numba.cuda.core.event as ev
+
+
+# Lock for the preventing multiple compiler execution
+class _CompilerLock(object):
+    def __init__(self):
+        self._lock = threading.RLock()
+
+    def acquire(self):
+        ev.start_event("numba.cuda:compiler_lock")
+        self._lock.acquire()
+
+    def release(self):
+        self._lock.release()
+        ev.end_event("numba.cuda:compiler_lock")
+
+    def __enter__(self):
+        self.acquire()
+
+    def __exit__(self, exc_val, exc_type, traceback):
+        self.release()
+
+    def is_locked(self):
+        is_owned = getattr(self._lock, "_is_owned")
+        if not callable(is_owned):
+            is_owned = self._is_owned
+        return is_owned()
+
+    def __call__(self, func):
+        @functools.wraps(func)
+        def _acquire_compile_lock(*args, **kwargs):
+            with self:
+                return func(*args, **kwargs)
+
+        return _acquire_compile_lock
+
+    def _is_owned(self):
+        # This method is borrowed from threading.Condition.
+        # Return True if lock is owned by current_thread.
+        # This method is called only if _lock doesn't have _is_owned().
+        if self._lock.acquire(0):
+            self._lock.release()
+            return False
+        else:
+            return True
+
+
+global_compiler_lock = _CompilerLock()
+
+
+def require_global_compiler_lock():
+    """Sentry that checks the global_compiler_lock is acquired."""
+    # Use assert to allow turning off this check
+    assert global_compiler_lock.is_locked()

numba_cuda/numba/cuda/core/compiler_machinery.py

@@ -7,7 +7,7 @@
 import inspect
 
 
-from numba.core.compiler_lock import global_compiler_lock
+from numba.cuda.core.compiler_lock import global_compiler_lock
 from numba.core import errors
 from numba.cuda.core import config
 from numba.cuda import utils

numba_cuda/numba/cuda/dispatcher.py

@@ -19,7 +19,7 @@
 from numba.cuda import serialize, utils
 from numba import cuda
 
-from numba.core.compiler_lock import global_compiler_lock
+from numba.cuda.core.compiler_lock import global_compiler_lock
 from numba.cuda.typeconv.rules import default_type_manager
 from numba.cuda.typing.templates import fold_arguments
 from numba.cuda.typing.typeof import Purpose, typeof

numba_cuda/numba/cuda/target.py

@@ -10,7 +10,7 @@
 
 from numba.cuda import types
 
-from numba.core.compiler_lock import global_compiler_lock
+from numba.cuda.core.compiler_lock import global_compiler_lock
 from numba.cuda.core.errors import NumbaWarning
 from numba.cuda.core.base import BaseContext
 from numba.cuda.typing import cmathdecl