From e75299a111fb879579b727bfb97c50a3dec7a935 Mon Sep 17 00:00:00 2001 From: Netanel Haber <58652339+netanel-haber@users.noreply.github.com> Date: Wed, 14 Jan 2026 11:26:48 +0200 Subject: [PATCH] Fix issues/16714: Revert comment out of `tl.debug_barrier()` in causal_conv1d_triton (#16899) Co-authored-by: Yi Zhang <1109276519@qq.com> --- .../srt/layers/attention/mamba/causal_conv1d_triton.py | 6 +++--- test/registered/layers/mamba/test_causal_conv1d.py | 6 ------ 2 files changed, 3 insertions(+), 9 deletions(-) diff --git a/python/sglang/srt/layers/attention/mamba/causal_conv1d_triton.py b/python/sglang/srt/layers/attention/mamba/causal_conv1d_triton.py index 247ae4326..c82f4d730 100644 --- a/python/sglang/srt/layers/attention/mamba/causal_conv1d_triton.py +++ b/python/sglang/srt/layers/attention/mamba/causal_conv1d_triton.py @@ -186,7 +186,7 @@ def _causal_conv1d_fwd_kernel( # continuous batching ) mask = (idx_tokens_conv < state_len)[:, None] & (idx_feats < dim)[None, :] - # tl.debug_barrier() # NOTE: use this due to bug in Triton compiler + tl.debug_barrier() # NOTE: use this due to bug in Triton compiler tl.store(conv_states_ptrs_target, new_conv_state, mask) else: @@ -221,7 +221,7 @@ def _causal_conv1d_fwd_kernel( # continuous batching ) # token-index # token-index # feature-index loaded_x = tl.load(x_ptrs, mask_x, 0.0) - # tl.debug_barrier() # need this due to the bug in tl.where not enforcing this when data is the result of another tl.load + tl.debug_barrier() # need this due to the bug in tl.where not enforcing this when data is the result of another tl.load new_conv_state = tl.where( mask, conv_state, loaded_x ) # BUG in 'tl.where' which requires a barrier before this @@ -728,7 +728,7 @@ def _causal_conv1d_update_kernel( & (idx_feats < dim)[None, :] ) # token-index # token-index # feature-index loaded_x = tl.load(x_ptrs, mask_x, 0.0) - # tl.debug_barrier() + tl.debug_barrier() new_conv_state = tl.where(mask, conv_state, loaded_x) diff --git a/test/registered/layers/mamba/test_causal_conv1d.py b/test/registered/layers/mamba/test_causal_conv1d.py index c50384f16..728c96a67 100644 --- a/test/registered/layers/mamba/test_causal_conv1d.py +++ b/test/registered/layers/mamba/test_causal_conv1d.py @@ -196,9 +196,6 @@ def test_causal_conv1d_update_with_batch_gather( if not torch.cuda.is_available(): pytest.skip("CUDA device not available") - # TODO: Fix Triton kernel bug - see issue #16714 - pytest.skip("Triton kernel bug - see issue #16714") - device = "cuda" rtol, atol = (3e-4, 1e-3) if itype == torch.float32 else (3e-3, 5e-3) if itype == torch.bfloat16: @@ -279,9 +276,6 @@ def test_causal_conv1d_varlen( if not torch.cuda.is_available(): pytest.skip("CUDA device not available") - # TODO: Fix Triton kernel bug - see issue #16714 - pytest.skip("Triton kernel bug - see issue #16714") - device = "cuda" torch.cuda.empty_cache() rtol, atol = (3e-4, 1e-3) if itype == torch.float32 else (3e-3, 5e-3)