[DarkSharpness]

(auto generated by Claude)

Motivation

TensorMatcher::with_device<Codes...>(sym) and with_dtype<Ts...>(sym) called set_options even with an empty pack, which replaced the symbol's options with an empty list (= anything allowed). The common launcher idiom device.set_options<kDLCUDA>(); TensorMatcher(...).with_device(device).verify(x); therefore never checked the device type, so a CPU tensor passed. About 200 of 260 SymbolicDevice declarations under python/sglang/kernels/jit/csrc/ used this idiom. The same bug affected the dtype restriction in distributed/nvlink_comm.cuh.

Introduced in #14570, which added the set_options call to the symbol overloads; before that (#13453) they left the symbol's options untouched, so every launcher written with this idiom since has had an unchecked device type.

Modifications

• tensor.h: with_device(sym) / with_dtype(sym) only override options when the pack is non-empty. set_options now re-checks a value that is already bound, so a runtime DLDevice / DLDataType passed together with a type list (e.g. .with_device<kDLGPU>(x.device())) is checked too; that path was also unchecked before.
• Replace hard-coded kDLCUDA / kDLCUDAHost with kDLGPU / kDLGPUHost in JIT launchers and headers (jit/csrc, jit/include, kda_kernels/csrc), so the check is correct on ROCm. moe/paged_experts.cuh needed using sglang::kDLGPU; because it only imports sglang::host.
• Left as kDLCUDA because they are CUDA-only and don't include sgl_kernel/utils.cuh: jit/csrc/moe/tvm_ffi_utils.h (cudaGetDevice + CUDA stream helper), the vendored ops/lora/moe/trtllm_lora_temp/data/csrc/trtllm_fused_moe_kernel_launcher.cu, and the device-name table in tensor.h.
• add-jit-kernel skill examples updated to kDLGPU, plus a note that options set on a symbol stay in force.
• New test/registered/unit/kernels/test_tensor_matcher.py: JIT-compiles a small launcher and checks that a CPU tensor or a wrong dtype is rejected through both the symbol path and the runtime-value path. All 4 cases fail on the old header.

Fallout

No launcher started rejecting its inputs: no test anywhere hit not in the allowed options, so no call site or declared option needed changing.

Note: moe/expert_specialization/es_sm100_mxfp8_blockscaled_{group_quant,moe_group_gemm}.cuh are not referenced by any JIT loader and already fail to compile on main (host is not visible there). They were only touched by the rename.

Accuracy / tests

• All test/registered/kernels/ops/** and test/registered/unit/kernels/** on 1x B200. Every failure reproduces on main with identical counts (paged_mqa_metadata, q8kv8_sparse_prefill_backend, tilelang_dsa_zero_rope, model_fast_paths, diffusion test_rope, hy4_decode_kernels, lanczos, moe_topk_softmax, per_token_group_quant_8bit_v2). The rest are multi-GPU or ROCm-only tests that can't run on a single CUDA GPU.
• Multi-GPU communication tests on 4x H200: all_reduce_residual_norm, deterministic_reduce_scatter, fp4_allgather_moe_reduction, tp_qknorm, custom_all_reduce, symm_mem_all_gather pass. ar_fusion and beam_all_reduce need 8 GPUs (left to CI).
• Every changed header compiled standalone for sm_100a on this branch vs main: no new failures.
• ROCm not run locally; relying on AMD CI.

Follow-up (not in this PR)

Hand-written checks that should become matcher expressions (RuntimeCheck(t.dim() == ...), stride / contiguity / alignment checks, raw dtype .code / .bits, TVM_FFI_ICHECK) remain in about 90 files under jit/csrc, e.g. distributed/custom_all_reduce.cuh, gemm/fp8_blockwise/fp8_blockwise_scaled_mm_sm120.cuh, elementwise/fused_qknorm_rope.cuh, inkling/inkling_all_reduce.cuh, deepseek_v4/hash_topk.cuh.

🤖 Generated with Claude Code

---

CI States

Latest PR Test (Base): 🚫 Run #40964193122
Latest PR Test (Extra): ❌ Run #40964193139
Latest PR Test (AMD ROCm 10): 🚫 Run #40964192940

[claude]

Claude Code Review

This repository is configured for manual code reviews. Comment @claude review for a one-time review, or @claude review always to subscribe this PR to a review on every future push.

_{Tip: disable this comment in your organization's Code Review settings.}