(fix:cuda) CUDA portability improved\nbuild system now checks for valid compiler compatible with nvcc and fails if that is not found and cuda is selected. Further, compiler choices are propegated throughout build along with CFLAGS and CXXFLAGS

build-config/mfem/meson.build

@@ -273,6 +273,57 @@ endif
 mfem_consumer_cpp_std = (
   mfem_features.get('raja') or mfem_features.get('umpire')
 ) ? 'c++20' : 'c++17'
+cc_command = cc.cmd_array()
+cxx_command = cpp.cmd_array()
+known_compiler_launchers = ['ccache', 'sccache', 'distcc']
+cc_executable = ''
+cc_launchers = []
+cc_fixed_args = []
+foreach command_part : cc_command
+  if cc_executable == '' and fs.name(command_part) in known_compiler_launchers
+    cc_launchers += [command_part]
+  elif cc_executable == ''
+    cc_executable = command_part
+  else
+    cc_fixed_args += [command_part]
+  endif
+endforeach
+cxx_executable = ''
+cxx_launchers = []
+cxx_fixed_args = []
+foreach command_part : cxx_command
+  if cxx_executable == '' and fs.name(command_part) in known_compiler_launchers
+    cxx_launchers += [command_part]
+  elif cxx_executable == ''
+    cxx_executable = command_part
+  else
+    cxx_fixed_args += [command_part]
+  endif
+endforeach
+if cc_executable == '' or cxx_executable == ''
+  error('Could not identify the C/C++ compiler executables in Meson compiler commands.')
+endif
+cuda_compiler_path = ''
+if mfem_has_cuda
+  cuda_compiler_path = dependency_has_nvcc ? (dependency_prefix / 'bin' / 'nvcc') : nvcc_program.full_path()
+  cuda_check = files('../../tools/check_cuda_toolchain.py')
+  cuda_check_args = []
+  foreach arg : cc_fixed_args + get_option('c_args') + platform_c_args
+    cuda_check_args += ['--c-arg=' + arg]
+  endforeach
+  foreach arg : cxx_fixed_args + get_option('cpp_args') + platform_cpp_args
+    cuda_check_args += ['--cxx-arg=' + arg]
+  endforeach
+  cuda_check_result = run_command(
+    python_build, cuda_check, '--nvcc', cuda_compiler_path,
+    '--cc', cc_executable, '--cxx', cxx_executable,
+    '--standard', mfem_consumer_cpp_std.substring(3), cuda_check_args, check: false,
+  )
+  if cuda_check_result.returncode() != 0
+    error(cuda_check_result.stdout() + cuda_check_result.stderr())
+  endif
+  message(cuda_check_result.stdout().strip())
+endif
 system_mfem = disabler()
 mfem_provider = 'source bundle'
 if effective_allow_preinstalled
@@ -428,37 +479,7 @@ else
   mfem_runtime_prefix = meson.current_build_dir() / bundle_output_name
   bundle_work_dir = meson.current_build_dir() / '_bundle' / 'work'
   bundle_builder = files('../../tools/build_mfem_bundle.py')
