diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.c b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.c
index 91d08dafcca1..1ac78643fed8 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.c
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.c
@@ -989,6 +989,274 @@ static int make_spec(int idx, rocke_attention_tiled_2d_spec_t* s)
         s->use_transposed_mask_once = true;
         break;
 
+    case 59:
+        s->head_size = 64;
+        s->block_size = 32;
+        s->num_query_heads = 64;
+        s->num_kv_heads = 8;
+        s->dtype = "bf16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 1;
+        s->num_warps = 4;
+        s->has_waves_per_eu = true;
+        s->waves_per_eu = 2;
+        s->use_v_double_buffer = true;
+        s->has_tile_size = true;
+        s->tile_size = 64;
+        s->block_m_per_warp = 32;
+        s->use_mfma_32x32 = true;
+        s->use_transposed_qk_32x32 = true;
+        s->use_transposed_scalar_state = true;
+        s->use_transposed_mask_once = true;
+        s->use_transposed_half_local_pv = true;
+        s->use_mfma32_skip_legacy_qreg = true;
+        s->use_transposed_mask_limit = true;
+        s->use_fast_paged_kv_desc = true;
+        break;
+
+    case 60:
+        s->head_size = 64;
+        s->block_size = 32;
+        s->num_query_heads = 32;
+        s->num_kv_heads = 32;
+        s->dtype = "fp16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 1;
+        s->num_warps = 4;
+        s->has_waves_per_eu = true;
+        s->waves_per_eu = 2;
+        s->use_v_double_buffer = true;
+        s->use_staggered_iter_wait = true;
+        s->has_tile_size = true;
+        s->tile_size = 128;
+        s->block_m_per_warp = 32;
+        s->use_mfma_32x32 = true;
+        s->use_transposed_qk_32x32 = true;
+        s->use_transposed_scalar_state = true;
+        s->use_transposed_mask_once = true;
+        s->use_transposed_half_local_pv = true;
+        s->use_mfma32_skip_legacy_qreg = true;
+        s->use_transposed_mask_limit = true;
+        break;
+
+    case 61:
+        s->head_size = 64;
+        s->block_size = 32;
+        s->num_query_heads = 8;
+        s->num_kv_heads = 8;
+        s->dtype = "bf16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 1;
+        s->num_warps = 2;
+        s->use_v_double_buffer = true;
+        break;
+
+    case 62:
+        s->head_size = 64;
+        s->block_size = 32;
+        s->num_query_heads = 8;
+        s->num_kv_heads = 8;
+        s->dtype = "bf16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 1;
+        s->num_warps = 2;
+        s->block_m_per_warp = 32;
+        s->use_mfma_32x32 = true;
+        s->use_v_double_buffer = true;
+        s->use_staggered_iter_wait = true;
+        break;
+
+    case 63:
+        s->head_size = 128;
+        s->block_size = 32;
+        s->num_query_heads = 8;
+        s->num_kv_heads = 8;
+        s->dtype = "fp16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 1;
+        s->num_warps = 1;
+        s->use_v_double_buffer = true;
+        s->use_early_v_schedule = true;
+        s->use_sched_barrier = true;
+        break;
+
+    case 64:
+        s->head_size = 256;
+        s->block_size = 32;
+        s->num_query_heads = 64;
+        s->num_kv_heads = 8;
+        s->dtype = "bf16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 2;
+        s->num_warps = 2;
+        s->has_waves_per_eu = true;
+        s->waves_per_eu = 2;
+        s->has_tile_size = true;
+        s->tile_size = 64;
+        s->block_m_per_warp = 32;
+        s->use_mfma_32x32 = true;
+        s->use_transposed_qk_32x32 = true;
+        s->use_kq_lds_pad = true;
+        s->kq_lds_pad_halves = 16;
+        s->use_transposed_scalar_state = true;
+        s->use_transposed_mask_once = true;
+        s->use_transposed_half_local_pv = true;
+        s->use_transposed_mask_limit = true;
+        s->use_k_single_buffer = true;
+        s->use_q_direct_reg = true;
+        s->use_softmax_mfma_interleave = true;
+        s->softmax_interleave_mode = 2;
+        s->use_mask_phase_split = true;
+        break;
+
+    case 65:
+        s->head_size = 256;
+        s->block_size = 32;
+        s->num_query_heads = 64;
+        s->num_kv_heads = 8;
+        s->dtype = "bf16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 2;
+        s->num_warps = 2;
+        s->has_waves_per_eu = true;
+        s->waves_per_eu = 2;
+        s->has_tile_size = true;
+        s->tile_size = 64;
+        s->block_m_per_warp = 32;
+        s->use_mfma_32x32 = true;
+        s->use_transposed_qk_32x32 = true;
+        s->use_kq_lds_pad = true;
+        s->kq_lds_pad_halves = 16;
+        s->use_transposed_scalar_state = true;
+        s->use_transposed_mask_once = true;
+        s->use_transposed_half_local_pv = true;
+        s->use_transposed_mask_limit = true;
+        s->use_k_single_buffer = true;
+        s->use_q_direct_reg = true;
+        s->use_softmax_mfma_interleave = true;
+        s->softmax_interleave_mode = 2;
+        s->use_mask_phase_split = true;
+        s->use_transposed_invariant_hoist = true;
+        break;
+
+    case 66:
+        s->head_size = 128;
+        s->block_size = 32;
+        s->num_query_heads = 64;
+        s->num_kv_heads = 8;
+        s->dtype = "bf16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 2;
+        s->num_warps = 4;
+        s->has_tile_size = true;
+        s->tile_size = 64;
+        s->block_m_per_warp = 32;
+        s->use_mfma_32x32 = true;
+        s->use_transposed_qk_32x32 = true;
+        s->use_transposed_scalar_state = true;
+        s->use_transposed_mask_once = true;
+        s->use_transposed_half_local_pv = true;
+        s->use_transposed_mask_limit = true;
+        s->use_q_direct_reg = true;
+        break;
+
+    case 67:
+        s->head_size = 128;
+        s->block_size = 32;
+        s->num_query_heads = 64;
+        s->num_kv_heads = 8;
+        s->dtype = "fp16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 1;
+        s->num_warps = 2;
+        s->has_tile_size = true;
+        s->tile_size = 64;
+        s->block_m_per_warp = 32;
+        s->use_mfma_32x32 = true;
+        s->use_transposed_qk_32x32 = true;
+        s->use_transposed_scalar_state = true;
+        s->use_transposed_mask_once = true;
+        s->use_transposed_half_local_pv = true;
+        s->use_transposed_mask_limit = true;
+        s->use_k_single_buffer = true;
+        s->use_kq_lds_pad = true;
+        s->use_mask_phase_split = true;
+        s->use_softmax_mfma_interleave = true;
+        s->softmax_interleave_mode = 2;
+        s->softmax_interleave_groups = 2;
+        break;
+
+    case 68:
+        s->head_size = 256;
+        s->block_size = 16;
+        s->num_query_heads = 32;
+        s->num_kv_heads = 32;
+        s->dtype = "bf16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 1;
+        s->num_warps = 2;
+        s->has_waves_per_eu = true;
+        s->waves_per_eu = 2;
+        s->has_tile_size = true;
+        s->tile_size = 64;
+        s->block_m_per_warp = 32;
+        s->use_mfma_32x32 = true;
+        s->use_transposed_qk_32x32 = true;
+        s->use_kq_lds_pad = true;
+        s->kq_lds_pad_halves = 16;
+        s->use_transposed_scalar_state = true;
+        s->use_transposed_mask_once = true;
+        s->use_transposed_half_local_pv = true;
+        s->use_transposed_mask_limit = true;
+        s->use_k_single_buffer = true;
+        s->use_q_direct_reg = true;
+        s->use_softmax_mfma_interleave = true;
+        s->softmax_interleave_mode = 2;
+        s->use_mask_phase_split = true;
+        break;
+
+    case 69:
+        s->head_size = 64;
+        s->block_size = 32;
+        s->num_query_heads = 64;
+        s->num_kv_heads = 8;
+        s->dtype = "bf16";
+        s->use_sinks = false;
+        s->sliding_window = 0;
+        s->has_softcap = false;
+        s->num_seqs = 2;
+        s->num_warps = 4;
+        s->has_tile_size = true;
+        s->tile_size = 64;
+        s->block_m_per_warp = 32;
+        s->use_mfma_32x32 = true;
+        s->use_transposed_qk_32x32 = true;
+        s->use_transposed_scalar_state = true;
+        s->use_transposed_half_local_pv = true;
+        s->kv_storage_dtype = "fp8e4m3";
+        s->use_fp8_mfma_qk = true;
+        break;
+
     default:
         return -1;
     }
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.py b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.py
index ad0a91caf51a..39fc2e3ae2c1 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx950_attention_tiled_2d_emit.py
@@ -964,6 +964,262 @@
         use_transposed_invariant_hoist=True,
         use_transposed_mask_once=True,
     ),
