[BUG]: cuda_bindings/examples globalToShmemAsyncCopy_test.py "catastrophic error" is masked by pytest skipped

[rwgk]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Silent Failure

Component

cuda.bindings

Describe the bug

The "catastrophic error" below was overlooked in the the preparation for SWQA testing, because it's masked as skipped.

How to Reproduce

(WinDevVenv) PS C:\Users\rgrossekunst\forked\cuda-python-private\unreleased-13.0\cuda_bindings> python -m pytest -s -vv examples\
================================================= test session starts =================================================
platform win32 -- Python 3.13.5, pytest-8.4.1, pluggy-1.6.0 -- C:\Users\rgrossekunst\forked\cuda-python-private\unreleased-13.0\WinDevVenv\scripts\python.exe
cachedir: .pytest_cache
benchmark: 5.1.0 (defaults: timer=time.perf_counter disable_gc=False min_rounds=5 min_time=0.000005 max_time=1.0 calibration_precision=10 warmup=False warmup_iterations=100000)
rootdir: C:\Users\rgrossekunst\forked\cuda-python-private\unreleased-13.0\cuda_bindings\examples
configfile: pytest.ini
plugins: benchmark-5.1.0
collecting ... sourceCode.cu(2): catastrophic error: cannot open source file "cuda/barrier"
      #include <cuda/barrier>
                             ^

1 catastrophic error detected in the compilation of "sourceCode.cu".
Compilation terminated.

CUDA error code=6(b'NVRTC_ERROR_COMPILATION')
collected 1 item

examples\3_CUDA_Features\globalToShmemAsyncCopy_test.py::main SKIPPED (Automation filter against incompatible
kernel)

================================================= 1 skipped in 0.49s ==================================================

Expected behavior

We need to be more selective in catching and processing this kind of error.

Operating System

No response

nvidia-smi output

No response

[rwgk]

This is what we have at the moment (globalToShmemAsyncCopy_test.py):

def checkKernelCompiles():
    kernel_headers = """\
    #line __LINE__
    #if __CUDA_ARCH__ >= 700
    #include <cuda/barrier>
    #endif
    #include <cooperative_groups.h>
    #include <cooperative_groups/reduce.h>
    #include <cuda/pipeline>
    """
    try:
        common.KernelHelper(kernel_headers, findCudaDevice())
    except:
        # Filters out test from automation for two reasons
        # 1. Headers are not found
        # 2. Incompatible device
        return False
    return True

In common.KernelHelper.__init__() near the top of the function:

        CUDA_HOME = os.getenv("CUDA_HOME")
        if CUDA_HOME is None:
            CUDA_HOME = os.getenv("CUDA_PATH")
        if CUDA_HOME is None:
            raise RuntimeError("Environment variable CUDA_HOME or CUDA_PATH is not set")

[rwgk]

Idea for making this less likely to mask actual errors:

• If "Environment variable CUDA_HOME and CUDA_PATH are both not set":
  • either pytest error, or
  • pytest skipped with CUDA_HOME/CUDA_PATH not set as the SKIPPED message.
• Otherwise (if CUDA_HOME is set):
  • Query the compute capability of the GPU, generate pytest skipped if too low

In all other cases, to not skip the test, so that real errors can bubble up properly (pytest error).

[rwgk]

The idea in my previous comment was implemented on an internal branch, locally tested, and forwarded to SWQA for retesting.

[leofang]

Closing, this is fixed as part of #792.