-  bundle_auxiliary_files = files('../../tools/ceed_jit_source_root_relocatable.c')
+  bundle_auxiliary_files = files('../../tools/ceed_jit_source_root_relocatable.c', '../../tools/check_cuda_toolchain.py')
-  cc_command = cc.cmd_array()
-  cxx_command = cpp.cmd_array()
-  known_compiler_launchers = ['ccache', 'sccache', 'distcc']
-  cc_executable = ''
-  cc_launchers = []
-  cc_fixed_args = []
-  foreach command_part : cc_command
-    if cc_executable == '' and fs.name(command_part) in known_compiler_launchers
-      cc_launchers += [command_part]
-    elif cc_executable == ''
-      cc_executable = command_part
-    else
-      cc_fixed_args += [command_part]
-    endif
-  endforeach
-  cxx_executable = ''
-  cxx_launchers = []
-  cxx_fixed_args = []
-  foreach command_part : cxx_command
-    if cxx_executable == '' and fs.name(command_part) in known_compiler_launchers
-      cxx_launchers += [command_part]
-    elif cxx_executable == ''
-      cxx_executable = command_part
-    else
-      cxx_fixed_args += [command_part]
-    endif
-  endforeach
-  if cc_executable == '' or cxx_executable == ''
-    error('Could not identify the C/C++ compiler executables in Meson compiler commands.')
-  endif
   bundle_command = [
     python_build,
     bundle_builder,
@@ -473,6 +494,7 @@ else
     '--host-system', host_machine.system(),
     '--cross-build', meson.is_cross_build() ? 'true' : 'false',
     '--library-kind', is_wasm ? 'static' : 'shared',
+    '--cuda-compiler', cuda_compiler_path,
     '--cuda-arch', get_option('mfem_cuda_arch'),
     '--hip-arch', get_option('mfem_hip_arch'),
     '--precision', get_option('mfem_precision'),

justfile

@@ -36,6 +36,9 @@ benchmark_max_applications := env_var_or_default("MFEM_BENCH_MAX_APPLICATIONS",
 benchmark_omp_threads := env_var_or_default("MFEM_BENCH_OMP_THREADS", "8")
 benchmark_mpi_ranks := env_var_or_default("MFEM_BENCH_MPI_RANKS", "4")
 
+cuda_cc := env_var_or_default("MFEM_CUDA_CC", env_var_or_default("CC", "cc"))
+cuda_cxx := env_var_or_default("MFEM_CUDA_CXX", env_var_or_default("CXX", "c++"))
+
 default:
     @just --list
 
@@ -124,7 +127,7 @@ cuda:
     if [[ -f "{{ cuda_dir }}/meson-private/coredata.dat" ]]; then
       setup_args=(--reconfigure "${setup_args[@]}")
     fi
-    "{{ meson_bin }}" setup "${setup_args[@]}"
+    CC="{{ cuda_cc }}" CXX="{{ cuda_cxx }}" "{{ meson_bin }}" setup "${setup_args[@]}"
     "{{ meson_bin }}" compile -C "{{ cuda_dir }}"
     "{{ meson_bin }}" test -C "{{ cuda_dir }}" --print-errorlogs
 

tests/meson.build

@@ -161,3 +161,9 @@ if get_option('build_tests') and get_option('build_python')
     timeout: 120,
   )
 endif
+
+if get_option('build_tests')
+  test('cuda-toolchain-diagnostics', python_build,
+    args: files('test_cuda_toolchain.py'),
+  )
+endif

tests/test_cuda_toolchain.py

@@ -0,0 +1,88 @@
+"""CUDA diagnostics and compiler propagation, without requiring a toolkit/GPU."""
+import os
+from pathlib import Path
+import subprocess
+import sys
+import tempfile
+from types import SimpleNamespace
+import unittest
+from unittest.mock import patch
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / 'tools'))
+from check_cuda_toolchain import check_toolchain
+from build_mfem_bundle import mpich_cuda_architectures, prepare_hypre_source
+
+
+class CudaToolchainTest(unittest.TestCase):
+    def test_probe_checks_both_compilers_without_running_gpu_code(self):
+        calls = []
+
+        def run(command, **kwargs):
+            calls.append(command)
+            if '--version' not in command:
+                self.assertNotIn('NVCC_APPEND_FLAGS', kwargs['env'])
+                self.assertNotIn('NVCC_CCBIN', kwargs['env'])
+            return subprocess.CompletedProcess(command, 0, 'CUDA test toolkit', '')
+
+        with patch.dict(os.environ, {'NVCC_APPEND_FLAGS': '-allow-unsupported-compiler',
+                                     'NVCC_CCBIN': 'wrong-host'}), patch('subprocess.run', side_effect=run):
+            check_toolchain('nvcc', 'gcc-compatible', 'g++-compatible')
+        self.assertEqual(len(calls), 3)
+        self.assertEqual(calls[1][2], 'gcc-compatible')
+        self.assertIn('-c', calls[1])  # A C driver need not link the C++ runtime.
+        self.assertEqual(calls[2][2], 'g++-compatible')
+        self.assertNotIn('-c', calls[2])
+
+    def test_mpich_architectures_follow_toolkit_and_user_selection(self):
+        self.assertEqual(mpich_cuda_architectures('86-real;90-virtual;86', 'nvcc'), '86,90')
+        self.assertEqual(mpich_cuda_architectures('native', 'nvcc'), 'auto')
+        with patch('build_mfem_bundle.output', return_value='sm_75\nsm_80\nsm_86\nsm_90'):
+            self.assertEqual(mpich_cuda_architectures('all-major', 'nvcc'), '75,80,90')
+            self.assertEqual(mpich_cuda_architectures('all', 'nvcc'), '75,80,86,90')
+        with self.assertRaisesRegex(RuntimeError, 'Invalid CUDA architecture'):
+            mpich_cuda_architectures('garbage', 'nvcc')
+
+    def test_hypre_patch_preserves_archive_and_older_cuda_branch(self):
+        with tempfile.TemporaryDirectory() as directory:
+            root = Path(directory)
+            original = root / 'original'
+            utilities = original / 'src' / 'utilities'
+            utilities.mkdir(parents=True)
+            header = ('#include <thrust/execution_policy.h>\n'
+                      '#define HYPRE_THRUST_NOT(pred) thrust::not1(pred)\n'
+                      'using namespace thrust::placeholders;\n')
+            for name in ('device_utils.h', '_hypre_utilities.hpp'):
+                (utilities / name).write_text(header)
+            memory = ''.join(
+                f'HYPRE_CUDA_CALL( cudaMemPrefetchAsync(ptr, size, {device},\n'
+                '                                            hypre_HandleComputeStream(hypre_handle())) );\n'
+                for device in ('hypre_HandleDevice(hypre_handle())', 'cudaCpuDeviceId'))
+            (utilities / 'memory.c').write_text(memory)
+            args = SimpleNamespace(hypre_source=original, work_dir=root / 'work')
+            self.assertEqual(prepare_hypre_source(args, {'cuda': False}), original)
+            patched = prepare_hypre_source(args, {'cuda': True})
+            self.assertEqual((utilities / 'memory.c').read_text(), memory)
+            updated = (patched / 'src/utilities/memory.c').read_text()
+            self.assertIn('#if CUDART_VERSION >= 13000', updated)
+            self.assertIn('cudaMemLocationTypeHost, 0', updated)
+            self.assertIn('cudaMemPrefetchAsync(ptr, size, cudaCpuDeviceId,', updated)
+            self.assertIn('#define HYPRE_THRUST_NOT(pred) thrust::not1(pred)',
+                          (patched / 'src/utilities/device_utils.h').read_text())
+            prepare_hypre_source(args, {'cuda': True})
+            self.assertEqual((patched / 'src/utilities/memory.c').read_text(), updated)
+
+    def test_cxx_failure_reports_toolkit_host_and_recovery(self):
+        results = [subprocess.CompletedProcess([], 0, 'CUDA test toolkit', ''),
+                   subprocess.CompletedProcess([], 0, '', ''),
+                   subprocess.CompletedProcess([], 1, '', 'unsupported Microsoft Visual Studio version')]
+        with patch('subprocess.run', side_effect=results):
+            with self.assertRaises(RuntimeError) as caught:
+                check_toolchain('nvcc.exe', 'cl.exe', 'cl.exe')
+        message = str(caught.exception)
+        for expected in ['C++ compiler: cl.exe', 'CUDA test toolkit',
+                         'unsupported Microsoft Visual Studio version', 'fresh Meson build directory']:
+            self.assertIn(expected, message)
+
+
+if __name__ == '__main__':
+    unittest.main()

