[Kernel] Wait for the PDL dependency before loading router bias (Triton + radix)

python/sglang/kernels/jit/csrc/moe/route_radix.cuh

@@ -142,9 +142,10 @@ SGL_DEVICE void route_radix_block(const RouteRadixParams& params, typename Large
     // radix math below is fp32 either way — only the load width differs.
     AlignedVector<packed_t<TScore>, kVecSize / 2> scores_vec;
 
-    // prefetch bias (frozen weight) before the PDL wait
-    bias_vec.load(params.bias, tx);
+    // Bias may be produced by a preceding cast or fill kernel (the caller
+    // does not guarantee a frozen weight), so wait before loading either input.
     PDLWaitPrimary<kUsePDL>();
+    bias_vec.load(params.bias, tx);
     scores_vec.load(scores, tx);
 
 #pragma unroll

python/sglang/kernels/ops/moe/moe_fused_gate.py

@@ -127,17 +127,19 @@ def _router_triton_kernel(
     mask_m = offs_m < M
     mask_n = offs_n < N
 
-    # Prefetch a real bias before the PDL wait. Plain softmax routing has no
-    # bias, so keep the zero value in registers rather than materializing and
-    # clearing a device tensor for every routing call.
+    # PDL may start this grid before prior kernel stores are visible. Bias can
+    # be produced by a preceding cast or fill kernel, so wait before loading
+    # either bias or scores.
+    if USE_PDL:
+        tl.extra.cuda.gdc_wait()
+
+    # Plain softmax routing has no bias, so keep the zero value in registers
+    # rather than materializing and clearing a device tensor per call.
     if HAS_BIAS:
         bias = tl.load(bias_ptr + offs_n, mask=mask_n, other=0.0).to(tl.float32)
     else:
         bias = tl.zeros([BLOCK_N], dtype=tl.float32)
 
-    if USE_PDL:
-        tl.extra.cuda.gdc_wait()
-
     row_ptr = scores_ptr + offs_m[:, None] * stride_sm + offs_n[None, :] * stride_sn
     mask2d = mask_m[:, None] & mask_n[None, :]
     scores = tl.load(row_ptr, mask=mask2d, other=0.0).to(