fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert fails to launch on C550: mcLaunchKernelExC returns mcErrorInvalidValue

[cheersluvs]

Summary

fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert cannot be launched on a MetaX C550. Every
call fails with:

[MCR][E]mc_runtime_api.cpp :382 : mcLaunchKernelExC: Returned mcErrorInvalidValue
torch.AcceleratorError: CUDA error: invalid argument

The kernel itself is correct — forwarding the launch to mcLaunchKernel makes it run and
produce bit-identical results. The failure is in the launch path, which is used by only two
operators in the whole library.

Note that mcLaunchKernelExC does not raise through the torch API; it logs and returns an
error, and the poisoned context surfaces later at an unrelated CUDA call. That makes this easy
to misattribute.

Environment

device	MetaX C550 (compute capability 8.0, 64-lane warps, 65536 B shared memory/block)
mcoplib	0.4.6 (Build_Maca_Version = 3.7.1.5, Vllm Op Version = 0.21.0)
torch	2.8.0+metax3.7.1.4
MACA	3.7.1.5

Reproduction

import torch, mcoplib._C

nt, nh, bs = 64, 64, 64
nb = (nt + bs - 1) // bs + 1
q  = torch.randn(nt, nh, 512, dtype=torch.bfloat16, device="cuda")
kv = torch.randn(nt, 512, dtype=torch.bfloat16, device="cuda")
kc = torch.zeros(nb, bs * 584, dtype=torch.uint8, device="cuda")
sm = torch.arange(nt, dtype=torch.int64, device="cuda")
ps = torch.arange(nt, dtype=torch.int64, device="cuda")
cs = torch.randn(4096, 64, dtype=torch.float32, device="cuda")

torch.ops._C.fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert(q, kv, kc, sm, ps, cs, 1e-6, bs)
torch.cuda.synchronize()

Fails identically at num_tokens = 1, 4, 64, 1024, 8192; num_heads = 64 and 128;
cache_block_size = 16 and 64; and for both the c10::Half and c10::BFloat16
instantiations. Host-side validation passes first (passing a bf16 k_cache correctly raises
k_cache must be uint8), so the arguments are accepted and the failure is at launch.

Rebuilding just this translation unit from the mcoplib_0.4.7 source with the MACA toolchain
(/opt/maca/tools/cu-bridge/bin/nvcc, -gencode=arch=compute_80,code=sm_80 -DUSE_MACA -DNV_ARCH_A100) reproduces the failure, so it is not specific to the shipped wheel.

Diagnosis

The launch configuration is valid. From the config struct passed to wcudaLaunchKernelExC:

grid = (520,1,1)   block = (256,1,1)   dynamicSmemBytes = 0   numAttrs = 0

The difference between this operator and the ones that work is the launch API:

launcher	launch call	works on C550
vllm::deepseek_v4_fused_ops::launchFusedDeepseekV4QNormRopeKVRopeInsert<T>	mcLaunchKernel	yes
vllm::deepseek_v4_fused_ops::launchFusedDeepseekV4QNormRopeKVRopeQuantInsert<T>	wcudaLaunchKernelExC	no

These are in the same source file and the same namespace. In mcoplib/_C.abi3.so,
wcudaLaunchKernelExC appears at only 4 call sites against ~5200 mcLaunchKernel calls,
covering exactly two operators:

• fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert (<Half> and <BFloat16>)
• dsv3_fused_a_gemm (invokeFusedAGemm<__maca_bfloat16, 7168, 2112, {8,16}>)

At compute capability 8.0 the extended call carries numAttrs = 0, so it requests nothing the
classic launch cannot express. Upstream vLLM uses the extended form only to attach a
Hopper-only PDL attribute, gated on sm_version >= 90.

Confirmation that the launch path is the cause

Interposing wcudaLaunchKernelExC with LD_PRELOAD and forwarding it to mcLaunchKernel
(same grid, block, shared memory and stream) makes the kernel run correctly:

[shim] grid=(520,1,1) block=(256,1,1) smem=0 attrs=0 -> mcLaunchKernel
q_changed=True cache_written=True

Output was compared against an independent Triton implementation on identical inputs:
FP8 cache bit-identical (0 of 74752 bytes differ), identical nonzero byte coverage, q
within one bf16 ULP. So the kernel and its device image are fine on this hardware.

Ruled out

• Shared memory — the config sets dynamicSmemBytes = 0.
• Block/grid dimensions — block is a constant (256,1,1); grid stays small at all shapes.
• The sm_version gate — cmp $0x4f (79) is the sm_80+ early-out and cmp $0x5a (90) only
  toggles numAttrs; both behave correctly at 80.
• FP8 support — torch.ops._C.static_scaled_fp8_quant runs correctly on this device, and
  Triton's tl.float8e4nv is bit-exact against torch here.