+    # V double-buffer (V[i+1] prefetched into the alternate V_lds slot). idx59 is
+    # the bf16 D64 GQA-8 single-batch prefill spec _tiled_spec_from_problem
+    # selects; idx60 the fp16 MHA one (tile 128) plus the staggered iter-start
+    # wait. idx61-63 cover the other QK/PV paths: narrow 16x16, non-transposed
+    # 32x32 with the staggered wait, and early-V + sched_barrier at D128.
+    59: dict(
+        head_size=64,
+        block_size=32,
+        num_query_heads=64,
+        num_kv_heads=8,
+        dtype="bf16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=1,
+        num_warps=4,
+        waves_per_eu=2,
+        use_v_double_buffer=True,
+        tile_size=64,
+        block_m_per_warp=32,
+        use_mfma_32x32=True,
+        use_transposed_qk_32x32=True,
+        use_transposed_scalar_state=True,
+        use_transposed_mask_once=True,
+        use_transposed_half_local_pv=True,
+        use_mfma32_skip_legacy_qreg=True,
+        use_transposed_mask_limit=True,
+        use_fast_paged_kv_desc=True,
+    ),
+    60: dict(
+        head_size=64,
+        block_size=32,
+        num_query_heads=32,
+        num_kv_heads=32,
+        dtype="fp16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=1,
+        num_warps=4,
+        waves_per_eu=2,
+        use_v_double_buffer=True,
+        use_staggered_iter_wait=True,
+        tile_size=128,
+        block_m_per_warp=32,
+        use_mfma_32x32=True,
+        use_transposed_qk_32x32=True,
+        use_transposed_scalar_state=True,
+        use_transposed_mask_once=True,
+        use_transposed_half_local_pv=True,
+        use_mfma32_skip_legacy_qreg=True,
+        use_transposed_mask_limit=True,
+    ),
+    61: dict(
+        head_size=64,
+        block_size=32,
+        num_query_heads=8,
+        num_kv_heads=8,
+        dtype="bf16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=1,
+        num_warps=2,
+        use_v_double_buffer=True,
+    ),
+    62: dict(
+        head_size=64,
+        block_size=32,
+        num_query_heads=8,
+        num_kv_heads=8,
+        dtype="bf16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=1,
+        num_warps=2,
+        block_m_per_warp=32,
+        use_mfma_32x32=True,
+        use_v_double_buffer=True,
+        use_staggered_iter_wait=True,
+    ),
+    63: dict(
+        head_size=128,
+        block_size=32,
+        num_query_heads=8,
+        num_kv_heads=8,
+        dtype="fp16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=1,
+        num_warps=1,
+        use_v_double_buffer=True,
+        use_early_v_schedule=True,
+        use_sched_barrier=True,
+    ),
+    # D256 fast path (_d256_gfx950_spec_overrides): idx64 is the bf16 GQA-8 spec
+    # _tiled_spec_from_problem selects (direct-to-register Q, K-slab LDS pad,
+    # mask phase split, interleave mode 2), idx65 the same with the invariant
+    # hoist, idx68 the MHA block-16 shape. idx66-67 take the levers apart at
+    # D128: direct-register Q alone; K-slab pad + phase split + mode 2 (2 groups).
+    64: dict(
+        head_size=256,
+        block_size=32,
+        num_query_heads=64,
+        num_kv_heads=8,
+        dtype="bf16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=2,
+        num_warps=2,
+        waves_per_eu=2,
+        tile_size=64,
+        block_m_per_warp=32,
+        use_mfma_32x32=True,
+        use_transposed_qk_32x32=True,
+        use_kq_lds_pad=True,
+        kq_lds_pad_halves=16,
+        use_transposed_scalar_state=True,
+        use_transposed_mask_once=True,
+        use_transposed_half_local_pv=True,
+        use_transposed_mask_limit=True,
+        use_k_single_buffer=True,
+        use_q_direct_reg=True,
+        use_softmax_mfma_interleave=True,
+        softmax_interleave_mode=2,
+        use_mask_phase_split=True,
+    ),
+    65: dict(
+        head_size=256,
+        block_size=32,
+        num_query_heads=64,
+        num_kv_heads=8,
+        dtype="bf16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=2,
+        num_warps=2,
+        waves_per_eu=2,
+        tile_size=64,
+        block_m_per_warp=32,
+        use_mfma_32x32=True,
+        use_transposed_qk_32x32=True,
+        use_kq_lds_pad=True,
+        kq_lds_pad_halves=16,
+        use_transposed_scalar_state=True,
+        use_transposed_mask_once=True,
+        use_transposed_half_local_pv=True,
+        use_transposed_mask_limit=True,
+        use_k_single_buffer=True,
+        use_q_direct_reg=True,
+        use_softmax_mfma_interleave=True,
+        softmax_interleave_mode=2,
+        use_mask_phase_split=True,
+        use_transposed_invariant_hoist=True,
+    ),
+    66: dict(
+        head_size=128,
+        block_size=32,
+        num_query_heads=64,
+        num_kv_heads=8,
+        dtype="bf16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=2,
+        num_warps=4,
+        tile_size=64,
+        block_m_per_warp=32,
+        use_mfma_32x32=True,
+        use_transposed_qk_32x32=True,
+        use_transposed_scalar_state=True,
+        use_transposed_mask_once=True,
+        use_transposed_half_local_pv=True,
+        use_transposed_mask_limit=True,
+        use_q_direct_reg=True,
+    ),
+    67: dict(
+        head_size=128,
+        block_size=32,
+        num_query_heads=64,
+        num_kv_heads=8,
+        dtype="fp16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=1,
+        num_warps=2,
+        tile_size=64,
+        block_m_per_warp=32,
+        use_mfma_32x32=True,
+        use_transposed_qk_32x32=True,
+        use_transposed_scalar_state=True,
+        use_transposed_mask_once=True,
+        use_transposed_half_local_pv=True,
+        use_transposed_mask_limit=True,
+        use_k_single_buffer=True,
+        use_kq_lds_pad=True,
+        use_mask_phase_split=True,
+        use_softmax_mfma_interleave=True,
+        softmax_interleave_mode=2,
+        softmax_interleave_groups=2,
+    ),
+    68: dict(
+        head_size=256,
+        block_size=16,
+        num_query_heads=32,
+        num_kv_heads=32,
+        dtype="bf16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=1,
+        num_warps=2,
+        waves_per_eu=2,
+        tile_size=64,
+        block_m_per_warp=32,
+        use_mfma_32x32=True,
+        use_transposed_qk_32x32=True,
+        use_kq_lds_pad=True,
+        kq_lds_pad_halves=16,
+        use_transposed_scalar_state=True,
+        use_transposed_mask_once=True,
+        use_transposed_half_local_pv=True,
+        use_transposed_mask_limit=True,
+        use_k_single_buffer=True,
+        use_q_direct_reg=True,
+        use_softmax_mfma_interleave=True,
+        softmax_interleave_mode=2,
+        use_mask_phase_split=True,
+    ),
+    # fp8-in-LDS K (native-fp8 MFMA QK loader) on the transposed 32x32 path: the
+    # transposed QK reads K through the shared dequantizing K reader.
+    69: dict(
+        head_size=64,
+        block_size=32,
+        num_query_heads=64,
+        num_kv_heads=8,
+        dtype="bf16",
+        use_sinks=False,
+        sliding_window=0,
+        has_softcap=False,
+        num_seqs=2,
+        num_warps=4,
+        tile_size=64,
+        block_m_per_warp=32,
+        use_mfma_32x32=True,
+        use_transposed_qk_32x32=True,
+        use_transposed_scalar_state=True,
+        use_transposed_half_local_pv=True,
+        kv_storage_dtype="fp8e4m3",
+        use_fp8_mfma_qk=True,
+    ),
 }
 
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_binding_spec_fields.py b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_binding_spec_fields.py
index 597ba931eee6..17b0fc77848c 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_binding_spec_fields.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_binding_spec_fields.py
@@ -30,9 +30,6 @@
     "use_k_hbm_direct",
     "use_causal_mask_phase_split",
     "use_additive_bias",
-    "use_kq_lds_pad",
-    "kq_lds_pad_halves",
-    "use_mask_phase_split",
     "causal_top_left",
     "right_bound",
 }
@@ -85,7 +82,7 @@ def test_gfx950_2d_default_spec_round_trips(engine):
 
 @pytest.mark.parametrize(
     "field, value",
-    [("use_additive_bias", True), ("kq_lds_pad_halves", 16), ("right_bound", -1)],
+    [("use_additive_bias", True), ("causal_top_left", True), ("right_bound", -1)],
 )
 def test_2d_binding_rejects_unported_fields(engine, field, value):
     d = asdict(Gfx950Spec2D(**_REQUIRED_2D, dtype="bf16"))
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py
index 25c06ee34c11..4ddaebd592db 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py
@@ -22,7 +22,6 @@
     _canonical_arch_for_selection,
     _enable_gfx942_flash_k_sliced_ring,
     _enable_gfx942_flash_mask_limit,
-    _enable_v_double_buffer,
     _gfx942_4warp_fast,
     _mask_label,
     _reject_mask_on_arch,
@@ -263,8 +262,6 @@ def _tiled_2d_not_ported(p: UnifiedAttentionProblem, arch: str):
         return "tiled 2D additive bias is not ported to C++ yet"
     if not p.default_mask:
         return f"tiled 2D {_mask_label(p)} masks are not ported to C++ yet"
-    if _enable_v_double_buffer(p, arch):
-        return "tiled 2D V double-buffer is not ported to C++ yet"
     return None
 
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_selector_cpp_parity.py b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_selector_cpp_parity.py
index 866d361629ff..f71a1c0a5819 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_selector_cpp_parity.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_selector_cpp_parity.py
@@ -223,9 +223,6 @@ def test_selectors_follow_env(engine, monkeypatch, env):
     "use_additive_bias": "additive bias is not in the C++ tiled-2D spec",
     "causal_top_left": "the top-left mask band is not in the C++ tiled-2D spec",
     "right_bound": "the right-bound mask band is not in the C++ tiled-2D spec",
-    "use_mask_phase_split": "the D256 fast path's mask phase split is not in the C++ tiled-2D spec",
-    "use_kq_lds_pad": "the D256 fast path's K/Q LDS padding is not in the C++ tiled-2D spec",
-    "kq_lds_pad_halves": "the D256 fast path's K/Q LDS padding is not in the C++ tiled-2D spec",
 }
 
 # gfx942 knobs carried by the C spec struct both arches share; the gfx950
@@ -260,18 +257,7 @@ class _Gap:
 
 # Spec drift the C++ side still has to port. Remove a row (and its xfail
 # below) when the port lands.
-_KNOWN_SPEC_GAPS = {
-    "d256_gfx950_overrides": _Gap(
-        "the gfx950 D256 fast-path spec overrides are not applied by the C++ spec builder",
-        frozenset(au._d256_gfx950_spec_overrides()),
-        lambda p, arch: au._d256_gfx950_fast(p, arch),
-    ),
-    "softmax_mfma_interleave": _Gap(
-        "the C++ spec builder never sets use_softmax_mfma_interleave",
-        frozenset({"use_softmax_mfma_interleave"}),
-        lambda p, arch: au._enable_softmax_mfma_interleave(p, arch),
-    ),
-}
+_KNOWN_SPEC_GAPS: dict[str, _Gap] = {}
 
 
 def _norm(v):
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp
index f1edee4ea98b..d9761c911e08 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp
@@ -848,9 +848,6 @@ const unported_field k_t2d_unported[] = {
     {"use_k_hbm_direct", false, 0},
     {"use_causal_mask_phase_split", false, 0},
     {"use_additive_bias", false, 0},
-    {"use_kq_lds_pad", false, 0},
-    {"kq_lds_pad_halves", false, 8},
-    {"use_mask_phase_split", false, 0},
     {"causal_top_left", false, 0},
     {"right_bound", false, 0},
 };
@@ -937,6 +934,9 @@ void t2d_read(const py::dict& d, Store& st, rocke_attention_tiled_2d_spec_t& s)
     s.use_q_direct_global = a_bool(d, "use_q_direct_global", s.use_q_direct_global);
     s.use_global_load_lds_k = a_bool(d, "use_global_load_lds_k", s.use_global_load_lds_k);
     s.use_q_major_grid = a_bool(d, "use_q_major_grid", s.use_q_major_grid);
+    s.use_kq_lds_pad = a_bool(d, "use_kq_lds_pad", s.use_kq_lds_pad);
+    s.kq_lds_pad_halves = a_int(d, "kq_lds_pad_halves", s.kq_lds_pad_halves);
+    s.use_mask_phase_split = a_bool(d, "use_mask_phase_split", s.use_mask_phase_split);
     std::string v;
     if(a_str(d, "dtype", v))
         s.dtype = st.keep(v);
@@ -1564,6 +1564,9 @@ static py::dict spec_to_dict(const rocke_attention_tiled_2d_spec_t& s)
         = s.kv_cache_policy ? py::cast(std::string(s.kv_cache_policy)) : py::none();
     d["use_global_load_lds_k"] = s.use_global_load_lds_k;
     d["use_q_major_grid"] = s.use_q_major_grid;
+    d["use_kq_lds_pad"] = s.use_kq_lds_pad;
+    d["kq_lds_pad_halves"] = s.kq_lds_pad_halves;
+    d["use_mask_phase_split"] = s.use_mask_phase_split;
     return d;
 }
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_helper_rocke.instances.gfx942.attention_tiled_2d.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_helper_rocke.instances.gfx942.attention_tiled_2d.h
index de8da22bdbe9..37bb1ab1dbcc 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_helper_rocke.instances.gfx942.attention_tiled_2d.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_helper_rocke.instances.gfx942.attention_tiled_2d.h
@@ -107,34 +107,25 @@ typedef struct rocke_attention_tiled_2d_spec
     bool use_register_pv; /* False */
 
     /* Deep K+V prefetch / V-double-buffer schedule controls (Python mirror).
-     * The gfx950 C emitter currently builds only the depth-2 single/early-V
-     * schedule; these fields are carried so the spec is byte-identical with the
-     * Python dataclass and the provider/dispatcher can pass them through. A
-     * builder that sees use_v_double_buffer / use_staggered_iter_wait / kv_ring
-     * _depth!=2 / use_q_reread it does not yet emit must reject in validate(),
-     * NOT silently ignore (mirrors the Python __post_init__). */
+     * The gfx950 C emitter builds the V double-buffer and the staggered
+     * iter-start wait; it rejects kv_ring_depth!=2 and use_q_reread rather than
+     * silently emit the depth-2 schedule. */
     bool use_v_double_buffer; /* False */
     int kv_ring_depth; /* 2 */
     bool use_staggered_iter_wait; /* False */
     bool use_q_reread; /* False */
-    /* VGPR-frugal BLOCK_M=128 body -- gather the per-lane Q32 MFMA operand
-     * straight from global into VGPRs (no Q_lds), freeing 32 KB LDS so
-     * BLOCK_M=128/T=64 fits the 32 KB / 2-WG/CU budget with K+V single-buffer.
-     * Carried for spec byte-identity with the Python dataclass; the gfx950 C
-     * twin does not yet emit this body and rejects it in validate() (mirrors
-     * Python __post_init__ -- reject, do not silently ignore). Default False. */
+    /* gfx950: gather the per-lane Q32 MFMA operand straight from global into
+     * VGPRs (no Q_lds), freeing the Q LDS slab. Transposed-32x32 path only. */
     bool use_q_direct_reg; /* False */
     /* Lever-3 sched_barrier fence after the QK MFMA cluster (CK-Tile-derived).
      * Emitted by the gfx950 C twin (independent of the V schedule). */
     bool use_sched_barrier; /* False */
     int sched_barrier_mask; /* 0 */
 
