[Refactor][NFC] Vendor-in compiler_lock for future CUDA-specific changes (#565)

This change vendors in the global_compiler_lock that is used to protect
the core compiler machinery 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.cuda.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,85 @@
+# 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
+from numba.cuda import HAS_NUMBA
+
+if HAS_NUMBA:
+    from numba.core.compiler_lock import (
+        global_compiler_lock as _numba_compiler_lock,
+    )
+else:
+    _numba_compiler_lock = None
+
+
+# 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 __call__(self, func):
+        @functools.wraps(func)
+        def _acquire_compile_lock(*args, **kwargs):
+            with self:
+                return func(*args, **kwargs)
+
+        return _acquire_compile_lock
+
+
+_numba_cuda_compiler_lock = _CompilerLock()
+
+
+# Wrapper that coordinates both numba and numba-cuda compiler locks
+class _DualCompilerLock(object):
+    """Wrapper that coordinates both the numba-cuda and upstream numba compiler locks."""
+
+    def __init__(self, cuda_lock, numba_lock):
+        self._cuda_lock = cuda_lock
+        self._numba_lock = numba_lock
+
+    def acquire(self):
+        self._numba_lock.acquire()
+        self._cuda_lock.acquire()
+
+    def release(self):
+        self._cuda_lock.release()
+        self._numba_lock.release()
+
+    def __enter__(self):
+        self.acquire()
+
+    def __exit__(self, exc_val, exc_type, traceback):
+        self.release()
+
+    def __call__(self, func):
+        @functools.wraps(func)
+        def _acquire_compile_lock(*args, **kwargs):
+            with self:
+                return func(*args, **kwargs)
+
+        return _acquire_compile_lock
+
+
+# Create the global compiler lock, wrapping both locks if numba is available
+if HAS_NUMBA:
+    global_compiler_lock = _DualCompilerLock(
+        _numba_cuda_compiler_lock, _numba_compiler_lock
+    )
+else:
+    global_compiler_lock = _numba_cuda_compiler_lock

numba_cuda/numba/cuda/core/compiler_machinery.py

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

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.cuda import HAS_NUMBA
-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