diff --git a/pretuned_kernels/fused_qk_norm_rope/_helion_aot_fused_qk_norm_rope_cuda_sm100.py b/pretuned_kernels/fused_qk_norm_rope/_helion_aot_fused_qk_norm_rope_cuda_sm100.py
index f045e7a22..13fc63c28 100644
--- a/pretuned_kernels/fused_qk_norm_rope/_helion_aot_fused_qk_norm_rope_cuda_sm100.py
+++ b/pretuned_kernels/fused_qk_norm_rope/_helion_aot_fused_qk_norm_rope_cuda_sm100.py
@@ -96,47 +96,47 @@ def key_fused_qk_norm_rope(*args) -> int:
 def autotune_fused_qk_norm_rope(*args) -> dict:
     """Select the optimal config for the given arguments."""
     _C = [
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', '', 'last', 'first', '', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[0, 1, 2]], 'l2_groupings': [2], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'last', 'first', 'first', ''], 'num_warps': 4, 'num_stages': 3, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'first', '', 'first', 'last', 'first', 'first', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [1], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [1], 'range_warp_specializes': [False], 'range_multi_buffers': [True], 'range_flattens': [True], 'load_eviction_policies': ['last', 'first', '', 'last', 'first', 'last', '', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': ['pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_interleaved', 'num_sm_multiplier': 4, 'maxnreg': 128},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'first', '', 'first', 'last', 'first', 'first', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'first', '', 'first', 'last', 'first', 'first', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [8], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [True], 'range_flattens': [False], 'load_eviction_policies': ['last', 'first', 'last', 'last', 'first', 'last', '', 'last'], 'num_warps': 8, 'num_stages': 2, 'indexing': ['tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 4, 'maxnreg': 128},
-        {'block_sizes': [4], 'loop_orders': [[0, 1, 2]], 'l2_groupings': [2], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', 'last', 'first', 'first', 'first', '', 'first'], 'num_warps': 1, 'num_stages': 3, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'first', '', 'first', 'last', 'first', 'first', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', '', '', '', 'last', 'last', 'first', 'first'], 'num_warps': 4, 'num_stages': 3, 'indexing': ['pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'first', '', 'first', 'last', 'first', 'first', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [8], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'first', 'first', 'first', 'last', 'last', 'last', 'last'], 'num_warps': 8, 'num_stages': 3, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [4], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', 'first', '', 'first', '', '', 'first'], 'num_warps': 2, 'num_stages': 8, 'indexing': ['tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', 'last', 'first', '', 'first', 'first', '', 'first'], 'num_warps': 2, 'num_stages': 4, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', '', '', '', 'last', 'last', 'first', 'first'], 'num_warps': 4, 'num_stages': 3, 'indexing': ['pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [4], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', 'first', '', 'first', '', '', 'first'], 'num_warps': 2, 'num_stages': 8, 'indexing': ['tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', 'last', 'last', '', 'last', 'last', '', 'first'], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [4], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', 'first', '', 'first', '', '', 'first'], 'num_warps': 2, 'num_stages': 8, 'indexing': ['tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [4], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', 'first', '', 'first', '', '', 'first'], 'num_warps': 2, 'num_stages': 8, 'indexing': ['tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'first', '', 'first', 'last', 'first', 'first', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', 'last', 'first', '', 'last', 'first', 'first', 'first'], 'num_warps': 2, 'num_stages': 7, 'indexing': ['tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [4], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', 'first', '', 'first', '', '', 'first'], 'num_warps': 2, 'num_stages': 8, 'indexing': ['tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', '', 'last', 'last', 'first', ''], 'num_warps': 1, 'num_stages': 7, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', '', 'last', 'first', '', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [4], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', 'first', '', 'first', '', '', 'first'], 'num_warps': 2, 'num_stages': 8, 'indexing': ['tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', 'last', 'first', '', 'first', 'first', '', 'first'], 'num_warps': 2, 'num_stages': 4, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', 'last', 'last', '', 'last', 'last', '', 'first'], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', 'last', 'last', 'first', 'first'], 'num_warps': 1, 'num_stages': 1, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'first', 'first', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [4], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['last', '', 'first', 'first', '', 'first', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'first', 'first', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [False], 'range_multi_buffers': [True], 'range_flattens': [False], 'load_eviction_policies': ['', 'last', '', '', 'last', 'last', 'first', ''], 'num_warps': 1, 'num_stages': 7, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_interleaved', 'num_sm_multiplier': 32},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [False], 'range_multi_buffers': [True], 'range_flattens': [False], 'load_eviction_policies': ['', 'last', '', '', 'last', 'last', 'first', ''], 'num_warps': 1, 'num_stages': 7, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_interleaved', 'num_sm_multiplier': 32},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [False], 'range_multi_buffers': [True], 'range_flattens': [False], 'load_eviction_policies': ['', 'first', 'last', '', 'first', '', 'first', ''], 'num_warps': 1, 'num_stages': 8, 'indexing': ['tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_interleaved', 'num_sm_multiplier': 32},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_multi_buffers': [True], 'range_flattens': [None], 'load_eviction_policies': ['', '', '', '', '', '', 'last', 'last'], 'num_warps': 1, 'num_stages': 6, 'indexing': ['pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [1], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [1], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [1], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [1], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [1], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [1], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', 'last', 'last', 'first', 'first'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'first', 'first', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [4], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['last', '', 'first', 'first', '', 'first', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'first', 'first', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'last', 'last', 'last', 'last', 'last', 'last'], 'num_warps': 1, 'num_stages': 1, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [False], 'range_multi_buffers': [True], 'range_flattens': [False], 'load_eviction_policies': ['', 'first', 'last', '', 'first', '', 'first', ''], 'num_warps': 1, 'num_stages': 8, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_interleaved', 'num_sm_multiplier': 32},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [None], 'range_multi_buffers': [True], 'range_flattens': [None], 'load_eviction_policies': ['', '', '', '', '', '', 'last', 'last'], 'num_warps': 1, 'num_stages': 6, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_warp_specializes': [False], 'range_multi_buffers': [None], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', '', '', '', 'last', ''], 'num_warps': 1, 'num_stages': 6, 'indexing': 'pointer', 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
     ]
     return _C[key_fused_qk_norm_rope(*args)]