-    /* softmax-window MFMA interleave (Python gfx950 lever). The default
-     * (off) is byte-identical; when on it emits a single iglp_opt at the loop
-     * top (mode 0/1) or sched_group_barrier groups (mode 2) to schedule the
-     * QK/PV MFMAs into the softmax VALU window. Carried for spec parity with the
-     * Python dataclass; the gfx950 C twin rejects it (the body it rides on,
-     * use_q_direct_reg, is itself unported). */
+    /* gfx950 softmax-window MFMA interleave. When on it emits a single
+     * iglp_opt at the loop top (mode 0/1) or softmax_interleave_groups
+     * sched_group_barrier pairs after the QK cluster (mode 2) to schedule the
+     * QK/PV MFMAs into the softmax VALU window. */
     bool use_softmax_mfma_interleave; /* False */
     int softmax_interleave_mode; /* 1 */
     int softmax_interleave_groups; /* 4 */
@@ -178,6 +169,14 @@ typedef struct rocke_attention_tiled_2d_spec
     const char* kv_cache_policy; /* "stream" */
     bool use_global_load_lds_k; /* False */
     bool use_q_major_grid; /* False */
+    /* gfx950-only (gfx942 never reads them). K_lds pad of kq_lds_pad_halves
+     * halves per async-DMA slab, breaking the row-aliased bank conflict on the
+     * QK K read; wired for the K single-buffer path only. */
+    bool use_kq_lds_pad; /* False */
+    int kq_lds_pad_halves; /* 8 */
+    /* gfx950-only: split the KV loop into a mask-free full-tile phase and a
+     * masked boundary phase (transposed-32x32 mask-limit path, no SW). */
+    bool use_mask_phase_split; /* False */
 } rocke_attention_tiled_2d_spec_t;
 
 /* Materialise every defaulted field (the dataclass defaults). The required
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx950_attention_tiled_2d_internal.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx950_attention_tiled_2d_internal.h
index 3e5902944206..32d620a01f09 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx950_attention_tiled_2d_internal.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx950_attention_tiled_2d_internal.h
@@ -130,6 +130,10 @@ typedef struct rocke_gfx950_attn2d_build_ctx
     /* K single-buffer (1 K slot) so T=64 fits the 2-WG/CU LDS budget at
      * HD=128. The next-K prefetch is deferred to after the PV-wait barrier. */
     bool K_SINGLE_BUFFER;
+    /* V double-buffer: V[i+1] is prefetched into the alternate V_lds slot
+     * during iter i. STAGGER_ITER_WAIT (V_DOUBLE_BUF only) issues K before V
+     * and partial-waits the iter-start drain so V[i] stays in flight. */
+    bool V_DOUBLE_BUF, STAGGER_ITER_WAIT;
     /* Lever-3 (CK-Tile-derived) sched_barrier fence after the QK MFMA cluster. */
     bool USE_SCHED_BARRIER;
     int SCHED_BARRIER_MASK;
@@ -137,6 +141,12 @@ typedef struct rocke_gfx950_attn2d_build_ctx
      * modes 0/1 emit iglp_opt(mode); mutually exclusive with USE_SCHED_BARRIER. */
     bool USE_SOFTMAX_INTERLEAVE;
     int SOFTMAX_INTERLEAVE_MODE;
+    int SOFTMAX_INTERLEAVE_GROUPS; /* mode 2: (MFMA, VALU) group-barrier pairs */
+    /* Q gathered straight from global into VGPRs: no Q_lds, no Q aliasing. */
+    bool Q_DIRECT_REG;
+    /* K_lds slab pad in halves (0 = off) and the rows one wave's async-DMA
+     * slab fills; K_lds is then [K_BUFS, T/KQ_SLAB_ROWS, KQ_SLAB_ROWS*HD+pad]. */
+    int KQ_LDS_PAD, KQ_SLAB_ROWS;
     /* fp8 K/V cache predicates */
     bool KV_FP8, FP8_MFMA_QK, FP8_MFMA_PV;
     bool REGISTER_PV, USE_MFMA_32X32, TRANSPOSED_QK_32X32;
@@ -254,6 +264,7 @@ typedef struct rocke_gfx950_attn2d_build_ctx
     rocke_value_t* qk_scale_v; /* derived from scale_p            */
     rocke_value_t* pv_fp8_scale_v; /* fdiv(v_scale, 240) when FP8_MFMA_PV (line 1149) */
     rocke_value_t* sw_const_v; /* b.const_i32(SLIDING_WINDOW)     */
+    rocke_value_t* z8_v; /* b.zero_vec(dtype, 8), the masked-Q fill */
 
     /* ---- paged-KV byte descriptor (full transform DAG, lines 1163-1630) ---- */
     rocke_tensor_descriptor_t* q_desc; /* output/query [token,head,dim]   */
@@ -264,6 +275,7 @@ typedef struct rocke_gfx950_attn2d_build_ctx
     rocke_value_t* lane_half_base_v; /* tid * KV_HALVES_PER_LANE        */
     rocke_value_t* zero_soff_v; /* const_i32(0) soffset            */
     rocke_value_t* wave_lds_off_i64_v; /* wave_lds_offset_i64             */
+    rocke_value_t* wave_lds_off_K_i64_v; /* slab-padded K wave offset (KQ_LDS_PAD) */
     rocke_value_t* v_wave_lds_off_i64_v; /* v_wave_lds_offset_i64           */
     rocke_value_t* K_lds_addr_v; /* ptrtoint K_lds                  */
     rocke_value_t* V_lds_addr_v; /* ptrtoint V_lds                  */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp
index f91de89db70f..4155f007dc2f 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp
@@ -186,11 +186,6 @@ static bool reject_tiled_2d_not_ported(const rocke_unified_attention_problem_t*
         gate_reason(reason, reason_cap, "tiled 2D %s masks are not ported to C++ yet", label);
         return true;
     }
-    if(rocke_unified_attn_enable_v_double_buffer(p, arch))
-    {
-        gate_reason(reason, reason_cap, "tiled 2D V double-buffer is not ported to C++ yet");
-        return true;
-    }
     return false;
 }
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_selectors.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_selectors.cpp
index df30d42b8f2a..19eace5cf93e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_selectors.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_selectors.cpp
@@ -1501,8 +1501,7 @@ static rocke_attention_tiled_2d_spec_t
     return spec_base_2d_generic(p, arch);
 }
 
-/* Python: spec_builder._spec_gfx950_generic(problem, arch), without the D256
- * fast-route overrides and the softmax-interleave fields. */
+/* Python: spec_builder._spec_gfx950_generic(problem, arch). */
 static rocke_attention_tiled_2d_spec_t
     spec_gfx950_generic(const rocke_unified_attention_problem_t* p, const char* arch)
 {
@@ -1518,6 +1517,39 @@ static rocke_attention_tiled_2d_spec_t
         int t_eff_ks = s.has_tile_size ? s.tile_size : s.block_size;
         s.use_k_single_buffer = enable_d128_small_tile(p, arch) && (block_m_ks <= t_eff_ks);
     }
+    if(enable_softmax_mfma_interleave(p, arch))
+    {
+        s.use_softmax_mfma_interleave = true;
+        s.softmax_interleave_mode = 1;
+    }
+
+    /* D256 bf16 prefill fast route: pin the 32x32 transposed stack, the
+     * direct-to-register Q gather, the mask phase split, the K-slab LDS pad and
+     * the mode-2 softmax<->MFMA interleave on top of the gated geometry
+     * (Python _d256_gfx950_spec_overrides). */
+    if(d256_gfx950_fast(p, arch))
+    {
+        s.use_mfma_32x32 = true;
+        s.use_transposed_qk_32x32 = true;
+        s.use_q_direct_reg = true;
+        s.use_transposed_half_local_pv = true;
+        s.use_transposed_scalar_state = true;
+        s.use_transposed_mask_once = true;
+        s.use_transposed_mask_limit = true;
+        s.use_mask_phase_split = true;
+        s.use_register_pv = false;
+        s.use_k_single_buffer = true;
+        s.use_v_double_buffer = false;
+        s.use_early_v_schedule = false;
+        s.use_sched_barrier = false;
+        s.use_softmax_mfma_interleave = true;
+        s.softmax_interleave_mode = 2;
+        s.softmax_interleave_groups = 4;
+        s.use_fast_paged_kv_desc = false;
+        s.use_mfma32_skip_legacy_qreg = false;
+        s.use_kq_lds_pad = true;
+        s.kq_lds_pad_halves = 16;
+    }
     return s;
 }
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d.cpp
index 074fe6fd4d94..a7a1087db21c 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_tiled_2d.cpp
@@ -128,6 +128,9 @@ rocke_attention_tiled_2d_spec_t rocke_attention_tiled_2d_spec_default(void)
     s.kv_cache_policy = "stream";
     s.use_global_load_lds_k = false;
     s.use_q_major_grid = false;
+    s.use_kq_lds_pad = false;
+    s.kq_lds_pad_halves = 8;
+    s.use_mask_phase_split = false;
     return s;
 }
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_body_pv_epilogue.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_body_pv_epilogue.cpp
index 0a8b47281437..a6d54dfb0feb 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_body_pv_epilogue.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_body_pv_epilogue.cpp
@@ -38,6 +38,7 @@
 #include "rocke/instance_gfx950_attention_tiled_2d_internal.h"
 
 #include <assert.h>
