[Feat] Add an experimental cuda nvvm backend

[sjfeng1999]

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[Copilot]

Pull request overview

Adds an experimental CUDA/NVVM backend stack to FlyDSL, extending the compiler/runtime and MLIR lowering pipeline so kernels can be lowered through NVVM and executed via a CUDA driver-API JIT runtime, with corresponding tests/examples and backend-aware test running.

Changes:

• Introduces the NVVM backend: FlyNVVM dialect + SM80 atom types, Fly→NVVM conversion pass, CUDA JIT runtime wrappers, and Python bindings.
• Adds Python-side CUDA support: compile backend (cuda), CUDA device runtime discovery, and fx.nvvm expr namespace.
• Updates build/test/docs: backend selection in CMake/scripts, backend-gated MLIR FileCheck, new unit tests, and reorganized examples.

Reviewed changes

Copilot reviewed 48 out of 53 changed files in this pull request and generated 5 comments.

Show a summary per file

File	Description
tests/unit/test_universal_atomic.py	Makes ROCDL-only sync-scope parametrization collection-safe for non-ROCDL builds.
tests/unit/test_tdm_mcast_add_gfx1250.py	Skips module import on non-ROCDL builds to avoid collection-time failures.
tests/unit/test_gfx1250_atoms.py	Adjusts markers to reflect target-dialect requirements.
tests/unit/test_cuda_backend.py	Adds backend-agnostic unit coverage for CUDA backend discovery and pipeline shape.
tests/unit/test_backend_cmake_defaults.py	Updates expected allowed backend strings to include nvvm.
tests/mlir/Conversion/fly-to-nvvm/mma_atom.mlir	Adds FileCheck coverage for NVVM MMA atom lowering.
tests/mlir/Conversion/fly-to-nvvm/copy_atom.mlir	Adds FileCheck coverage for NVVM cp.async + ldmatrix lowering.
scripts/run_tests.sh	Makes pytest/example selection backend-aware; skips MLIR tests for disabled backends.
scripts/build.sh	Allows passing FLYDSL_BACKENDS through to CMake from the build script.
README.md	Updates documented repo layout and example locations for ROCm/CUDA split.
python/mlir_flydsl/dialects/FlyNVVM.td	Adds TableGen entry for FlyNVVM Python bindings generation.
python/mlir_flydsl/dialects/fly_nvvm.py	Adds Python dialect module for FlyNVVM bindings.
python/mlir_flydsl/CMakeLists.txt	Builds/installs FlyNVVM Python bindings and extension when NVVM backend is enabled.
python/flydsl/runtime/device_runtime/cuda.py	Implements CUDA driver-API based runtime queries and arch detection.
python/flydsl/runtime/device_runtime/init.py	Registers CUDA runtime kind mapping and exports CUDA runtime class.
python/flydsl/expr/nvvm/universal.py	Adds Python factories for NVVM atom types (mma.sync, cp.async, ldmatrix).
python/flydsl/expr/nvvm/init.py	Adds fx.nvvm namespace with cp.async group barriers and atom exports.
python/flydsl/expr/init.py	Adds lazy backend-module entry for nvvm.
python/flydsl/compiler/jit_function.py	Generalizes ISA dumping to reuse the backend’s gpu-module-to-binary fragment.
python/flydsl/compiler/backends/cuda.py	Adds CUDA compile backend pipeline definition and toolkit detection.
lib/Runtime/CUDA/FlyCudaRuntimeWrappers.cpp	Implements CUDA mgpu* JIT runtime wrappers compatible with the host launcher IR.
lib/Runtime/CUDA/CMakeLists.txt	Builds CUDA JIT runtime and enforces CUDA 12+ header/toolkit requirements.
lib/Runtime/CMakeLists.txt	Adds CUDA runtime subdirectory when nvvm backend is enabled.
lib/Dialect/FlyNVVM/SM80/MmaAtom.cpp	Implements SM80 mma.sync.aligned atom type behavior and lowering.
lib/Dialect/FlyNVVM/SM80/CopyAtom.cpp	Implements SM80 cp.async and SM75+ ldmatrix copy atoms.
lib/Dialect/FlyNVVM/Dialect.cpp	Registers FlyNVVM dialect types.
lib/Dialect/FlyNVVM/CMakeLists.txt	Builds the FlyNVVM dialect library.
lib/Conversion/FlyToNVVM/FlyToNVVM.cpp	Adds Fly→NVVM conversion pass implementation (type conversion + op lowering).
lib/Conversion/FlyToNVVM/CMakeLists.txt	Builds the Fly→NVVM conversion library.
lib/CAPI/Dialect/FlyNVVM/FlyNVVMDialect.cpp	Exposes FlyNVVM dialect/pass registration via C API.
lib/CAPI/Dialect/FlyNVVM/CMakeLists.txt	Builds the FlyNVVM C API library.
lib/Bindings/Python/FlyNVVMExtension.cpp	Adds nanobind extension for FlyNVVM atom type constructors.
include/flydsl/Dialect/FlyNVVM/IR/MmaAtom.td	Declares FlyNVVM SM80 MMA atom type in TableGen.
include/flydsl/Dialect/FlyNVVM/IR/Dialect.td	Declares FlyNVVM dialect and common typedef base classes.
include/flydsl/Dialect/FlyNVVM/IR/Dialect.h	Adds C++ header for FlyNVVM dialect/types.
include/flydsl/Dialect/FlyNVVM/IR/CopyAtom.td	Declares FlyNVVM SM80/SM75 copy atom types in TableGen.
include/flydsl/Dialect/FlyNVVM/IR/CMakeLists.txt	Adds TableGen rules for FlyNVVM dialect/types.
include/flydsl/Dialect/FlyNVVM/IR/Atom.td	Aggregates FlyNVVM atom type definitions.
include/flydsl/Dialect/FlyNVVM/CMakeLists.txt	Wires FlyNVVM include subdirectory into the build.
include/flydsl/Conversion/FlyToNVVM/Passes.td	Declares the convert-fly-to-nvvm pass.
include/flydsl/Conversion/FlyToNVVM/FlyToNVVM.h	Declares Fly→NVVM pass entry points.
include/flydsl/Conversion/FlyToNVVM/CMakeLists.txt	Adds TableGen rules for Fly→NVVM pass declarations.
include/flydsl-c/FlyNVVMDialect.h	Adds C API header for FlyNVVM dialect/pass/plugin registration.
examples/rocm/03-preshuffle_gemm.py	Adds ROCm preshuffle GEMM example under new examples/rocm/ layout.
examples/rocm/02-MFMA.py	Adds ROCm MFMA example under new examples/rocm/ layout.
examples/rocm/01-BufferCopy.py	Adds ROCm BufferCopy example under new examples/rocm/ layout.
examples/notebooks/README.md	Updates notebook README to remove outdated example references.
examples/cuda/01-MmaSync.py	Adds CUDA example for NVVM mma.sync-aligned atom usage.
examples/02-gather_scatter.py	Adds target-neutral gather/scatter example based on offset tensors.
cmake/FlyDSLBackends.cmake	Allows nvvm as an enabled backend option in the CMake cache.
cmake/backends/nvvm.cmake	Adds NVVM backend descriptor to self-register build properties.
CLAUDE.md	Updates contributor docs to include NVVM backend and example layout.
.gitignore	Removes some ignored artifact patterns (notably benchmark CSV patterns).

Suppressed comments (1)

README.md:364

• This paragraph still references the old example filenames (02-tiledCopy.py, 03-tiledMma.py, 04-preshuffle_gemm.py). Update it to match the current examples/rocm/ filenames so users can follow the references.

See `examples/rocm/` for AMD examples including tiled copy (`02-tiledCopy.py`), tiled MMA (`03-tiledMma.py`), and preshuffle GEMM (`04-preshuffle_gemm.py`), and `examples/cuda/` for the NVIDIA NVVM examples.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.