• Environment/ABI — torch is 2.8.0+metax3.7.1.4, matching mcoplib's torch2.8 build, and
  torch.ops._C.silu_and_mul works in the same process.

Suggested fix

Either implement cudaLaunchKernelExC in the MACA compatibility layer, or emit a classic
mcLaunchKernel when numAttrs == 0.

Upstream vLLM already has a non-PDL path for exactly this situation, currently guarded on
USE_ROCM:

#else
  // ROCm: use standard kernel launch syntax (no PDL/stream serialization)
  fusedDeepseekV4QNormRopeKVRopeQuantInsertKernel<scalar_t_in>
      <<<grid, kBlockSize, 0, stream>>>(...);
#endif

One caveat: in our from-source rebuild, plain <<<...>>> syntax still lowered to
wcudaLaunchKernelExC for this kernel while 21 other launches in the same translation unit
lowered to mcLaunchKernel. So source-level changes alone may not be sufficient, and what
selects the extended form per kernel may be worth checking on your side.

Impact

dsv3_fused_a_gemm uses the same launch path. It currently fails earlier, at its own
required CUDA ARCH >= SM_90 host check, so the launch problem is masked — but it would hit
the same issue on any device that clears that guard.

[cheersluvs]

Follow-up after reading the source in this repository across releases. It looks like this is already fixed on your side, in 0.4.9 — and if so, the remaining question is whether the underlying compatibility-layer defect was fixed with it.

What the source shows

op/vllm/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu, the quant launcher:

branch	launch call
mcoplib-0.4.6 (the wheel this report is against)	cudaLaunchKernelEx (line 408)
mcoplib-0.4.7	cudaLaunchKernelEx (lines 564, 573)
mcoplib-0.4.8	cudaLaunchKernelEx (lines 592, 601)
mcoplib-0.4.9	none — plain <<<grid, kBlockSize, 0, stream>>>
mcoplib-0.4.10	none — plain <<<grid, kBlockSize, 0, stream>>>

In 0.4.9 and 0.4.10 the classic launch is unconditional, not inside the #ifndef USE_ROCM / #else pair that carried it in the earlier releases. So the operator no longer goes through the extended launch API at all, which is exactly the change suggested in the report above.

Could you confirm this was deliberate, and which released wheel first carries it? That would let users on C550 move to a working version rather than patching around it.

The likely root cause, which may still be present

In 0.4.6 through 0.4.8 the launcher does this:

cudaLaunchAttribute attrs[1];
attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;   // Hopper-only
attrs[0].val.programmaticStreamSerializationAllowed = 1;
config.attrs    = attrs;
config.numAttrs = (sm_version >= 90) ? 1 : 0;                       // 0 on C550
cudaLaunchKernelEx(&config, ...);

Two properties of this stand out, and either would produce mcErrorInvalidValue:

1. config.attrs stays non-NULL when numAttrs == 0. It points at a fully populated, Hopper-only attribute. Per the CUDA contract numAttrs is the count of valid entries, so a conforming runtime must not read attrs[0] here. A wrapper that validates or iterates attrs before checking numAttrs would see a Hopper-only attribute id on an sm_80 device and reject the launch.
2. cudaLaunchConfig_t config; is never zero-initialized — only CUDA's six fields are assigned. If MACA's struct carries any field beyond those six, it holds uninitialized stack on every call.

Moving this one operator to <<<...>>> sidesteps both, but neither is specific to this operator: any future use of the extended launch API on MACA would hit the same thing. The only other ExC user in _C.abi3.so, invokeFusedAGemm<__maca_bfloat16, 7168, 2112, {8,16}> (_C::dsv3_fused_a_gemm), fails earlier at its own required CUDA ARCH >= SM_90 host check, so there is still no working ExC launch on this device to compare against.

Dumping the full cudaLaunchConfig_t at the wrapper entry — in particular whether attrs is non-NULL while numAttrs == 0, and the value of sizeof(cudaLaunchConfig_t) against CUDA's 56 — would settle which of the two it is.

Please also note the caveat already recorded in the report above: in our own single-translation-unit rebuild, a plain <<<...>>> launch still lowered to wcudaLaunchKernelExC for this kernel while 21 other launches in the same file lowered to mcLaunchKernel. We have not re-verified that observation and cannot currently reproduce the build, so treat it as unconfirmed — but if it holds, the 0.4.9 source change alone may not be sufficient on this hardware, and would be worth testing on a C550 before closing.

---

This comment has been revised twice. An earlier version claimed the 0.4.7 source used a plain <<<...>>> launch; that was drawn from our own notes and is incorrect — 0.4.7 uses cudaLaunchKernelEx, as tabulated above. The table is read directly from this repository.