+#include <stdio.h>
 
 /* ============================================================ *
  *  Local dtype-dispatch wrappers mirroring the Python module aliases
@@ -277,6 +278,14 @@ void rocke_gfx950_attn2d_emit_pv_bucket(rocke_gfx950_attn2d_build_ctx_t* ctx,
         rocke_b_s_waitcnt(b, 0, 0, -1);
         rocke_b_sync(b);
     }
+    else if(ctx->V_DOUBLE_BUF)
+    {
+        /* V[i] (cur_buf) is the oldest pending stream in both V double-buffer
+         * schedules (or already drained at iter start); leave this iter's
+         * K[i+1] + V[i+1] prefetches in flight. */
+        rocke_b_s_waitcnt(b, 2 * kv_calls_per_tile, 2 * kv_calls_per_tile, -1);
+        rocke_b_sync(b);
+    }
     else if(ctx->K_SINGLE_BUFFER)
     {
         /* single K slot. Only V[i] is in flight here (next-K was deferred).
@@ -307,7 +316,7 @@ void rocke_gfx950_attn2d_emit_pv_bucket(rocke_gfx950_attn2d_build_ctx_t* ctx,
             /* Transposed PV: O^T = V^T @ P^T via register P^T (Python 3221-3341).
              * v_buf and use_hi are emitted ONCE here (Python 3231-3232), before
              * the per-N acc-scale + apply loop, and reused by every apply call. */
-            ctx->pv_v_buf_v = rocke_b_const_i32(b, 0);
+            ctx->pv_v_buf_v = ctx->V_DOUBLE_BUF ? ctx->cur_buf : rocke_b_const_i32(b, 0);
             ctx->pv_use_hi_v = rocke_b_cmp_eq(b, ctx->lane_half32_v, rocke_b_const_i32(b, 1));
             for(n = 0; n < ctx->ACC_N_TILES; ++n)
             {
@@ -349,7 +358,7 @@ void rocke_gfx950_attn2d_emit_pv_bucket(rocke_gfx950_attn2d_build_ctx_t* ctx,
         else
         {
             /* Transitional 32x32 PV from the P_lds bridge (Python 3342-3412). */
-            rocke_value_t* v_buf = rocke_b_const_i32(b, 0);
+            rocke_value_t* v_buf = ctx->V_DOUBLE_BUF ? ctx->cur_buf : rocke_b_const_i32(b, 0);
             for(n = 0; n < ctx->ACC_N_TILES; ++n)
             {
                 rocke_value_t* scaled[ROCKE_GFX950_ATTN2D_MAX_REGS_PER_LANE];
@@ -459,7 +468,7 @@ void rocke_gfx950_attn2d_emit_pv_bucket(rocke_gfx950_attn2d_build_ctx_t* ctx,
                 rocke_value_t* ncb_16 = rocke_b_const_i32(b, 16);
                 n_col_base = rocke_b_add(b, rocke_b_mul(b, ncb_n, ncb_16), ctx->pv_tr_reader->col);
             }
-            v_buf = rocke_b_const_i32(b, 0);
+            v_buf = ctx->V_DOUBLE_BUF ? ctx->cur_buf : rocke_b_const_i32(b, 0);
 
             for(k = 0; k < ctx->PV_K_ITERS; ++k)
             {
@@ -675,14 +684,56 @@ void rocke_gfx950_attn2d_emit_pv_bucket(rocke_gfx950_attn2d_build_ctx_t* ctx,
  *  emitter, then leave. Returns the loop handle; the epilogue reads the
  *  rewritten carry from kvloop.op->results.
  * ============================================================ */
+/* One scf_for_iter phase over [lo, hi): enter the body, unpack the carry the
+ * way the Python body does (lines 2495-2508), run the full body emitter with the
+ * given skip_mask, then leave. */
+static rocke_for_t drive_kv_phase(rocke_gfx950_attn2d_build_ctx_t* ctx,
+                                  rocke_value_t* lo,
+                                  rocke_value_t* hi,
+                                  const rocke_iter_arg_t* iter_args,
+                                  const char* iv_name,
+                                  bool skip_mask)
+{
+    rocke_ir_builder_t* b = ctx->b;
+    int ml = ctx->ml_count;
+    int num_accs = ctx->ACC_N_TILES * ctx->ACC_M_ATOMS;
+    int i;
+    rocke_for_t loop = rocke_b_scf_for_iter(
+        b, lo, hi, ctx->kv_step, iter_args, ctx->iter_args_count, iv_name, false, true);
+
+    rocke_b_region_enter(b, loop.body);
+    ctx->kv_tile_iv = loop.iv;
+    for(i = 0; i < ctx->SOFTMAX_STATE_SLOTS; ++i)
+    {
+        ctx->m_cur[i] = loop.iter_vars[2 * i];
+        ctx->l_cur[i] = loop.iter_vars[2 * i + 1];
+    }
+    for(i = 0; i < num_accs; ++i)
+        ctx->acc_cur[i] = loop.iter_vars[ml + i];
+    ctx->cur_buf = loop.iter_vars[ml + num_accs];
+    ctx->skip_mask = skip_mask;
+
+    rocke_gfx950_attn2d_emit_kv_body(ctx);
+
+    rocke_b_region_leave(b);
+    return loop;
+}
+
+/* ============================================================ *
+ *  rocke_gfx950_attn2d_drive_kv_loop   (Python lines 4442-4486)
+ *
+ *  scf_for_iter over [tile_start, tile_end) step kv_step with the named
+ *  iter_args carry. With the mask phase split (no-SW transposed mask-limit
+ *  path) the range splits at the first tile that needs causal / prefix
+ *  masking: a mask-free full-tile loop, then the masked boundary loop, with the
+ *  carry threaded through. Returns the last loop; the epilogue reads the
+ *  rewritten carry from its results.
+ * ============================================================ */
 rocke_for_t rocke_gfx950_attn2d_drive_kv_loop(rocke_gfx950_attn2d_build_ctx_t* ctx)
 {
     rocke_ir_builder_t* b = ctx->b;
     rocke_iter_arg_t iter_args[ROCKE_GFX950_ATTN2D_MAX_ITER_ARGS];
-    rocke_for_t kvloop;
     int i;
-    int ml = ctx->ml_count;
-    int num_accs = ctx->ACC_N_TILES * ctx->ACC_M_ATOMS;
 
     for(i = 0; i < ctx->iter_args_count; ++i)
     {
@@ -690,33 +741,38 @@ rocke_for_t rocke_gfx950_attn2d_drive_kv_loop(rocke_gfx950_attn2d_build_ctx_t* c
         iter_args[i].init = ctx->iter_args[i];
     }
 
-    kvloop = rocke_b_scf_for_iter(b,
-                                  ctx->tile_start,
-                                  ctx->tile_end,
-                                  ctx->kv_step,
-                                  iter_args,
-                                  ctx->iter_args_count,
-                                  "kv_tile",
-                                  false,
-                                  true);
+    if(!(ctx->spec->use_mask_phase_split && ctx->TRANSPOSED_MASK_LIMIT && ctx->SLIDING_WINDOW == 0))
+        return drive_kv_phase(ctx, ctx->tile_start, ctx->tile_end, iter_args, "kv_tile", false);
 
-    rocke_b_region_enter(b, kvloop.body);
-    ctx->kv_tile_iv = kvloop.iv;
-    /* Unpack the carry the same way the Python body does (lines 2495-2508). */
-    for(i = 0; i < ctx->SOFTMAX_STATE_SLOTS; ++i)
-    {
-        ctx->m_cur[i] = kvloop.iter_vars[2 * i];
-        ctx->l_cur[i] = kvloop.iter_vars[2 * i + 1];
-    }
-    for(i = 0; i < num_accs; ++i)
-        ctx->acc_cur[i] = kvloop.iter_vars[ml + i];
-    ctx->cur_buf = kvloop.iter_vars[ml + num_accs];
-    ctx->skip_mask = false;
+    /* A tile needs no masking iff its last key is within the causal bound of
+     * the block's first query row and below max_seq_prefix_len, so the first
+     * masked tile is min((context_len + qb_start_pos + 1) / T,
+     * max_seq_prefix_len / T), clamped into [tile_start, tile_end]. */
+    rocke_value_t* min_causal_lim = rocke_b_add(b, ctx->context_len, ctx->qb_start_pos);
+    rocke_value_t* causal_p1 = rocke_b_add(b, min_causal_lim, rocke_b_const_i32(b, 1));
+    rocke_value_t* full_by_causal = rocke_b_div(b, causal_p1, rocke_b_const_i32(b, ctx->T));
+    rocke_value_t* full_by_prefix
+        = rocke_b_div(b, ctx->max_seq_prefix_len_v, rocke_b_const_i32(b, ctx->T));
+    rocke_value_t* split_raw = rocke_b_select(
+        b, rocke_b_cmp_lt(b, full_by_causal, full_by_prefix), full_by_causal, full_by_prefix);
+    rocke_value_t* split_lo = rocke_b_select(
+        b, rocke_b_cmp_lt(b, split_raw, ctx->tile_start), ctx->tile_start, split_raw);
+    rocke_value_t* split
+        = rocke_b_select(b, rocke_b_cmp_lt(b, split_lo, ctx->tile_end), split_lo, ctx->tile_end);
 
-    rocke_gfx950_attn2d_emit_kv_body(ctx);
+    rocke_for_t loop1 = drive_kv_phase(ctx, ctx->tile_start, split, iter_args, "kv_tile_f", true);
 
-    rocke_b_region_leave(b);
-    return kvloop;
+    /* The iter-arg names seed the loop's SSA names, so the second loop's must
+     * differ from the first's. */
+    rocke_iter_arg_t iter_args2[ROCKE_GFX950_ATTN2D_MAX_ITER_ARGS];
+    for(i = 0; i < ctx->iter_args_count; ++i)
+    {
+        char name[64];
+        snprintf(name, sizeof(name), "%s_b", ctx->iter_args_names[i]);
+        iter_args2[i].name = rocke_arena_strdup(&b->arena, name);
+        iter_args2[i].init = loop1.op->results[i];
+    }
+    return drive_kv_phase(ctx, split, ctx->tile_end, iter_args2, "kv_tile_b", false);
 }
 
 /* ============================================================ *
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_body_qk_softmax.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_body_qk_softmax.cpp
index 92562732a9d2..8d26982a9e62 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_body_qk_softmax.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_body_qk_softmax.cpp
@@ -166,8 +166,7 @@ static rocke_value_t* fh_read_k8_mfma_operand(rocke_gfx950_attn2d_build_ctx_t* c
                                               rocke_value_t* k_row,
                                               rocke_value_t* k_off)
 {
-    rocke_value_t* idx[3] = {buf, k_row, k_off};
-    return rocke_b_smem_load_vN(ctx->b, ctx->K_lds, idx, 3, ctx->dtype, 8);
+    return rocke_gfx950_attn2d_read_k8_mfma_operand(ctx, buf, k_row, k_off);
 }
 
 static rocke_value_t* fh_apply_softcap(rocke_gfx950_attn2d_build_ctx_t* ctx, rocke_value_t* s)
@@ -463,6 +462,46 @@ void rocke_gfx950_attn2d_reset_softmax_scratch(void)
     memset(p_regs_f32_buf, 0, sizeof(p_regs_f32_buf));
 }
 
+/* Post-QK prefetch issue: once QK no longer reads VMEM, start the loads the
+ * rest of the iteration (or the next one) needs. The branch order mirrors the
+ * Python builder; the PV wait counts depend on which stream is older. */
+static void emit_post_qk_issue(rocke_gfx950_attn2d_build_ctx_t* ctx, rocke_value_t* safe_next_tile)
+{
+    rocke_value_t* cur_buf = ctx->cur_buf;
+    rocke_value_t* nxt_buf = ctx->nxt_buf_v;
+    if(ctx->GROUPED_KV2)
+    {
+        rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
+        rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, cur_buf);
+    }
+    else if(ctx->EARLY_V_SCHEDULE)
+    {
+        rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, nxt_buf);
+    }
+    else if(ctx->V_DOUBLE_BUF && ctx->STAGGER_ITER_WAIT)
+    {
+        /* K first, so K[i+1] is the older stream at the next iter-start wait. */
+        rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, nxt_buf);
+        rocke_gfx950_attn2d_issue_v(ctx, safe_next_tile, nxt_buf);
+    }
+    else if(ctx->V_DOUBLE_BUF)
+    {
+        rocke_gfx950_attn2d_issue_v(ctx, safe_next_tile, nxt_buf);
+        rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, nxt_buf);
+    }
+    else if(ctx->K_SINGLE_BUFFER)
+    {
+        /* single K slot: issue V[i] now; DEFER the next-K prefetch to after the
+         * PV-wait barrier (avoids WAR-racing QK[i] ds_reads). */
+        rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
+    }
+    else
+    {
+        rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
+        rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, nxt_buf);
+    }
+}
+
 /* ===================================================================== *
  *  Front half of the QK + mask + softmax body.
  *
@@ -583,8 +622,18 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
     }
     bool transposed_path = (ctx->USE_MFMA_32X32 && ctx->TRANSPOSED_QK_32X32);
 
-    /* ---- wait for current K + LDS barrier ---- */
-    rocke_b_s_waitcnt(b, /*vmcnt=*/0, /*lgkmcnt=*/0, /*expcnt=*/-1);
+    /* ---- wait for current K + LDS barrier ---- *
+     * Staggered: the previous iter issued K[i] then V[i] into cur_buf, so drain
+     * only the older K[i] stream and leave V[i] in flight until the PV wait. */
+    if(ctx->STAGGER_ITER_WAIT)
+    {
+        const int kv_calls_per_tile = (ctx->T * ctx->HD) / (ctx->THREADS * 8);
+        rocke_b_s_waitcnt(b, kv_calls_per_tile, kv_calls_per_tile, /*expcnt=*/-1);
+    }
+    else
+    {
+        rocke_b_s_waitcnt(b, /*vmcnt=*/0, /*lgkmcnt=*/0, /*expcnt=*/-1);
+    }
     rocke_b_sync(b);
 
     /* STEP 2 lever (modes 0/1): canned iglp interleave for the whole loop body,
@@ -595,11 +644,17 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
        && (ctx->SOFTMAX_INTERLEAVE_MODE == 0 || ctx->SOFTMAX_INTERLEAVE_MODE == 1))
         rocke_b_iglp_opt(b, ctx->SOFTMAX_INTERLEAVE_MODE);
 
-    /* Early-V schedule: V_lds is single-buffered and the iter-start full drain
-     * guarantees the previous PV's V reads retired, so the current V copy can be
-     * issued before QK to overlap with QK + softmax (Python 2566-2570). */
+    /* Early-V schedule. Double-buffered V: V[i] is already resident, so
+     * prefetch V[i+1] into the alternate slot for the whole iter. Single-
+     * buffered V: the iter-start full drain retired the previous PV's V reads,
+     * so the current V copy is issued before QK to overlap QK + softmax. */
     if(ctx->EARLY_V_SCHEDULE)
-        rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, ctx->cur_buf);
+    {
+        if(ctx->V_DOUBLE_BUF)
+            rocke_gfx950_attn2d_issue_v(ctx, safe_next_tile, ctx->nxt_buf_v);
+        else
+            rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, ctx->cur_buf);
+    }
     /* GROUPED_KV2: prefetch the second K tile into nxt_buf while QK0 reads
      * cur_buf (Python 2571-2575). */
     if(ctx->GROUPED_KV2)
@@ -629,6 +684,21 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
             }
             ST32_n[n] = acc32;
         }