diff --git a/pretuned_kernels/fused_qk_norm_rope/_helion_aot_fused_qk_norm_rope_cuda_sm90.py b/pretuned_kernels/fused_qk_norm_rope/_helion_aot_fused_qk_norm_rope_cuda_sm90.py
index 450e7427f..c33d2dd3e 100644
--- a/pretuned_kernels/fused_qk_norm_rope/_helion_aot_fused_qk_norm_rope_cuda_sm90.py
+++ b/pretuned_kernels/fused_qk_norm_rope/_helion_aot_fused_qk_norm_rope_cuda_sm90.py
@@ -99,50 +99,50 @@ def key_fused_qk_norm_rope(*args) -> int:
 def autotune_fused_qk_norm_rope(*args) -> dict:
     """Select the optimal config for the given arguments."""
     _C = [
-        {'block_sizes': [1], 'loop_orders': [[1, 0, 2]], 'l2_groupings': [64], 'range_unroll_factors': [4], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [True], 'load_eviction_policies': ['', '', 'first', '', '', 'last', '', 'first'], 'num_warps': 1, 'num_stages': 3, 'indexing': ['tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 2},
-        {'block_sizes': [1], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', '', 'first', 'first', 'last', '', 'last', ''], 'num_warps': 1, 'num_stages': 1, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [2], 'loop_orders': [[2, 0, 1]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'first', 'first', 'last', 'last', 'first'], 'num_warps': 2, 'num_stages': 1, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [1], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', '', '', 'first', 'first', 'last', 'first', 'last'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [1], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', '', 'last', 'first', 'first', 'first', 'last', ''], 'num_warps': 1, 'num_stages': 8, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'first', '', '', 'last', '', '', 'first'], 'num_warps': 1, 'num_stages': 8, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [1], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', '', '', 'last', 'last', '', 'last'], 'num_warps': 1, 'num_stages': 4, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', '', 'last', '', 'first', '', 'first', 'last'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'first', 'last', 'last', 'last', ''], 'num_warps': 1, 'num_stages': 2, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [1], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [0], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'first', 'first', 'last', 'first', 'first', 'last', ''], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer'], 'pid_type': 'flat', 'atomic_indexing': [], 'range_warp_specializes': [], 'range_num_stages': []},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'first', 'last', 'last', 'last', ''], 'num_warps': 1, 'num_stages': 2, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', 'last', 'first', 'last', 'last', 'last', ''], 'num_warps': 1, 'num_stages': 2, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [1], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', 'last', '', '', 'last', 'last', '', 'last'], 'num_warps': 1, 'num_stages': 4, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', '', 'last', 'first', 'last', 'last', 'last', ''], 'num_warps': 1, 'num_stages': 8, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', 'last', '', 'first', 'last', 'last', 'first', 'first'], 'num_warps': 1, 'num_stages': 4, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', '', '', 'first', 'last', '', 'first', 'first'], 'num_warps': 1, 'num_stages': 3, 'indexing': ['pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', '', 'first', 'last', '', 'last', 'last'], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer'], 'pid_type': 'flat', 'atomic_indexing': [], 'range_warp_specializes': [], 'range_num_stages': []},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['last', '', 'last', 'last', 'first', 'last', 'first', ''], 'num_warps': 1, 'num_stages': 2, 'indexing': ['tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', 'last', 'last', 'first', 'first'], 'num_warps': 1, 'num_stages': 1, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [4], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', 'first', 'last', '', 'last', 'last', 'last', 'first'], 'num_warps': 2, 'num_stages': 2, 'indexing': ['pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['first', 'last', 'last', 'last', 'last', 'first', 'first', ''], 'num_warps': 2, 'num_stages': 6, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', 'last', 'last', 'first', 'first'], 'num_warps': 1, 'num_stages': 1, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', '', '', '', 'last', 'last', 'last', ''], 'num_warps': 2, 'num_stages': 5, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', '', '', 'first', '', 'first', 'first'], 'num_warps': 2, 'num_stages': 2, 'indexing': ['tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', '', '', '', 'last', 'last', 'last', ''], 'num_warps': 2, 'num_stages': 5, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'flat'},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'first', 'first', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'first', 'first', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'first', 'first', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'first', 'first', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [4], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [True], 'load_eviction_policies': ['first', '', 'first', 'first', 'last', 'last', '', ''], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_interleaved', 'num_sm_multiplier': 64, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'first', 'first', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', '', 'last', 'first', 'first', 'last', 'first'], 'num_warps': 1, 'num_stages': 5, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [0], 'range_multi_buffers': [None], 'range_flattens': [None], 'load_eviction_policies': ['', 'last', 'first', 'last', '', 'last', '', 'first'], 'num_warps': 1, 'num_stages': 6, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer'], 'pid_type': 'persistent_interleaved', 'num_sm_multiplier': 128, 'maxnreg': 64, 'atomic_indexing': [], 'range_warp_specializes': []},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', 'last', '', 'last', 'first', ''], 'num_warps': 1, 'num_stages': 4, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 64, 'maxnreg': 256},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [16], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [True], 'load_eviction_policies': ['', 'first', 'first', 'last', '', '', 'last', 'first'], 'num_warps': 1, 'num_stages': 1, 'indexing': ['pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 64, 'maxnreg': 128},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [16], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [True], 'load_eviction_policies': ['', 'first', 'first', 'last', '', '', 'last', 'first'], 'num_warps': 1, 'num_stages': 1, 'indexing': ['pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 64, 'maxnreg': 128},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [False], 'load_eviction_policies': ['', '', '', 'last', '', 'last', 'first', ''], 'num_warps': 1, 'num_stages': 4, 'indexing': ['pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'tensor_descriptor', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 64, 'maxnreg': 256},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', '', '', '', 'first'], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', '', '', '', 'first'], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', '', '', '', 'first'], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', '', '', '', 'first'], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', '', '', '', 'first'], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_multi_buffers': [None], 'range_flattens': [True], 'load_eviction_policies': ['', 'last', 'last', '', 'last', 'last', '', ''], 'num_warps': 1, 'num_stages': 7, 'indexing': ['pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer', 'pointer'], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64, 'atomic_indexing': [], 'range_warp_specializes': []},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', '', '', '', 'first'], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
-        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'load_eviction_policies': ['', 'first', 'last', '', '', '', '', 'first'], 'num_warps': 1, 'num_stages': 2, 'indexing': ['pointer', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'pointer', 'pointer', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor', 'tensor_descriptor'], 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [1], 'loop_orders': [[1, 0, 2]], 'l2_groupings': [64], 'range_unroll_factors': [4], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [True], 'num_warps': 1, 'num_stages': 3, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 2},
+        {'block_sizes': [1], 'loop_orders': [[0, 2, 1]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 1, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[2, 0, 1]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 2, 'num_stages': 1, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [1], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [1], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 8, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 8, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [1], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 4, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [1], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [0], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'pid_type': 'flat', 'atomic_indexing': [], 'range_warp_specializes': [], 'range_num_stages': []},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [1], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 4, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 8, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 4, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 3, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [32], 'range_unroll_factors': [0], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'pid_type': 'flat', 'atomic_indexing': [], 'range_warp_specializes': [], 'range_num_stages': []},
+        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 1, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [4], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 2, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 2, 'num_stages': 6, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [4], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 1, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 2, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [8], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 2, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [4], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [64], 'range_unroll_factors': [0], 'range_warp_specializes': [], 'range_num_stages': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 2, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'flat'},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [4], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [True], 'num_warps': 1, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'persistent_interleaved', 'num_sm_multiplier': 64, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [64], 'range_unroll_factors': [2], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 5, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 32, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [0], 'range_multi_buffers': [None], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 6, 'pid_type': 'persistent_interleaved', 'num_sm_multiplier': 128, 'maxnreg': 64, 'atomic_indexing': [], 'range_warp_specializes': []},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [False], 'num_warps': 1, 'num_stages': 4, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 64, 'maxnreg': 256},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [16], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [True], 'num_warps': 1, 'num_stages': 1, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 64, 'maxnreg': 128},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [16], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [None], 'range_flattens': [True], 'num_warps': 1, 'num_stages': 1, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 64, 'maxnreg': 128},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [16], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [False], 'num_warps': 1, 'num_stages': 4, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 64, 'maxnreg': 256},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[1, 2, 0]], 'l2_groupings': [1], 'range_unroll_factors': [0], 'range_multi_buffers': [None], 'range_flattens': [True], 'num_warps': 1, 'num_stages': 7, 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64, 'atomic_indexing': [], 'range_warp_specializes': []},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
+        {'block_sizes': [2], 'loop_orders': [[2, 1, 0]], 'l2_groupings': [2], 'range_unroll_factors': [1], 'range_warp_specializes': [], 'range_multi_buffers': [False], 'range_flattens': [None], 'num_warps': 1, 'num_stages': 2, 'atomic_indexing': [], 'pid_type': 'persistent_blocked', 'num_sm_multiplier': 128, 'maxnreg': 64},
     ]
     return _C[key_fused_qk_norm_rope(*args)]
