Warn and skip native configuration on unsupported devices (#1502)

## Summary

Follow up on #1499 so consumers can call Liger's public configuration
API without duplicating GPU eligibility checks. This unblocks removing
those checks from verl-project/verl#7945.

- Keep LCK strict: `liger_cute_kernels.configure` raises the exported
`UnsupportedDeviceError` when it cannot support the requested device,
without loading or initializing native resources. Successful LCK setup
retains its `None` return.
- Catch that specific exception in
**`liger_kernel.ops.configure.configure`**, emit a warning, and return
`False`. Warning/fallback policy lives in Liger-Kernel, not LCK or
downstream consumers.
- Apply the same contract to FLSCE and MoE configuration. On unsupported
NVIDIA architectures, retain collective plan validation before LCK
raises the typed error.
- Return `True` from the Liger-Kernel frontend after successful native
setup. Older native packages without the new exception export remain
usable; their genuine setup errors still propagate.
- Keep actual setup failures visible: invalid capacities, mismatched
collective plans, dependency failures, and conflicting NVSHMEM ownership
are not converted into fallback results.

Missing optional LCK already warns and returns `False`; this extends
that behavior to unsupported hardware. Operator frontends continue to
own implementation selection. Skipping setup does not make an explicitly
selected native-only operator usable, and this PR does not change CUDA
kernels or operator fallback algorithms.

## Validation

- 143 focused tests passed on each of H200 and B300: 89 Liger
configuration contract tests and 54 VeRL integration regressions.
- Coverage includes supported Hopper/Blackwell setup, direct LCK
exceptions without warnings, public Liger exception-to-warning
conversion, both operator sections, no native initialization on skipped
setup, collective hardware disagreement, invalid capacity rejection,
legacy native return compatibility, and VeRL delegation without hardware
probes.
- Runtime, value, and OS errors remain unchanged through the public
boundary with both old and new native package interfaces.
- Unsupported architectures and non-NVIDIA devices are covered with
host-side mocks, not physical hardware runs.
- Repository Ruff checks and formatting passed.

Compatible Python and native packages containing this change must be
released before downstream users can rely on the new unsupported-device
behavior. The broader VeRL performance campaign was not rerun.

---------

Co-authored-by: kolehma8 <[email protected]>
Co-authored-by: Copilot <[email protected]>

liger_cute_kernels/README.md

@@ -26,6 +26,12 @@ Use `liger_cute_kernels.configure(...)` to initialize NVSHMEM once and prepare
 FLSCE and/or MoE through one public setup boundary. This API requires the LCK
 Python package containing this extension; older wheels do not expose it.
 Use a matching native build exposing the TP and EP context entry points.
+Successful native setup returns `None`. Unsupported devices raise
+`liger_cute_kernels.UnsupportedDeviceError` without loading or initializing
+native resources. The `liger_kernel.ops.configure.configure` frontend catches
+that specific error, emits a warning, and returns `False`. Consumers should
+use that Liger-Kernel frontend rather than inspect device capabilities or
+call LCK directly; operator frontends continue to own fallback selection.
 
 ```python
 import torch
@@ -211,11 +217,16 @@ export this configuration API.
 Ordinary operator imports do not expose or load this optional setup API.
 Importing `liger_kernel.ops.configure`, importing its function or classes, and
 constructing configuration objects all work without LCK installed. LCK is
-loaded only when `configure` is called. If LCK is absent, the function emits a
+loaded only when `configure` is called. If LCK is absent or does not support
+the device, the function emits a
 warning and returns `False` without changing the runtime; other Liger kernels
-remain usable. Successful native setup returns `True`. Failures inside an
-installed LCK package still propagate. All bootstrap ranks must use a consistent
-installation. Skipping setup does not enable an explicitly selected native-only
+remain usable. This warning conversion belongs to Liger-Kernel: LCK itself
+raises `UnsupportedDeviceError`. Successful setup through the Liger-Kernel
+frontend returns `True`. Other failures inside an
+installed LCK package still propagate, including invalid capacity, inconsistent
+collective plans, and conflicting NVSHMEM ownership. All bootstrap ranks must
+use a consistent installation and matching hardware. Skipping setup does not
+enable an explicitly selected native-only
 operator or change backend selection.
 The operations' existing `.apply(...)` interfaces and fallback dispatch remain
 unchanged.

liger_cute_kernels/liger_cute_kernels/__init__.py

@@ -23,6 +23,7 @@ from __future__ import annotations
 
 from .configuration import FusedLinearCrossEntropyConfig
 from .configuration import MoEConfig
+from .configuration import UnsupportedDeviceError
 from .configuration import configure
 
-__all__ = ["FusedLinearCrossEntropyConfig", "MoEConfig", "configure"]
+__all__ = ["FusedLinearCrossEntropyConfig", "MoEConfig", "UnsupportedDeviceError", "configure"]

liger_cute_kernels/liger_cute_kernels/configuration.py

@@ -12,6 +12,10 @@ import torch.distributed as dist
 _OperatorSpec = tuple[tuple[int, ...], tuple[int, ...]]
 
 
+class UnsupportedDeviceError(RuntimeError):
+    """The requested device cannot use LCK's native runtime."""
+
+
 @dataclass(frozen=True)
 class FusedLinearCrossEntropyConfig:
     """Maximum local FLSCE problem; ``group`` names a registered process group.
@@ -216,6 +220,11 @@ def configure(
 ) -> None:
     """Configure the shared NVSHMEM runtime and optional FLSCE/MoE resources.
 
+    Unsupported devices raise ``UnsupportedDeviceError`` without initializing
+    native resources. Liger-Kernel's public configuration frontend converts
+    this error into a warning and leaves fallback selection to its operators.
+    Invalid configurations and native setup failures still raise.
+
     Collective on ``bootstrap_group`` (default: torch WORLD), including on
     repeated calls. All ranks must supply the same sections and capacities;
     their local TP/EP groups must form uniform strided partitions of that world.
@@ -252,15 +261,13 @@ def configure(
     from . import tvm_ffi
 
     global _configuration
+    device = torch.device(device if device is not None else "cuda")
+    if device.type != "cuda" or not torch.cuda.is_available() or torch.version.hip is not None:
+        raise UnsupportedDeviceError("LCK requires an NVIDIA CUDA device")
     if not dist.is_initialized():
         raise RuntimeError("torch.distributed must be initialized before configure")
-    if not torch.cuda.is_available() or torch.version.hip is not None:
-        raise RuntimeError("LCK configure requires an NVIDIA CUDA device")
     bootstrap_group = bootstrap_group if bootstrap_group is not None else dist.group.WORLD
     bootstrap_ranks = _ranks(bootstrap_group)
-    device = torch.device(device if device is not None else "cuda")
-    if device.type != "cuda":
-        raise ValueError("LCK configure requires a CUDA device")
     if device.index is None:
         device = torch.device("cuda", torch.cuda.current_device())
 
@@ -277,19 +284,18 @@ def configure(
                 torch.cuda.get_device_capability(),
                 torch.cuda.get_device_properties(device).multi_processor_count,
             )
-            if hardware[0][0] not in (9, 10):
-                raise ValueError("LCK configure supports Hopper and Blackwell devices")
             groups, group_specs = _describe_process_groups(process_groups)
             flsce_spec = _describe_flsce(flsce, groups, bootstrap_group)
             moe_spec = _describe_moe(moe, groups, bootstrap_group)
-            if _configuration is None:
+            native_supported = hardware[0] == (9, 0) or hardware[0][0] == 10
+            if _configuration is None and native_supported:
                 if tvm_ffi.nvshmem_is_initialized():
                     raise RuntimeError(
                         "NVSHMEM is already initialized outside LCK configure; sharing with an unmanaged "
                         "runtime such as DeepEP V1 RDMA is unsupported"
                     )
                 tvm_ffi._load_module()
-            else:
+            elif _configuration is not None:
                 if _configuration.bootstrap_ranks != bootstrap_ranks or _configuration.device != device:
                     raise ValueError("LCK bootstrap group and device cannot change after configuration")
                 _check_existing("flsce", flsce_spec, _configuration.flsce)
@@ -342,6 +348,9 @@ def configure(
         for index in range(len(names)):
             _check_group_partition([plan[5][index][1] for plan in plans], bootstrap_ranks)
 
+        if not native_supported:
+            raise UnsupportedDeviceError("LCK supports Hopper and Blackwell devices")
+
         torch.cuda.synchronize(device)
         if _configuration is None:
             nvshmem.init_from_pg(bootstrap_group)

liger_cute_kernels/test/test_configuration.py

@@ -21,6 +21,7 @@ import torch
 
 from liger_cute_kernels import FusedLinearCrossEntropyConfig
 from liger_cute_kernels import MoEConfig
+from liger_cute_kernels import UnsupportedDeviceError
 from liger_cute_kernels import configuration
 from liger_cute_kernels import configure
 from liger_cute_kernels import nvshmem
@@ -75,7 +76,7 @@ def runtime(monkeypatch):
 
 
 def test_shared_runtime_configures_both_operators_once(runtime):
-    runtime.configure(flsce=FLSCE, moe=MOE)
+    assert runtime.configure(flsce=FLSCE, moe=MOE) is None
     runtime.configure(flsce=FLSCE, moe=MOE)
     assert runtime.calls == [
         ("init", runtime.world),
@@ -620,7 +621,9 @@ def test_public_configuration_data_matches_lck_schema(public_configuration):
         assert public_fields == native_fields
 
 
-def test_public_module_configures_both_operators(runtime, public_configuration):
+@pytest.mark.parametrize("capability", [(9, 0), (10, 0), (10, 3)])
+def test_public_module_configures_both_operators(runtime, public_configuration, monkeypatch, capability):
+    monkeypatch.setattr(torch.cuda, "get_device_capability", lambda device=None: capability)
     assert (
         public_configuration.configure(
             process_groups={"tp": runtime.world, "ep": runtime.world},
@@ -635,6 +638,75 @@ def test_public_module_configures_both_operators(runtime, public_configuration):
     assert runtime.calls[-1] == ("moe", (2048, 512, 8, 2, 1, 2, 1, 19, 0))
 
 
+@pytest.mark.parametrize("capability", [(8, 0), (8, 9), (9, 1), (12, 0)])
+def test_public_configuration_skips_unsupported_architecture(
+    runtime, public_configuration, monkeypatch, capability, recwarn
+):
+    monkeypatch.setattr(torch.cuda, "get_device_capability", lambda device=None: capability)
+    monkeypatch.setattr(tvm_ffi, "_load_module", lambda: pytest.fail("must not load native code"))
+    monkeypatch.setattr(tvm_ffi, "nvshmem_is_initialized", lambda: pytest.fail("must not inspect native runtime"))
+    with pytest.raises(UnsupportedDeviceError, match="Hopper and Blackwell"):
+        runtime.configure(flsce=FLSCE, moe=MOE)
+    assert not recwarn
+    with pytest.warns(UserWarning, match="native configuration was skipped"):
+        assert (
+            public_configuration.configure(
+                bootstrap_group=runtime.world,
+                device="cuda:0",
+                flsce=public_configuration.FusedLinearCrossEntropyConfig(**asdict(FLSCE)),
+                moe=public_configuration.MoEConfig(**asdict(MOE)),
+            )
+            is False
+        )
+    assert runtime.calls == []
+    assert configuration._configuration is None
+
+
+@pytest.mark.parametrize(
+    ("device", "cuda_available", "hip"),
+    [("cpu", True, None), ("xpu:0", True, None), ("cuda:0", False, None), ("cuda:0", True, "6.3")],
+)
+def test_public_configuration_skips_non_nvidia_devices(
+    runtime, public_configuration, monkeypatch, device, cuda_available, hip, recwarn
+):
+    monkeypatch.setattr(torch.cuda, "is_available", lambda: cuda_available)
+    monkeypatch.setattr(torch.version, "hip", hip)
+    monkeypatch.setattr(torch.cuda, "device", lambda device: pytest.fail("must not enter a CUDA context"))
+    with pytest.raises(UnsupportedDeviceError, match="NVIDIA CUDA device"):
+        configure(device=device)
+    assert not recwarn
+    with pytest.warns(UserWarning, match="native configuration was skipped"):
+        assert public_configuration.configure(device=device) is False
+    assert runtime.calls == []
+    assert configuration._configuration is None
+
+
+def test_unsupported_architecture_still_rejects_invalid_capacity(runtime, monkeypatch):
+    monkeypatch.setattr(torch.cuda, "get_device_capability", lambda device=None: (8, 0))
+    with pytest.raises(RuntimeError, match="max_tokens must be a positive int32"):
+        runtime.configure(flsce=replace(FLSCE, max_tokens=0))
+    assert runtime.calls == []
+
+
+def test_mixed_architectures_reject_collectively_without_initialization(runtime, monkeypatch):
+    monkeypatch.setattr(torch.cuda, "get_device_capability", lambda device=None: (8, 0))
+
+    def gather(output, value, group):
+        peer = list(value)
+        peer[3] = ((9, 0), 132)
+        output[:] = [value, tuple(peer)]
+
+    monkeypatch.setattr(configuration.dist, "all_gather_object", gather)
+    with pytest.raises(RuntimeError, match="matching hardware"):
+        runtime.configure(flsce=FLSCE, moe=MOE)
+    assert runtime.calls == []
+
+
+def test_public_configuration_preserves_legacy_native_success(public_configuration, monkeypatch):
+    monkeypatch.setattr(importlib, "import_module", lambda name: SimpleNamespace(configure=lambda **kwargs: None))
+    assert public_configuration.configure() is True
+
+
 def test_backend_discovery_does_not_import_optional_lck(public_ops):
     for name in ("configure", "FusedLinearCrossEntropyConfig", "MoEConfig"):
         assert not hasattr(public_ops, name)
@@ -690,14 +762,21 @@ def test_configuration_preserves_dependency_import_errors(public_configuration,
     assert caught.value is error
 
 
-def test_configuration_preserves_native_setup_errors(public_configuration, monkeypatch):
-    error = RuntimeError("native setup failed")
+@pytest.mark.parametrize("error_type", [RuntimeError, ValueError, OSError])
+@pytest.mark.parametrize("typed_error_available", [False, True])
+def test_configuration_preserves_native_setup_errors(
+    public_configuration, monkeypatch, error_type, typed_error_available
+):
+    error = error_type("native setup failed")
 
     def fail(**kwargs):
         raise error
 
-    monkeypatch.setattr(importlib, "import_module", lambda name: SimpleNamespace(configure=fail))
-    with pytest.raises(RuntimeError) as caught:
+    lck = SimpleNamespace(configure=fail)
+    if typed_error_available:
+        lck.UnsupportedDeviceError = UnsupportedDeviceError
+    monkeypatch.setattr(importlib, "import_module", lambda name: lck)
+    with pytest.raises(error_type) as caught:
         public_configuration.configure()
     assert caught.value is error
 

src/liger_kernel/ops/configure.py

@@ -1,8 +1,8 @@
 """Optional native runtime setup, separate from ordinary operator imports.
 
 Importing this module and constructing configuration data do not require LCK.
-When LCK is absent, ``configure`` warns and returns ``False`` without changing
-the runtime.
+When LCK is absent or the device is unsupported, ``configure`` warns and
+returns ``False`` without changing the runtime.
 """
 
 from __future__ import annotations
@@ -67,10 +67,12 @@ def configure(
 ) -> bool:
     """Collectively configure native resources using LCK's shared runtime API.
 
-    Returns ``False`` with a warning when optional LCK is absent, leaving
-    non-LCK kernels available. Returns ``True`` after native setup succeeds.
-    Failures inside an installed LCK package propagate rather than silently
-    disabling it. All bootstrap ranks must use a consistent installation.
+    Returns ``False`` with a warning when optional LCK is absent or its native
+    kernels do not support the device, leaving non-LCK kernels available.
+    Converts LCK's ``UnsupportedDeviceError`` into a warning and ``False``.
+    Returns ``True`` after native setup succeeds. Other failures inside an
+    installed LCK package propagate rather than silently disabling it.
+    All bootstrap ranks must use a consistent installation.
     See ``liger_cute_kernels.configure`` for collective setup, named process
     groups, capacity, and lifetime requirements.
     """
@@ -83,11 +85,21 @@ def configure(
             stacklevel=2,
         )
         return False
-    lck.configure(
-        process_groups=process_groups,
-        bootstrap_group=bootstrap_group,
-        device=device,
-        flsce=lck.FusedLinearCrossEntropyConfig(**asdict(flsce)) if flsce is not None else None,
-        moe=lck.MoEConfig(**asdict(moe)) if moe is not None else None,
-    )
+    # Older native packages have no typed eligibility error; their errors still propagate.
+    unsupported_device_errors = (lck.UnsupportedDeviceError,) if hasattr(lck, "UnsupportedDeviceError") else ()
+    try:
+        lck.configure(
+            process_groups=process_groups,
+            bootstrap_group=bootstrap_group,
+            device=device,
+            flsce=lck.FusedLinearCrossEntropyConfig(**asdict(flsce)) if flsce is not None else None,
+            moe=lck.MoEConfig(**asdict(moe)) if moe is not None else None,
+        )
+    except unsupported_device_errors as exc:
+        warnings.warn(
+            f"{exc}; native configuration was skipped. Liger's non-LCK kernels remain available.",
+            UserWarning,
+            stacklevel=2,
+        )
+        return False
     return True