+        /* Interleave mode 2: alternating (MFMA, VALU) sched_group_barrier
+         * groups at the QK->softmax boundary, so the scheduler spreads the QK
+         * and PV MFMAs across the softmax VALU window. */
+        if(ctx->USE_SOFTMAX_INTERLEAVE && ctx->SOFTMAX_INTERLEAVE_MODE == 2)
+        {
+            const int groups
+                = ctx->SOFTMAX_INTERLEAVE_GROUPS > 1 ? ctx->SOFTMAX_INTERLEAVE_GROUPS : 1;
+            const int total_mfma = QK_N_TILES * QK_K_ITERS;
+            const int mfma_per_group = total_mfma / groups > 1 ? total_mfma / groups : 1;
+            for(int gi = 0; gi < groups; ++gi)
+            {
+                rocke_b_sched_group_barrier(b, 0x008, mfma_per_group, 0);
+                rocke_b_sched_group_barrier(b, 0x002, 100, 0);
+            }
+        }
         if(ctx->GROUPED_KV2)
         {
             rocke_b_s_waitcnt(b, /*vmcnt=*/0, /*lgkmcnt=*/0, /*expcnt=*/-1);
@@ -660,9 +730,8 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
         rocke_value_t* st_scores1[ROCKE_GFX950_ATTN2D_MAX_N_TILES][16];
         int n_groups = ctx->GROUPED_KV2 ? 2 : 1;
 
-        /* ---- TRANSPOSED_MASK_LIMIT setup (Python 2791-2820). skip_mask is
-         * always False in this single-loop body, so the not-skip_mask branch is
-         * the only one wired. Algebraically identical to the per-element
+        /* ---- TRANSPOSED_MASK_LIMIT setup (Python 2791-2820), skipped in the
+         * mask-free full-tile phase (skip_mask). Algebraically identical to the per-element
          * land(row_ok, col_abs <= valid_tail): fold row_ok into the threshold
          * (row invalid -> threshold = -BIG) and pre-subtract the compile-time
          * row_off so the per-element add folds away into a single cmp_le. ---- */
@@ -670,7 +739,8 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
         rocke_value_t* st_row_half_base = NULL;
         for(int _r = 0; _r < 16; ++_r)
             st_mask_thresh[_r] = NULL;
-        if(ctx->TRANSPOSED_MASK_LIMIT)
+        const bool mask_limit_masked = ctx->TRANSPOSED_MASK_LIMIT && !ctx->skip_mask;
+        if(mask_limit_masked)
         {
             rocke_value_t* mlim_row_ok
                 = ctx->TRANSPOSED_INVARIANT_HOIST ? ctx->st_row_ok_hoist : st_row_ok_iter;
@@ -730,7 +800,7 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
             {
                 /* MASK_LIMIT per-(group,n) col_abs_base (Python 2862-2865). */
                 rocke_value_t* col_abs_base = NULL;
-                if(ctx->TRANSPOSED_MASK_LIMIT)
+                if(mask_limit_masked)
                     col_abs_base
                         = rocke_b_add(b,
                                       rocke_b_add(b, group_tile_off, rocke_b_const_i32(b, n * 32)),
@@ -742,7 +812,13 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
                     rocke_value_t* causal_lim = NULL;
                     rocke_value_t* col_abs = NULL;
                     rocke_value_t* m_ok = NULL;
-                    if(ctx->TRANSPOSED_MASK_LIMIT)
+                    if(ctx->TRANSPOSED_MASK_LIMIT && ctx->skip_mask)
+                    {
+                        /* Full-tile phase: every key of this tile is inside the
+                         * causal + prefix bound for every row, so no mask. */
+                        m_ok = NULL;
+                    }
+                    else if(ctx->TRANSPOSED_MASK_LIMIT)
                     {
                         /* col_abs_base + row_off <= valid_tail  <=>
                          * col_abs_base <= (valid_tail_eff - row_off), with row_ok
@@ -824,7 +900,8 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
                     {
                         s_scaled = rocke_b_fmul(b, fh_apply_softcap(ctx, s_scaled), rcp_ln2);
                     }
-                    rocke_value_t* score = rocke_b_select(b, m_ok, s_scaled, neg_inf);
+                    rocke_value_t* score
+                        = m_ok == NULL ? s_scaled : rocke_b_select(b, m_ok, s_scaled, neg_inf);
                     if(ctx->USE_ALIBI)
                     {
                         rocke_value_t* pos_off = rocke_b_sub(b, col_abs, ctx->context_len);
@@ -904,28 +981,7 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
             rocke_b_sched_barrier(b, ctx->SCHED_BARRIER_MASK);
 
         /* post-QK V/K issue. */
-        rocke_value_t* cur_buf = ctx->cur_buf;
-        rocke_value_t* nxt_buf = ctx->nxt_buf_v;
-        if(ctx->GROUPED_KV2)
-        {
-            rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
-            rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, cur_buf);
-        }
-        else if(ctx->EARLY_V_SCHEDULE)
-        {
-            rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, nxt_buf);
-        }
-        else if(ctx->K_SINGLE_BUFFER)
-        {
-            /* single K slot: issue V[i] now; DEFER the next-K prefetch to
-             * after the PV-wait barrier (avoids WAR-racing QK[i] ds_reads). */
-            rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
-        }
-        else
-        {
-            rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
-            rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, nxt_buf);
-        }
+        emit_post_qk_issue(ctx, safe_next_tile);
 
         for(int r = 0; r < SOFTMAX_STATE_SLOTS; ++r)
             alpha_regs[r] = rocke_b_exp2(b, rocke_b_fsub(b, ctx->m_cur[r], m_new[r]));
@@ -971,29 +1027,7 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
             rocke_b_sched_barrier(b, ctx->SCHED_BARRIER_MASK);
 
         /* post-QK V/K issue (Python 2966-2975). */
-        {
-            rocke_value_t* cur_buf = ctx->cur_buf;
-            rocke_value_t* nxt_buf = ctx->nxt_buf_v;
-            if(ctx->GROUPED_KV2)
-            {
-                rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
-                rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, cur_buf);
-            }
-            else if(ctx->EARLY_V_SCHEDULE)
-            {
-                rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, nxt_buf);
-            }
-            else if(ctx->K_SINGLE_BUFFER)
-            {
-                /* single K slot: V[i] now; defer next-K to post-PV barrier. */
-                rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
-            }
-            else
-            {
-                rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
-                rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, nxt_buf);
-            }
-        }
+        emit_post_qk_issue(ctx, safe_next_tile);
 
         /* mask / scale / softcap / alibi / qq-bias (Python 2998-3041). */
         for(int reg = 0; reg < REGS_PER_LANE; ++reg)
@@ -1156,28 +1190,7 @@ void rocke_gfx950_attn2d_emit_kv_body(rocke_gfx950_attn2d_build_ctx_t* ctx)
             rocke_b_sched_barrier(b, ctx->SCHED_BARRIER_MASK);
 
         /* post-QK V/K issue. */
-        rocke_value_t* cur_buf = ctx->cur_buf;
-        rocke_value_t* nxt_buf = ctx->nxt_buf_v;
-        if(ctx->GROUPED_KV2)
-        {
-            rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
-            rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, cur_buf);
-        }
-        else if(ctx->EARLY_V_SCHEDULE)
-        {
-            rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, nxt_buf);
-        }
-        else if(ctx->K_SINGLE_BUFFER)
-        {
-            /* single K slot: issue V[i] now; DEFER the next-K prefetch to
-             * after the PV-wait barrier (avoids WAR-racing QK[i] ds_reads). */
-            rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
-        }
-        else
-        {
-            rocke_gfx950_attn2d_issue_v(ctx, ctx->kv_tile_iv, cur_buf);
-            rocke_gfx950_attn2d_issue_k(ctx, safe_next_tile, nxt_buf);
-        }
+        emit_post_qk_issue(ctx, safe_next_tile);
 
         /* mask / scale / softcap / alibi / qq-bias. */
         rocke_value_t* masked[ROCKE_GFX950_ATTN2D_MAX_N_TILES]
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_descriptors_loads.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_descriptors_loads.cpp
index 82129cb1b3bd..d071d1ed2716 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_descriptors_loads.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_kv_descriptors_loads.cpp
@@ -175,6 +175,45 @@ static rocke_value_t* rocke__wave_lds_offset_i64(rocke_gfx950_attn2d_build_ctx_t
     return ctx->wave_lds_off_i64_v;
 }
 
+/* Slab-padded K strides: one wave-slab is WAVE_BYTES contiguous bytes, and
+ * KQ_LDS_PAD*2 bytes go between slabs (Python WAVE_BYTES_K / bytes_per_call_K /
+ * bytes_per_buf_K). Without the pad these equal the shared K/V strides. */
+static int rocke__wave_bytes_K(const rocke_gfx950_attn2d_build_ctx_t* ctx)
+{
+    return rocke__wave_bytes(ctx) + ctx->KQ_LDS_PAD * 2;
+}
+
+static int rocke__bytes_per_call_K(const rocke_gfx950_attn2d_build_ctx_t* ctx)
+{
+    return ctx->KQ_LDS_PAD ? ctx->NUM_WARPS * rocke__wave_bytes_K(ctx) : rocke__bytes_per_call(ctx);
+}
+
+static int rocke__bytes_per_buf_K(const rocke_gfx950_attn2d_build_ctx_t* ctx)
+{
+    return ctx->KQ_LDS_PAD ? (ctx->T / ctx->KQ_SLAB_ROWS) * rocke__wave_bytes_K(ctx)
+                           : rocke__bytes_per_buf(ctx);
+}
+
+/* wave_lds_offset_K_i64: its own SSA value under the slab pad (emitted right
+ * after wave_lds_offset_i64), else the shared wave offset. */
+static rocke_value_t* rocke__wave_lds_offset_K_i64(rocke_gfx950_attn2d_build_ctx_t* ctx)
+{
+    if(!ctx->KQ_LDS_PAD)
+        return rocke__wave_lds_offset_i64(ctx);
+    if(ctx->wave_lds_off_K_i64_v != NULL)
+        return ctx->wave_lds_off_K_i64_v;
+    if(ctx->NUM_WARPS == 1)
+    {
+        ctx->wave_lds_off_K_i64_v = rocke_b_const_i64(ctx->b, 0);
+        return ctx->wave_lds_off_K_i64_v;
+    }
+    rocke_value_t* off_i32 = rocke_b_to_sgpr_u32(
+        ctx->b,
+        rocke_b_mul(ctx->b, ctx->wave_id, rocke_b_const_i32(ctx->b, rocke__wave_bytes_K(ctx))));
+    ctx->wave_lds_off_K_i64_v = rocke_b_zext(ctx->b, off_i32, rocke_i64());
+    return ctx->wave_lds_off_K_i64_v;
+}
+
 /* ==========================================================================
  * Pre-loop: K/V buffer resources, byte derivation, paged-KV descriptor
  * (Python lines 1163-1630, the K/V-descriptor parts).
@@ -218,6 +257,7 @@ void rocke_gfx950_attn2d_emit_preloop(rocke_gfx950_attn2d_build_ctx_t* ctx)
      * otherwise their value numbers land later and shift the downstream stream. */
     (void)rocke__zero_soff(ctx);
     (void)rocke__wave_lds_offset_i64(ctx);