tools/build_mfem_bundle.py

@@ -16,6 +16,8 @@ import sys
 import tarfile
 from typing import Iterable, Mapping, Sequence
 
+from check_cuda_toolchain import check_toolchain
+
 
 CMAKE_FEATURES = {
     "mpi": "MFEM_USE_MPI",
@@ -89,6 +91,7 @@ def parse_args() -> argparse.Namespace:
     parser.add_argument("--host-system", required=True)
     parser.add_argument("--cross-build", choices=("true", "false"), required=True)
     parser.add_argument("--library-kind", choices=("static", "shared"), required=True)
+    parser.add_argument("--cuda-compiler", default="")
     parser.add_argument("--cuda-arch", default="native")
     parser.add_argument("--hip-arch", default="")
     parser.add_argument("--precision", choices=("single", "double"), default="double")
@@ -356,6 +359,18 @@ def common_cmake_args(
         "-DCMAKE_FIND_USE_PACKAGE_REGISTRY=OFF",
         f"-DBUILD_SHARED_LIBS={'ON' if args.library_kind == 'shared' else 'OFF'}",
     ]
+    if args.cuda_compiler:
+        cuda_standard = 20 if any(feature in args.feature for feature in ("raja=ON", "umpire=ON")) else 17
+        result.extend([
+            f"-DCUDAToolkit_ROOT={Path(args.cuda_compiler).parent.parent}",
+            f"-DCMAKE_CUDA_COMPILER={args.cuda_compiler}",
+            f"-DCMAKE_CUDA_HOST_COMPILER={args.cxx}",
+            f"-DCUDA_HOST_COMPILER={args.cxx}",
+            f"-DCUDA_NVCC_EXECUTABLE={args.cuda_compiler}",
+            f"-DCMAKE_CUDA_STANDARD={cuda_standard}",
+            "-DCMAKE_CUDA_STANDARD_REQUIRED=ON",
+            "-DCMAKE_CXX_STANDARD=17",
+        ])
     if c_compiler:
         result.append(f"-DCMAKE_C_COMPILER={c_compiler}")
         if c_compiler == args.cc and args.cc_launcher:
@@ -502,6 +517,25 @@ def prepare_mpich_source(args: argparse.Namespace) -> Path:
     return source
 
 
+def mpich_cuda_architectures(architecture: str, nvcc: str) -> str:
+    """Translate CMake architecture syntax without Yaksa's obsolete default list."""
+    if architecture == "native":
+        return "auto"
+    if architecture in {"all", "all-major"}:
+        supported = re.findall(r"sm_([0-9]+[a-z]?)", output([nvcc, "--list-gpu-code"]))
+        if architecture == "all-major":
+            supported = [sm for i, sm in enumerate(supported)
+                         if i == 0 or sm.endswith("0")]
+        if not supported:
+            raise RuntimeError("nvcc did not report any supported GPU architectures")
+        return ",".join(dict.fromkeys(supported))
+    architectures = [re.sub(r"-(real|virtual)$", "", sm)
+                     for sm in architecture.split(";")]
+    if not all(re.fullmatch(r"[0-9]+[a-z]?", sm) for sm in architectures):
+        raise RuntimeError(f"Invalid CUDA architecture list: {architecture}")
+    return ",".join(dict.fromkeys(architectures))
+
+
 def build_mpich(
     args: argparse.Namespace,
     features: Mapping[str, bool],
@@ -560,7 +594,9 @@ def build_mpich(
         nvcc = find_program_in_environment("nvcc", env)
         if not nvcc:
             raise RuntimeError("CUDA-enabled MPI requested but nvcc is unavailable")
-        configure.append(f"--with-cuda={Path(nvcc).resolve().parent.parent}")
+        configure.append("--with-cuda-sm=" + mpich_cuda_architectures(args.cuda_arch, args.cuda_compiler))
+        local_env["NVCC"] = shlex.join([args.cuda_compiler, "-ccbin", args.cxx, "-std=c++17"])
+        configure.append(f"--with-cuda={Path(args.cuda_compiler).resolve().parent.parent}")
     elif features["hip"]:
         hipcc = find_program_in_environment("hipcc", env)
         if not hipcc:
@@ -638,6 +674,79 @@ def build_metis(
     shutil.copy2(source / "LICENSE.txt", legal_dir / "metis-LICENSE.txt")
 
 
+def prepare_hypre_source(args: argparse.Namespace, features: Mapping[str, bool]) -> Path:
+    if not features["cuda"]:
+        return args.hypre_source
+    source = args.work_dir / "hypre-source"
+    shutil.copytree(args.hypre_source, source, dirs_exist_ok=True, symlinks=True)
+    # CCCL 3 no longer supplies these public aliases through transitive includes.
+    for name in ("_hypre_utilities.hpp", "device_utils.h"):
+        header = source / "src" / "utilities" / name
+        text = header.read_text()
+        anchor = "#include <thrust/execution_policy.h>"
+        if anchor not in text:
+            raise RuntimeError(f"Hypre Thrust include anchor missing in {header}")
+        header.write_text(text.replace(anchor, anchor + "\n#include <thrust/tuple.h>\n#include <thrust/pair.h>\n#include <thrust/iterator/reverse_iterator.h>"))
+    # Recent CCCL releases removed these deprecated Thrust helper types.
+    # Keep the older typed predicate contract,
+    # including the typedefs required by older Thrust's not1 adaptor.
+    identity = """
+#ifndef HYPRE_MESON_THRUST_IDENTITY
+#define HYPRE_MESON_THRUST_IDENTITY
+template <typename T, typename R> struct hypre_meson_unary_function
+{
+   typedef T argument_type;
+   typedef R result_type;
+};
+template <typename T, typename U, typename R> struct hypre_meson_binary_function
+{
+   typedef T first_argument_type;
+   typedef U second_argument_type;
+   typedef R result_type;
+};
+template <typename T> struct hypre_meson_identity
+{
+   typedef T argument_type;
+   typedef T result_type;
+   __host__ __device__ const T& operator()(const T& value) const { return value; }
+};
+#endif
+"""
+    for name in ("_hypre_utilities.hpp", "device_utils.h"):
+        header = source / "src" / "utilities" / name
+        text = header.read_text()
+        header.write_text(text.replace("using namespace thrust::placeholders;",
+                                       identity + "using namespace thrust::placeholders;"))
+    for file in (source / "src").rglob("*"):
+        if file.suffix not in {".c", ".h", ".hpp"}:
+            continue
+        text = file.read_text()
+        patched = text.replace("thrust::identity<", "hypre_meson_identity<")
+        patched = patched.replace("thrust::unary_function<", "hypre_meson_unary_function<")
+        patched = patched.replace("thrust::binary_function<", "hypre_meson_binary_function<")
+        if file.suffix == ".c":
+            patched = patched.replace("thrust::not1(", "HYPRE_THRUST_NOT(")
+        if patched != text:
+            file.write_text(patched)
+    memory = source / "src" / "utilities" / "memory.c"
+    text = memory.read_text()
+    for device, location_type in (("hypre_HandleDevice(hypre_handle())", "cudaMemLocationTypeDevice"),
+                                  ("cudaCpuDeviceId", "cudaMemLocationTypeHost")):
+        call = (f"HYPRE_CUDA_CALL( cudaMemPrefetchAsync(ptr, size, {device},\n"
+                "                                            hypre_HandleComputeStream(hypre_handle())) );")
+        if call not in text:
+            raise RuntimeError("Hypre CUDA prefetch patch anchor missing")
+        device_id = device if location_type == "cudaMemLocationTypeDevice" else "0"
+        replacement = (f"#if CUDART_VERSION >= 13000\n"
+                       f"      cudaMemLocation destination = {{{location_type}, {device_id}}};\n"
+                       "      HYPRE_CUDA_CALL( cudaMemPrefetchAsync(ptr, size, destination, 0,\n"
+                       "                                            hypre_HandleComputeStream(hypre_handle())) );\n"
+                       "#else\n      " + call + "\n#endif")
+        text = text.replace(call, replacement)
+    memory.write_text(text)
+    return source
+
+
 def build_hypre(
     args: argparse.Namespace,
     features: Mapping[str, bool],
@@ -651,6 +760,7 @@ def build_hypre(
         args,
         build_type,
         c_compiler=args.cc,
+        cxx_compiler=args.cxx,
     ) + [
         f"-DHYPRE_ENABLE_MPI={'ON' if features['mpi'] else 'OFF'}",
         f"-DHYPRE_ENABLE_OPENMP={'ON' if features['openmp'] else 'OFF'}",
@@ -678,7 +788,7 @@ def build_hypre(
         options.append(f"-DCMAKE_HIP_ARCHITECTURES={args.hip_arch}")
     cmake_build_install(
         args,
-        args.hypre_source / "src",
+        prepare_hypre_source(args, features) / "src",
         args.work_dir / "hypre",
         options,
         env,
@@ -729,8 +839,8 @@ def build_sundials(
     options = common_cmake_args(
         args,
         build_type,
-        c_compiler=mpicc if features["mpi"] else args.cc,
+        c_compiler=args.cc,
-        cxx_compiler=mpicxx if features["mpi"] else args.cxx,
+        cxx_compiler=args.cxx,
     ) + [
         f"-DSUNDIALS_ENABLE_MPI={'ON' if features['mpi'] else 'OFF'}",
         f"-DSUNDIALS_ENABLE_OPENMP={'ON' if features['openmp'] else 'OFF'}",
@@ -882,6 +992,15 @@ def build_libceed(
             f"EMSCRIPTEN={'1' if args.host_system == 'emscripten' else ''}",
         ]
     )
+    if features["cuda"]:
+        nvcc_flags = ['-ccbin', args.cxx, '-std=c++17',
+                      '-Xcompiler', optimization, '-Xcompiler', '-fPIC']
+        if ceed_cuda_arch:
+            nvcc_flags.append('-arch=' + ceed_cuda_arch)
+        command.extend([
+            f"NVCC={args.cuda_compiler}",
+            f"NVCCFLAGS={shlex.join(nvcc_flags)}",
+        ])
     if args.host_system not in {"darwin", "emscripten"}:
         command.append("LDLIBS=-ldl")
     run(command, env=local_env)
@@ -2204,6 +2323,13 @@ def main() -> int:
                 "dependency_prefix and the private output prefix must be disjoint"
             )
     features = parse_features(args.feature)
+    if features["cuda"]:
+        args.cuda_compiler = str(Path(args.cuda_compiler or shutil.which("nvcc") or "nvcc").resolve())
+        args.cc = shutil.which(args.cc) or args.cc
+        args.cxx = shutil.which(args.cxx) or args.cxx
+        print(check_toolchain(args.cuda_compiler, args.cc, args.cxx,
+                              standard="20" if features["raja"] or features["umpire"] else "17",
+                              c_args=args.c_arg, cxx_args=args.cxx_arg), flush=True)
     build_type = cmake_build_type(args.buildtype)
     compiler_identity = {
         "cc": output([*args.cc_launcher, args.cc, "--version"]).splitlines()[0],
@@ -2218,6 +2344,10 @@ def main() -> int:
             (args.source / "CMakeLists.txt").read_bytes()
         ).hexdigest(),
     }
+    if features["cuda"]:
+        fingerprint_data["cuda_version"] = output([args.cuda_compiler, "--version"])
+        fingerprint_data["cuda_check_sha256"] = hashlib.sha256(
+            Path(__file__).with_name("check_cuda_toolchain.py").read_bytes()).hexdigest()
     if features["ceed"]:
         ceed_jit_source = Path(__file__).with_name(
             "ceed_jit_source_root_relocatable.c"
@@ -2249,9 +2379,17 @@ def main() -> int:
     args.prefix.mkdir(parents=True, exist_ok=True)
 
     env = dict(os.environ)
+    if features["cuda"]:
+        env["CUDA_HOME"] = str(Path(args.cuda_compiler).parent.parent)
+        env["CUDA_PATH"] = env["CUDA_HOME"]
+        env["CUDACXX"] = args.cuda_compiler
+        env["CUDAHOSTCXX"] = args.cxx
+        env["NVCC_CCBIN"] = args.cxx
     path_entries = [os.fspath(args.prefix / "bin")]
     if args.dependency_prefix:
         path_entries.append(os.fspath(Path(args.dependency_prefix) / "bin"))
+    if features["cuda"]:
+        path_entries.append(str(Path(args.cuda_compiler).parent))
     path_entries.append(env.get("PATH", ""))
     env["PATH"] = os.pathsep.join(path_entries)
     env["CMAKE_PREFIX_PATH"] = os.pathsep.join(

tools/check_cuda_toolchain.py

@@ -0,0 +1,72 @@
+#!/usr/bin/env python3
+"""Probe the installed toolkit instead of maintaining a CUDA/compiler version table."""
+from __future__ import annotations
+
+import argparse
+import os
+from pathlib import Path
+import shlex
+import subprocess
+import tempfile
+
+
+def check_toolchain(nvcc: str, cc: str, cxx: str, *, standard: str = '17',
+                    c_args: list[str] | None = None,
+                    cxx_args: list[str] | None = None) -> str:
+    version = subprocess.run([nvcc, '--version'], text=True, capture_output=True, check=True).stdout.strip()
+    # Do not permit ambient flags to bypass the installed toolkit's version checks.
+    env = dict(os.environ)
+    for key in ('NVCC_PREPEND_FLAGS', 'NVCC_APPEND_FLAGS', 'NVCC_CCBIN'):
+        env.pop(key, None)
+    with tempfile.TemporaryDirectory(prefix='mfem-cuda-check-') as directory:
+        root = Path(directory)
+        source = root / 'probe.cu'
+        source.write_text('''#include <cuda_runtime.h>
+#include <thrust/device_vector.h>
+#include <cmath>
+__global__ void probe(int *p) { *p = 1; }
+int main() { int *p = nullptr; cudaMalloc(&p, sizeof(int));
+probe<<<1, 1>>>(p); return (int)cudaFree(p); }
+''')
+        for label, host, flags in [('C', cc, c_args or []), ('C++', cxx, cxx_args or [])]:
+            command = [nvcc, '--compiler-bindir', host, '--std=c++' + standard,
+                       str(source), '-o', str(root / ('probe.exe' if os.name == 'nt' else 'probe'))]
+            if label == 'C':
+                command.append('-c')
+            for flag in flags:
+                command += ['-Xcompiler', flag]
+            result = subprocess.run(command, env=env, text=True, capture_output=True)
+            if result.returncode:
+                raise RuntimeError(
+                    f'CUDA compatibility check failed for the selected {label} compiler: {host}\n'
+                    f'{version}\nCommand: {shlex.join(command)}\n'
+                    f'{result.stdout}{result.stderr}\n'
+                    'Select a C/C++ toolchain supported by this CUDA toolkit using CC and CXX '
+                    '(for example CC=gcc-15 CXX=g++-15), or install a compatible toolkit. '
+                    'Use a fresh Meson build directory when changing compilers; --reconfigure '
+                    'does not change cached compilers. For a CPU build use -Dmfem_cuda=disabled. '
+                    'The probe compiles and links but does not require a running GPU.'
+                )
+    return f'CUDA compatibility checks passed: C={cc}, C++={cxx}, CUDA C++{standard}\n{version}'
+
+
+def main() -> int:
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument('--nvcc', required=True)
+    parser.add_argument('--cc', required=True)
+    parser.add_argument('--cxx', required=True)
+    parser.add_argument('--standard', default='17')
+    parser.add_argument('--c-arg', action='append', default=[])
+    parser.add_argument('--cxx-arg', action='append', default=[])
+    args = parser.parse_args()
+    try:
+        print(check_toolchain(args.nvcc, args.cc, args.cxx, standard=args.standard,
+                              c_args=args.c_arg, cxx_args=args.cxx_arg))
+    except (RuntimeError, OSError, subprocess.SubprocessError) as error:
+        print(error)
+        return 1
+    return 0
+
+
+if __name__ == '__main__':
+    raise SystemExit(main())