diff --git a/pretuned_kernels/fused_qk_norm_rope/fused_qk_norm_rope.py b/pretuned_kernels/fused_qk_norm_rope/fused_qk_norm_rope.py
index 97fa9b332..b52b5e538 100644
--- a/pretuned_kernels/fused_qk_norm_rope/fused_qk_norm_rope.py
+++ b/pretuned_kernels/fused_qk_norm_rope/fused_qk_norm_rope.py
@@ -86,23 +86,7 @@ def fused_qk_norm_rope(
     for tile_m, tile_gn, tile_n in hl.tile(
         [num_tokens, qk_heads, head_dim], block_size=[1, None, head_dim]
     ):
-        x_blk = qkv[tile_m, tile_gn, tile_n].to(dtype=torch.float32)
-
-        rms = x_blk.pow(2).sum(dim=-1)
-        rms = torch.rsqrt(rms * (1.0 / head_dim) + eps)
-
         use_q_weight = (tile_gn.index < num_heads_q)[None, :, None]
-        w_blk = torch.where(
-            use_q_weight, q_weight[None, None, tile_n], k_weight[None, None, tile_n]
-        )
-
-        x_blk = (x_blk * rms[:, :, None]).to(qkv.dtype) * w_blk
-
-        qkv[tile_m, tile_gn, tile_n] = x_blk
-
-        pos_id = position_ids[tile_m]
-        cos_blk = cos_sin_cache[pos_id, hl.arange(embed_dim)]
-        sin_blk = cos_sin_cache[pos_id, hl.arange(embed_dim) + embed_dim]
 
         if is_neox:
             x1_offset = hl.arange(embed_dim)
@@ -111,8 +95,39 @@ def fused_qk_norm_rope(
             x1_offset = hl.arange(embed_dim) * 2
             x2_offset = x1_offset + 1
 
-        x1_blk = qkv[tile_m, tile_gn, x1_offset]
-        x2_blk = qkv[tile_m, tile_gn, x2_offset]
+        if rotary_dim < head_dim:
+            x_blk = qkv[tile_m, tile_gn, tile_n].to(dtype=torch.float32)
+            rms = torch.rsqrt(x_blk.pow(2).sum(dim=-1) * (1.0 / head_dim) + eps)
+            weight = torch.where(
+                use_q_weight,
+                q_weight[None, None, tile_n],
+                k_weight[None, None, tile_n],
+            )
+            x_blk = (x_blk * rms[:, :, None]).to(qkv.dtype) * weight
+            qkv[tile_m, tile_gn, tile_n] = x_blk
+            x1_blk = qkv[tile_m, tile_gn, x1_offset]
+            x2_blk = qkv[tile_m, tile_gn, x2_offset]
+        else:
+            x1_blk = qkv[tile_m, tile_gn, x1_offset].to(dtype=torch.float32)
+            x2_blk = qkv[tile_m, tile_gn, x2_offset].to(dtype=torch.float32)
+            rms = x1_blk.pow(2).sum(dim=-1) + x2_blk.pow(2).sum(dim=-1)
+            rms = torch.rsqrt(rms * (1.0 / head_dim) + eps)
+            x1_weight = torch.where(
+                use_q_weight,
+                q_weight[None, None, x1_offset],
+                k_weight[None, None, x1_offset],
+            )
+            x2_weight = torch.where(
+                use_q_weight,
+                q_weight[None, None, x2_offset],
+                k_weight[None, None, x2_offset],
+            )
+            x1_blk = (x1_blk * rms[:, :, None]).to(qkv.dtype) * x1_weight
+            x2_blk = (x2_blk * rms[:, :, None]).to(qkv.dtype) * x2_weight
+
+        pos_id = position_ids[tile_m]
+        cos_blk = cos_sin_cache[pos_id, hl.arange(embed_dim)]
+        sin_blk = cos_sin_cache[pos_id, hl.arange(embed_dim) + embed_dim]
 
         o1_blk = x1_blk * cos_blk[:, None, :] - x2_blk * sin_blk[:, None, :]
         o2_blk = x2_blk * cos_blk[:, None, :] + x1_blk * sin_blk[:, None, :]
diff --git a/test/test_pretuned_kernels.py b/test/test_pretuned_kernels.py
index c7b6a8f54..3963624aa 100644
--- a/test/test_pretuned_kernels.py
+++ b/test/test_pretuned_kernels.py
@@ -342,6 +342,7 @@ class ExpectedPerf:
     },
     "fused_qk_norm_rope": {
         "sm90": ExpectedPerf(helion_wins=21, total=21, geomean=7.2, wins_slack=2),
+        "sm100": ExpectedPerf(helion_wins=21, total=21, geomean=1.27, wins_slack=1),
     },
 }