+    (void)rocke__wave_lds_offset_K_i64(ctx);
 
     /* seq_base (1525) + block_table_max_idx (1538): the per-sequence base index
      * into block_tables and the global footprint bound; both wave-uniform SGPRs.
@@ -387,13 +427,13 @@ void rocke_gfx950_attn2d_issue_k_load_runtime(rocke_gfx950_attn2d_build_ctx_t* c
                                               rocke_value_t* buf_idx)
 {
     rocke_ir_builder_t* b = ctx->b;
-    int bytes_per_buf = rocke__bytes_per_buf(ctx);
-    int bytes_per_call = rocke__bytes_per_call(ctx);
+    int bytes_per_buf = rocke__bytes_per_buf_K(ctx);
+    int bytes_per_call = rocke__bytes_per_call_K(ctx);
     int kv_calls_per_tile = rocke__kv_calls_per_tile(ctx);
     int kv_halves_per_call = rocke__kv_halves_per_call(ctx);
 
     rocke_value_t* K_lds_addr = rocke__K_lds_addr(ctx);
-    rocke_value_t* wave_off = rocke__wave_lds_offset_i64(ctx);
+    rocke_value_t* wave_off = rocke__wave_lds_offset_K_i64(ctx);
     rocke_value_t* zero_soff = rocke__zero_soff(ctx);
     rocke_value_t* kv_block_bytes_c = rocke__kv_block_bytes_c(ctx);
     rocke_value_t* lane_half_base = rocke__lane_half_base(ctx);
@@ -474,15 +514,28 @@ void rocke_gfx950_attn2d_issue_v_load_runtime(rocke_gfx950_attn2d_build_ctx_t* c
     int kv_calls_per_tile = rocke__kv_calls_per_tile(ctx);
     int kv_halves_per_call = rocke__kv_halves_per_call(ctx);
 
-    (void)buf_idx; /* V is single-buffered; always slot 0 (Python 1699). */
     rocke_value_t* V_lds_addr = rocke__V_lds_addr(ctx);
     rocke_value_t* wave_off = rocke__wave_lds_offset_i64(ctx);
     rocke_value_t* zero_soff = rocke__zero_soff(ctx);
     rocke_value_t* kv_block_bytes_c = rocke__kv_block_bytes_c(ctx);
     rocke_value_t* lane_half_base = rocke__lane_half_base(ctx);
 
-    /* V_wave_base = V_lds_addr + wave_lds_offset (1700). */
-    rocke_value_t* V_wave_base = rocke_b_smem_ptr_add(b, V_lds_addr, wave_off);
+    /* V double-buffer: buf_idx picks the V_lds slot (T*HD*2 bytes apart);
+     * single-buffered V ignores it and always writes slot 0. */
+    rocke_value_t* V_wave_base;
+    if(ctx->V_DOUBLE_BUF)
+    {
+        rocke_value_t* v_buf_off_i64
+            = rocke_b_zext(b,
+                           rocke_b_mul(b, buf_idx, rocke_b_const_i32(b, rocke__bytes_per_buf(ctx))),
+                           rocke_i64());
+        V_wave_base
+            = rocke_b_smem_ptr_add(b, rocke_b_smem_ptr_add(b, V_lds_addr, v_buf_off_i64), wave_off);
+    }
+    else
+    {
+        V_wave_base = rocke_b_smem_ptr_add(b, V_lds_addr, wave_off);
+    }
 
     rocke_value_t* fast_block0 = NULL;
     rocke_value_t* fast_block1 = NULL;
@@ -599,6 +652,17 @@ rocke_value_t* rocke_gfx950_attn2d_read_k8_mfma_operand(rocke_gfx950_attn2d_buil
     idx[1] = k_row;
     idx[2] = k_off;
 
+    /* Slab-padded K_lds: map (row, col) to (slab, row_in_slab*HD + col) so the
+     * read hits the slot the DMA wrote. */
+    if(ctx->KQ_LDS_PAD && !ctx->FP8_MFMA_QK)
+    {
+        rocke_value_t* slab = rocke_b_div(b, k_row, rocke_b_const_i32(b, ctx->KQ_SLAB_ROWS));
+        rocke_value_t* row_in_slab = rocke_b_mod(b, k_row, rocke_b_const_i32(b, ctx->KQ_SLAB_ROWS));
+        rocke_value_t* row_off = rocke_b_mul(b, row_in_slab, rocke_b_const_i32(b, ctx->HD));
+        idx[1] = slab;
+        idx[2] = rocke_b_add(b, row_off, k_off);
+        return rocke_b_smem_load_vN(b, ctx->K_lds, idx, 3, ctx->dtype, 8);
+    }
     /* not K_FP8_MFMA: plain bf16 read (Python 2241-2242). */
     if(!ctx->FP8_MFMA_QK)
         return rocke_b_smem_load_vN(b, ctx->K_lds, idx, 3, ctx->dtype, 8);
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_public_entry_glue.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_public_entry_glue.cpp
index 859ece99e9d4..041de43d2169 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_public_entry_glue.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_public_entry_glue.cpp
@@ -449,30 +449,30 @@ static bool rocke_g950_build_ctx_init_local(rocke_gfx950_attn2d_build_ctx_t* ctx
         rocke_g950_fail(b, ROCKE_ERR_VALUE, "use_fast_paged_kv_desc is restricted to dtype='bf16'");
     }
 
-    /* ---- V-double-buffer / deep-ring / staggered-wait / q-reread.
-     * Python __post_init__ accepts these flags (and the emitter wires the
-     * depth-2 V-double-buffer + staggered iter-start schedule). The gfx950 C
-     * twin currently emits ONLY the depth-2 single/early-V schedule, so reject
-     * any spec that asks for the V-double-buffer family (and kv_ring_depth!=2 /
-     * q_reread) rather than silently emit the wrong schedule. This keeps the
-     * spec struct byte-identical with Python (provider passthrough) while making
-     * the C emitter's narrower coverage explicit. When these schedules are
-     * ported to C, replace these guards with the real emit + Python's exact
-     * __post_init__ messages. */
+    /* ---- V double-buffer / staggered iter-start wait (Python __post_init__). */
     if(spec->use_v_double_buffer)
     {
-        rocke_g950_fail(b,
-                        ROCKE_ERR_NOTIMPL,
-                        "use_v_double_buffer: gfx950 C twin emits only the depth-2 "
-                        "single/early-V schedule (not yet ported)");
+        if(spec->use_grouped_kv2_softmax)
+            rocke_g950_fail(b, ROCKE_ERR_VALUE, "use_v_double_buffer does not support grouped_kv2");
+        if(spec->kv_storage_dtype != NULL)
+            rocke_g950_fail(b, ROCKE_ERR_VALUE, "use_v_double_buffer v1 does not support FP8 KV");
     }
     if(spec->use_staggered_iter_wait)
     {
-        rocke_g950_fail(b,
-                        ROCKE_ERR_NOTIMPL,
-                        "use_staggered_iter_wait: gfx950 C twin emits only the depth-2 "
-                        "single/early-V schedule (not yet ported)");
+        if(!spec->use_v_double_buffer)
+            rocke_g950_fail(
+                b, ROCKE_ERR_VALUE, "use_staggered_iter_wait requires use_v_double_buffer");
+        if(spec->use_grouped_kv2_softmax)
+            rocke_g950_fail(
+                b, ROCKE_ERR_VALUE, "use_staggered_iter_wait does not support grouped_kv2");
+        if(spec->kv_storage_dtype != NULL)
+            rocke_g950_fail(
+                b, ROCKE_ERR_VALUE, "use_staggered_iter_wait v1 does not support FP8 KV");
     }
+
+    /* ---- deep ring / q-reread: the Python builder emits these schedules but
+     * the gfx950 C twin does not port them yet. Reject rather than silently
+     * emit the depth-2 schedule. */
     if(spec->kv_ring_depth != 2)
     {
         /* the Python builder wires a depth-3 deep K prefetch ring on the
@@ -490,28 +490,40 @@ static bool rocke_g950_build_ctx_init_local(rocke_gfx950_attn2d_build_ctx_t* ctx
                         ROCKE_ERR_NOTIMPL,
                         "use_q_reread: gfx950 C twin does not yet port the Q-reread path");
     }
+    /* ---- direct-to-register Q gather / mask phase split (Python __post_init__). */
     if(spec->use_q_direct_reg)
     {
-        /* the direct-to-register Q gather (Triton-style, frees Q_lds) is
-         * carried in the spec for byte-identity with the Python dataclass but
-         * the gfx950 C twin does not yet emit this body. Reject (do not
-         * silently ignore), matching the Python __post_init__ contract. */
-        rocke_g950_fail(
-            b,
-            ROCKE_ERR_NOTIMPL,
-            "use_q_direct_reg: gfx950 C twin does not yet port the direct-register Q gather");
+        if(!(spec->use_mfma_32x32 && spec->use_transposed_qk_32x32))
+            rocke_g950_fail(b,
+                            ROCKE_ERR_VALUE,
+                            "use_q_direct_reg requires the transposed-32x32 path "
+                            "(use_mfma_32x32 + use_transposed_qk_32x32)");
+        if(spec->use_q_reread)
+            rocke_g950_fail(b,
+                            ROCKE_ERR_VALUE,
+                            "use_q_direct_reg is mutually exclusive with use_q_reread "
+                            "(direct-reg never stages Q in LDS; re-read needs a Q_lds)");
+        if(spec->kv_storage_dtype != NULL)
+            rocke_g950_fail(b, ROCKE_ERR_VALUE, "use_q_direct_reg v1 does not support FP8 KV");
+        if(spec->use_fp8_mfma_qk)
+            rocke_g950_fail(b,
+                            ROCKE_ERR_VALUE,
+                            "use_q_direct_reg v1 does not support native-fp8 QK "
+                            "(Q stays in the working dtype)");
     }
-    if(spec->use_softmax_mfma_interleave && spec->softmax_interleave_mode == 2)
+    if(spec->use_mask_phase_split)
     {
-        /* Modes 0/1 (iglp_opt at the loop top) ARE ported below -- a single
-         * iglp_opt call after the iter-start K drain (mirrors Python 3093-3094).
-         * Mode 2 (alternating sched_group_barrier grouping at the QK->softmax
-         * boundary) is a distinct emission the gfx950 C twin does not yet port;
-         * reject it (do not silently ignore) to match the Python contract. */
-        rocke_g950_fail(b,
-                        ROCKE_ERR_NOTIMPL,
-                        "use_softmax_mfma_interleave mode 2 (sched_group_barrier grouping): "
-                        "gfx950 C twin ports iglp_opt modes 0/1 only");
+        if(!(spec->use_mfma_32x32 && spec->use_transposed_qk_32x32
+             && spec->use_transposed_mask_limit))
+            rocke_g950_fail(b,
+                            ROCKE_ERR_VALUE,
+                            "use_mask_phase_split requires the transposed-32x32 "
+                            "mask_limit combo path");
+        if(spec->sliding_window > 0)
+            rocke_g950_fail(b,
+                            ROCKE_ERR_VALUE,
+                            "use_mask_phase_split requires no sliding window "
+                            "(the peel assumes a monotone causal bound)");
     }
     if(spec->use_softmax_mfma_interleave && spec->use_sched_barrier)
     {
@@ -598,10 +610,13 @@ static bool rocke_g950_build_ctx_init_local(rocke_gfx950_attn2d_build_ctx_t* ctx
     ctx->EARLY_V_SCHEDULE = spec->use_early_v_schedule;
     ctx->AGPR_ALLOC_ZERO = spec->use_agpr_alloc_zero;
     ctx->K_SINGLE_BUFFER = spec->use_k_single_buffer; /* K single-buffer */
+    ctx->V_DOUBLE_BUF = spec->use_v_double_buffer;
+    ctx->STAGGER_ITER_WAIT = spec->use_staggered_iter_wait;
     ctx->USE_SCHED_BARRIER = spec->use_sched_barrier;
     ctx->SCHED_BARRIER_MASK = spec->sched_barrier_mask;
     ctx->USE_SOFTMAX_INTERLEAVE = spec->use_softmax_mfma_interleave;
     ctx->SOFTMAX_INTERLEAVE_MODE = spec->softmax_interleave_mode;
+    ctx->SOFTMAX_INTERLEAVE_GROUPS = spec->softmax_interleave_groups;
 
     /* ---- fp8 K/V cache predicates (Py783-797) ---- */
     const bool KV_FP8 = rocke_g950_streq(spec->kv_storage_dtype, "fp8e4m3");
@@ -815,19 +830,51 @@ static bool rocke_g950_build_ctx_init_local(rocke_gfx950_attn2d_build_ctx_t* ctx
     ctx->K_BUFS = ctx->K_SINGLE_BUFFER ? 1 : 2;
     ctx->K_TOTAL_BYTES = K_TOTAL_BYTES;
 
-    const bool Q_ALIAS_K = rocke_type_eq(K_LDS_DTYPE, dtype) && (Q_BYTES <= K_TOTAL_BYTES);
+    /* K_lds slab pad: one wave's async-DMA slab (1024 B) fills KQ_SLAB_ROWS
+     * rows; padding between slabs keeps the DMA write contiguous. The read
+     * remap exists only for the bf16 K path with a single K slot, so the pad
+     * is dropped under fp8 QK, unsupported geometry, or K double-buffering. */
+    {
+        int pad = spec->use_kq_lds_pad ? spec->kq_lds_pad_halves : 0;
+        int slab_rows = pad ? 512 / HD : 0;
+        if(pad
+           && (FP8_MFMA_QK || 512 % HD != 0 || slab_rows == 0 || T % slab_rows != 0
+               || pad % 8 != 0))
+        {
+            pad = 0;
+            slab_rows = 0;
+        }
+        if(pad && ctx->K_BUFS != 1)
+        {
+            pad = 0;
+            slab_rows = 0;
+        }
+        ctx->KQ_LDS_PAD = pad;
+        ctx->KQ_SLAB_ROWS = slab_rows;
+    }
+
+    bool Q_ALIAS_K = rocke_type_eq(K_LDS_DTYPE, dtype) && (Q_BYTES <= K_TOTAL_BYTES);
+    ctx->Q_DIRECT_REG = spec->use_q_direct_reg;
+    if(ctx->Q_DIRECT_REG)
+        Q_ALIAS_K = false; /* Q never touches LDS */
     const bool Q_USES_DUAL_SLOT = Q_ALIAS_K && (BLOCK_M > T);
     ctx->Q_DIRECT_GLOBAL = false; /* gfx950 has no Q_DIRECT_GLOBAL path */
     ctx->Q_ALIAS_K = Q_ALIAS_K;
     ctx->Q_USES_DUAL_SLOT = Q_USES_DUAL_SLOT;
 
     /* ---- K_lds smem_alloc (Py1038) ---- */
+    if(ctx->KQ_LDS_PAD)
+    {
+        int shp[3] = {ctx->K_BUFS, T / ctx->KQ_SLAB_ROWS, ctx->KQ_SLAB_ROWS * HD + ctx->KQ_LDS_PAD};
+        ctx->K_lds = rocke_b_smem_alloc(b, K_LDS_DTYPE, shp, 3, "Klds");
+    }
+    else
     {
         int shp[3] = {ctx->K_BUFS, T, HD}; /* 1 slot when K single-buffer */
         ctx->K_lds = rocke_b_smem_alloc(b, K_LDS_DTYPE, shp, 3, "Klds");
     }
 
-    const int V_BUFS = 1; /* single-buffer V (race-free, Py1039) */
+    const int V_BUFS = ctx->V_DOUBLE_BUF ? 2 : 1;
     ctx->V_BUFS = V_BUFS;
 
     /* ---- V_lds smem_alloc (Py1040-1060) ---- */
@@ -872,7 +919,11 @@ static bool rocke_g950_build_ctx_init_local(rocke_gfx950_attn2d_build_ctx_t* ctx
     }
 
     /* ---- Q_lds (Py1098-1104) ---- */
-    if(Q_ALIAS_K)
+    if(ctx->Q_DIRECT_REG)
+    {
+        ctx->Q_lds = NULL;
+    }
+    else if(Q_ALIAS_K)
     {
         ctx->Q_lds = ctx->K_lds;
     }
@@ -1105,6 +1156,7 @@ static rocke_status_t rocke_g950_attn2d_kernel_name(const rocke_attention_tiled_
             s->use_early_v_schedule ? "earlyv" : "",
             s->use_v_double_buffer ? "vdbuf" : "",
             s->use_k_single_buffer ? "ksb" : "", /* K single-buffer */
+            s->use_mask_phase_split ? "mphsplit" : "",
             s->use_staggered_iter_wait ? "stgw" : "",
             schedb_buf,
             smxil_buf,
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_q_load_gather.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_q_load_gather.cpp
index 81b1e09a4da5..2ff82ca4a8ae 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_q_load_gather.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx950/attention_tiled_2d_q_load_gather.cpp
@@ -103,12 +103,15 @@ void rocke_gfx950_attn2d_emit_q_load(rocke_gfx950_attn2d_build_ctx_t* ctx)
      * cooperative loop; emit it here to reproduce the exact IR order (it is not a
      * shared ctx field -- its sole consumer is the staged-store fallback below). */
     z8 = rocke_b_zero_vec(b, dtype, 8);
+    ctx->z8_v = z8;
 
     /* ---------------- Q -> LDS (cooperative vec8 chunks) ----------------
      * Q_VECS_PER_ROW = HD // 8 ; Q_VECS_PER_THREAD = (BLOCK_M * Q_VECS_PER_ROW)
      * // THREADS. Both are host-only ints. */
     Q_VECS_PER_ROW = ctx->HD / 8;
-    Q_VECS_PER_THREAD = (ctx->BLOCK_M * Q_VECS_PER_ROW) / ctx->THREADS;
+    /* direct-reg: no cooperative Q->LDS staging; the Q32 operand is gathered
+     * straight from global in emit_q_gather. */
+    Q_VECS_PER_THREAD = ctx->Q_DIRECT_REG ? 0 : (ctx->BLOCK_M * Q_VECS_PER_ROW) / ctx->THREADS;
 
     /* q_desc = TensorDescriptor.naive("Q", lengths=[1<<30, NUM_QH, HD],
      *          coord_names=("token","head","dim")). Stored on ctx for the
@@ -413,6 +416,40 @@ void rocke_gfx950_attn2d_emit_loop_bounds_and_inits(rocke_gfx950_attn2d_build_ct
     ctx->iter_args_count = n_args;
 }
 
+/* Python _read_q32_global(k): the direct-to-register Q32 operand gather. */
+static rocke_value_t* q_read_q32_global(rocke_gfx950_attn2d_build_ctx_t* ctx,
+                                        rocke_value_t* q32_row,
+                                        rocke_value_t* lane_half,
+                                        int k)
+{
+    rocke_ir_builder_t* b = ctx->b;
+    rocke_value_t* q32_col_base = rocke_b_const_i32(b, (int64_t)(k * 16));
+    rocke_value_t* q32_col
+        = rocke_b_add(b, q32_col_base, rocke_b_mul(b, lane_half, rocke_b_const_i32(b, 8)));
+    rocke_value_t* qg_pos = rocke_b_add(
+        b, ctx->qb_start_pos, rocke_b_div(b, q32_row, rocke_b_const_i32(b, ctx->NQK)));
+    rocke_value_t* qg_h_mul = rocke_b_mul(b, ctx->kv_head_idx, rocke_b_const_i32(b, ctx->NQK));
+    rocke_value_t* qg_h_mod = rocke_b_mod(b, q32_row, rocke_b_const_i32(b, ctx->NQK));
+    rocke_value_t* qg_h = rocke_b_add(b, qg_h_mul, qg_h_mod);
+    rocke_value_t* qg_ok_pos = rocke_b_cmp_lt(b, qg_pos, ctx->cur_batch_q_len);
+    rocke_value_t* qg_ok_h = rocke_b_cmp_lt(b, qg_h, rocke_b_const_i32(b, ctx->NUM_QH));
+    rocke_value_t* qg_mask = rocke_b_land(b, qg_ok_pos, qg_ok_h);
+    rocke_value_t* qg_pos_safe = rocke_b_select(b, qg_mask, qg_pos, rocke_b_const_i32(b, 0));
+    rocke_value_t* qg_h_safe = rocke_b_select(b, qg_mask, qg_h, rocke_b_const_i32(b, 0));
+    rocke_value_t* qg_off = NULL;
+    {
+        static const char* const in_names[3] = {"token", "head", "dim"};
+        rocke_value_t* in_values[3];
+        in_values[0] = rocke_b_add(b, ctx->cu_q_start, qg_pos_safe);
+        in_values[1] = qg_h_safe;
+        in_values[2] = rocke_b_const_i32(b, 0);
+        rocke_transforms_descriptor_offset(b, ctx->q_desc, in_names, in_values, 3, &qg_off, NULL);
+    }
+    rocke_value_t* v8
+        = rocke_b_global_load_vN(b, ctx->query, rocke_b_add(b, qg_off, q32_col), ctx->dtype, 8, 16);
+    return rocke_b_vector_select(b, rocke_b_vector_splat(b, qg_mask, 8), v8, ctx->z8_v);
+}
+
 /* ============================================================ *
  *  Per-lane Q -> VGPR gather (Python lines 2250-2363, excluding the trailing
  *  _issue_k(tile_start, 0) prefetch which the preloop bucket owns).
@@ -446,7 +483,7 @@ void rocke_gfx950_attn2d_emit_q_gather(rocke_gfx950_attn2d_build_ctx_t* ctx)
     ctx->q_regs_count = ctx->M_ATOMS_PER_WARP * ctx->QK_K_ITERS;
 
     /* if not (USE_MFMA_32X32 and SKIP_LEGACY_QREG): (line 2258) */
-    if(!(ctx->USE_MFMA_32X32 && ctx->SKIP_LEGACY_QREG))
+    if(!(ctx->USE_MFMA_32X32 && ctx->SKIP_LEGACY_QREG) && !ctx->Q_DIRECT_REG)
     {
         for(atom = 0; atom < ctx->M_ATOMS_PER_WARP; ++atom)
         {
@@ -573,7 +610,17 @@ void rocke_gfx950_attn2d_emit_q_gather(rocke_gfx950_attn2d_build_ctx_t* ctx)
             }
         }
 
-        for(k = 0; k < ctx->QK_K_ITERS; ++k)
+        /* direct-to-register Q gather: load this lane's 8-element Q32 operand
+         * straight from HBM, indexing the BLOCK_M x HD Q tile exactly as the
+         * LDS-staged read would. Out-of-range rows read zero, like the staging
+         * store's masked fill. */
+        if(ctx->Q_DIRECT_REG)
+        {
+            for(k = 0; k < ctx->QK_K_ITERS; ++k)
+                ctx->q32_regs[k] = q_read_q32_global(ctx, q32_row, lane_half, k);
+        }
+
+        for(k = 0; k < (ctx->Q_DIRECT_REG ? 0 : ctx->QK_K_ITERS); ++k)
         {
             rocke_value_t* q32_col;
             rocke_value_t* q32;
@@ -628,6 +675,10 @@ void rocke_gfx950_attn2d_emit_preloop_prefetch(rocke_gfx950_attn2d_build_ctx_t*
 
     /* Prefetch tile_start's K into buffer 0 BEFORE the loop (2384). */
     rocke_gfx950_attn2d_issue_k(ctx, ctx->tile_start, rocke_b_const_i32(B, 0));
+    /* V double-buffer: V[tile_start] goes into slot 0 too, so iter 0 finds it
+     * ready; each iter then prefetches V[i+1] into the alternate slot. */
+    if(ctx->V_DOUBLE_BUF)
+        rocke_gfx950_attn2d_issue_v(ctx, ctx->tile_start, rocke_b_const_i32(B, 0));
 
     /* cur_buf_init = const_i32(0); iter_args.append(("cur_buf", cur_buf_init))
      * (2399-2400). */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md
index 3332da5eea9e..f8135119f539 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/development/known_gaps.md
@@ -281,25 +281,26 @@ llvm23.
 
 ---
 
-## 8. gfx950 tiled 2D — V double-buffer and deep-prefetch schedules
+## 8. gfx950 tiled 2D — deep-prefetch ring and Q-reread schedules
 
-**Subsystem:** `instances/gfx950/attention_tiled_2d_public_entry_glue.cpp`, `instances/common/attention_unified_selectors.cpp`
+**Subsystem:** `instances/gfx950/attention_tiled_2d_public_entry_glue.cpp`
 
 The Python gfx950 tiled-2D builder emits these schedules. The C++ builder
-emits only the depth-2 single / early-V schedule and rejects each of these
-spec fields with `ROCKE_ERR_NOTIMPL`. These rejects are not faithful
-mirrors of Python behavior. They are deliberate rejects that stop C++
-from silently building a different schedule.
-
-The byte-identity gate stays green because no parity-emitter config sets any
-of these fields. The C++ selector **does** reach one of them:
-`rocke_unified_attn_tiled_spec_from_problem` sets `use_v_double_buffer` exactly
-as Python's `_enable_v_double_buffer` does (single-batch d64 bf16/fp16 prefill
-with `256 < max_seqlen_q <= 1024`). The selector is intentionally kept in step
-with Python, so selector parity holds. Building that spec in C++ fails loudly.
+rejects each of these spec fields with `ROCKE_ERR_NOTIMPL`. These rejects are
+not faithful mirrors of Python behavior. They are deliberate rejects that stop
+C++ from silently building a different schedule. The C++ selector never sets
+either field, and Python marks both experimental.
+
+The V double-buffer (`use_v_double_buffer`) and the staggered iter-start wait
+(`use_staggered_iter_wait`) are ported and pinned by the gfx950 tiled-2D parity
+emitter.
 
 | Spec field | Selected by the shared selector? | What unblocks it |
 |---|---|---|
+| `kv_ring_depth != 2` | No | Port the 3-slot deep K prefetch ring. |
+| `use_q_reread` | No | Port the Q-reread path. |
+
+---|---|---|
 | `use_v_double_buffer` (and `use_sched_barrier`, which gates on it) | Yes, see above | Port the V[i+1] prefetch slots and the double-buffer PV schedule, then drop the reject. |
 | `use_staggered_iter_wait` | No | Port the staggered-wait schedule together with the above. |
 | `kv_ring_depth != 2` | No | Port the 3-slot deep K prefetch ring. |
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp
index c46cd6e1940b..5ee241292538 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp
@@ -284,7 +284,8 @@ static void test_supports_tiled_2d(void)
     check_tiled_2d(
         &p, "gfx942", false, "gfx942 tiled 2D flash mask-limit is not ported to C++ yet");
 
-    /* gfx950: single-batch fp16 D64 MHA picks the V double-buffer. */
+    /* gfx950: single-batch fp16 D64 MHA picks the V double-buffer, which the
+     * C++ builder builds. */
     p = decode_problem();
     p.dtype = "fp16";
     p.head_size = 64;
@@ -292,7 +293,7 @@ static void test_supports_tiled_2d(void)
     p.num_seqs = 1;
     p.max_seqlen_q = 512;
     p.total_q = 512;
-    check_tiled_2d(&p, "gfx950", false, "tiled 2D V double-buffer is not ported to C++ yet");
+    check_tiled_2d(&p, "gfx950", true, "supported");
 
     p = decode_problem();
     p.dtype = "fp16";
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json
index bd9d69c9c737..da32b647b59c 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json
@@ -1361,7 +1361,18 @@
       "56": "9bc58bd47c222fbaf412aad80d84a82eb27ffaa7497cc873c2a93d1ddd3192cb",
       "57": "343ab03ba213210079cc94cc21c835829d1f4772ac6f2b0577fab39779a86e15",
       "58": "3e67da9078dfe812e834eb595cf6463669bd265c6d6a6547d30178fba736bc6e",
+      "59": "52004d3d845f050a446a5f6487c73260386b3784dfaeec9d4bf7d92457d9a3e8",
       "6": "555adce2cbd9bb832d640131769903e5d639db70bc07461eea5c8a59c2358718",
+      "60": "8298f45da3527ac6f2abdf1f91c176b20bbeb121d89cbd7ccfb0f7db5d346d49",
+      "61": "e2fca2f115d1f3803e0a7191ecae8592eb2f4f64764134023301c2f1327bdb8e",
+      "62": "d4715fd1a9a66de53ed4cb8b40501c65b747c7a5312a292e67fd96b1c2784c5f",
+      "63": "35f5e3b35661cae0cdc6b8754b79a2f81f69fb402445de1a24010b9217b6964b",
+      "64": "bce57098d15ca0d17f26eba3327969bacc35584b59353716b3379712f9c57239",
+      "65": "533a0bf77dd549341ea92f502a6b4f5b8bd213af5c0198c03bb0c9fa1cb6c575",
+      "66": "b0fd587c73e803297f8686b388d2695cf40ef1d77b237c4f7a05955dd2ed2a8b",
+      "67": "a6a6db53ee92c088897becf02724fb61bcd9c09951bf5ebb56303baeb0e2318d",
+      "68": "58107eead082bbf925d9bad75329114554140afa9f406bacf59e0d4076a14b30",
+      "69": "00eb4118ff9353fd230681d8fc026ecc5a42b81af70637d602d366235a381ace",
       "7": "bbcd9758625b990d8d78cbc25e11756950e85d7db44c1b5050724f0efd721b1f",
       "8": "cea4c8a1c2e10ab7d794c144103859fb72be23ab8cd875339df07d9fb6965aa4",
       "9": "7180bac09e314b919c16242218d43935eaf463660488b5f7986fdefab664dca5"
@@ -3022,7 +3033,18 @@
         "56": "ce2854f67d3b8af76aa7a2946be2417217a743092b7b49538d175fbe9665b592",
         "57": "643f555d5c1d0bec4a7399b1d724d4089424d99c6c9bd74fcf1435815e7ef6a1",
         "58": "21eed5600fb30fa4a91786bbc5084450cbdd768fced6610f97dd97d8f4490344",
+        "59": "051adfb47d98025680f271519ea973990da86384944135e7b7e0fc3ea7b7ab98",
         "6": "a8e770f914300d621e67fe470600bf63f5368f0c298b0390fd439cd96d0cc247",
+        "60": "934398a5ada88742bb88138f10172a572633015e6b7e50a98b552e530391facd",
+        "61": "b341efea64e30cca391484b5dc7bb56fbafa6f80e78e68125e5948a1ec557bf8",
+        "62": "0f5a55831bc6bf1d68351dde26e3fc89c71f8373dccb2548cbb82e345eed6638",
+        "63": "5bcb6006f685ece4028286cc541d18fc462c5542ff2c921196ff01a0badb564d",
+        "64": "b299a0f957188f62ccfd62e76bf248e85fa2965150b2466d3deb4aee01232950",
+        "65": "269e7872b6e16f1cf085b9dee74f95cc26778095fe8c42441bddb2d255f2a59c",
+        "66": "a52f4a83cc9c694fe4ae38b61e40eb47e37f9f2a70c0b611cf01fc71fc8e690c",
+        "67": "1d1242b9488401bc2ded3f74f32bf55196fce84c9896d686aa285690822f80c9",
+        "68": "0c332a3950d36f84cae5998f47eb555be3a744fe3633788939bc1c2149fad245",
+        "69": "656e405b7ce38bc8687e9bf36fbd2fcd54b8e7f2f8ceaccd5cf4ed0552cab3d9",
         "7": "758aec903672f77f8a5ddc109e3571302baa37470d709972d829547d7eba4783",
         "8": "43299e6f6d1b5c05ea7ab4b81ddaf3693b0e87ce51bedaf0f5d6d70855707566",
         "9": "49bf3ded9d268ab03c2a8d6d371b2e6dca582177e26dfd89a58167626e276594"
@@ -4682,7 +4704,18 @@
         "56": "a80cca61009fe45886f0b1d44616a855da2f1b8b8e32ab6dd86f7e408a892d8b",
         "57": "f3d485d9fbcdbbdb52f10804b59d3dfbc5b2764121ddd9669ad3e127b440b0e2",
         "58": "da7253f146b9cb48484318f596ea49e04f0e7a195a5f8d78ac7cfe13a7750e69",
+        "59": "d07d12e98e56254e4fa45cebd84ed1beea9e048f423e5493eef8d72b06692284",
         "6": "9cfce58150ae462a04619701d8d43d3f2c31e3befadb1203c65d69a639dd948e",
+        "60": "37b7c0399a34acf98b033311d12650e8df355cb7377c63c4f12ba38028118d92",
+        "61": "1b47775f8430d4d1bcace93d9b342c8d590e6b2b2e087eadfa26007c1ebfc104",
+        "62": "437acbeed670fc361f975990ca93452dd15d03aecf24a1a1de089e343c93dbf7",
+        "63": "2a26a23e581316317455d48ba45a08442c71e8c5f9ed1328d9368dd69ca4f7c0",
+        "64": "943a702a2e5246a1bdf3888b697222e5979ceb5fca5825271cb987fa1f3474ca",
+        "65": "2f6df4f5b63d949e9de56b48db838c824b06c405b01dcb032e8113f49a7b4b87",
+        "66": "e913e50e5fcee8ac664b6ca7aa55bfc14ab962e5bae306129efa25f88e73620f",
+        "67": "1ad95e908f7a687230eb3493d9de4b35946e8d8e962ba3bf19c1d0e4e38c8ead",
+        "68": "5a6e10c622e6319a54fb60843e6ea0c368e1980cf983b3625898533826d816a5",
+        "69": "0c057d81fbdbe42fffe35f992118ed732ae2e4848d045e82352ea6980e7441ea",
         "7": "286f9016a8a3230e2e5d927897e81708b099862c61afabc890f50076e4ae0462",
         "8": "3fe933c10995d0eaa679ddb53b935f1f44eb9bb23f7f0fe68ff84a598c557d8c",
         "9": "44dc12f2553925df374df58403abbd0ee6e754c8987323d490199d742334b026"
@@ -6448,7 +6481,18 @@
         "56": "7de0bf5dbe9592f19070788b37399c98c36390020c17d764ccce00d2d6c157a3",
         "57": "1219fa054dfff27259a4852cb0e5729e809283361a52c1bb1df9d6d89967c0bf",
         "58": "c810f1b769742b345cf68b67f33a8478d38446ec76009a85b0474d32f1b4d839",
+        "59": "29ce57c20d39b8911baac4d703e603c11b8812ff0f271b87950e9a1d9046796f",
         "6": "cb2c0b480840f421ebb818e886e506b333f18969d09d722e30e2168d9910cbeb",
+        "60": "653ec23001efc9ca07d1e4a1c31d56af3e9b70d4f581c0d87dae90a47aa0f7ec",
+        "61": "d182544c14d2e39bfc8f59c71c08ffa2f2a3207592725a65982380e414cccfff",
+        "62": "ee1746b9cb7f8902bcec5f32bead78c2211dd393d4e64b38784930c49f5f6975",
+        "63": "7d2a9d6868a3db378978b2e868fbf80f2a211569722751918403720f54013049",
+        "64": "002302a16e49cbef2a9af663c2b9507a2f5951fc84b1d3cc9d7942b5d6eb19a0",
+        "65": "9688c58f25c18e14c771acfc4c4dedbda0bb96a17a4109fdb1e19ee0d3ad9803",
+        "66": "16909172085226a2a261bf382f7005de851ece6c8e680b25dec5483536aa6f04",
+        "67": "06d23e7388d84f3da4949c4d0683fe591c3e088a531ea1d83542cdd5ba88aaa6",
+        "68": "7c889fb5fed5784289fc62c0d7062659c669b11d9ac760cb404fb3f9d024a261",
+        "69": "8dff90c6b262a6a2ab6ed9baf51ebaca7123279f0d2788ef0b57819f15e0bf1f",
         "7": "79d7633c087a7eb374095998f7c72ecd22f707930db6343f54bc64e866f722ed",
         "8": "1b7d6305a39a83570adb2c90ba34ea90c090e5c7f9492fd1799886dfd4c6649c",
         "9": "11eac5a6a6b571b6bdc6e4eab286782607a2358ee8e26ba7d8a10f35464d109a"