From 711584506c5472823f690f2d49dcd3b788c9749f Mon Sep 17 00:00:00 2001 From: Varun Thumbe Date: Mon, 31 Aug 2026 00:27:02 -0700 Subject: [PATCH 1/3] cutlass commit update Signed-off-by: Varun Thumbe --- 3rdparty/cutlass | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/3rdparty/cutlass b/3rdparty/cutlass index 57e3cfb47a..cdcf8d86da 160000 --- a/3rdparty/cutlass +++ b/3rdparty/cutlass @@ -1 +1 @@ -Subproject commit 57e3cfb47a2d9e0d46eb6335c3dc411498efa198 +Subproject commit cdcf8d86daa9b417840fd99875a1b1af685d389d From 125674ef7a7e9b15f10908225dda141791988612 Mon Sep 17 00:00:00 2001 From: Varun Thumbe Date: Tue, 1 Sep 2026 07:57:49 +0000 Subject: [PATCH 2/3] change the namespace to different one Signed-off-by: Varun Thumbe --- ...safe_group_row_cast_col_hadamard_transform_cast_fusion.cu | 4 ++-- .../group_hadamard_transform_cast_fusion.cu | 4 ++-- .../group_row_cast_col_hadamard_transform_cast_fusion.cu | 4 ++-- .../hadamard_transform/hadamard_transform_cast_fusion.cu | 4 ++-- .../row_cast_col_hadamard_transform_cast_fusion.cu | 5 +++-- 5 files changed, 11 insertions(+), 10 deletions(-) diff --git a/transformer_engine/common/hadamard_transform/graph_safe_group_row_cast_col_hadamard_transform_cast_fusion.cu b/transformer_engine/common/hadamard_transform/graph_safe_group_row_cast_col_hadamard_transform_cast_fusion.cu index 0f2456c975..7619b1a4c5 100644 --- a/transformer_engine/common/hadamard_transform/graph_safe_group_row_cast_col_hadamard_transform_cast_fusion.cu +++ b/transformer_engine/common/hadamard_transform/graph_safe_group_row_cast_col_hadamard_transform_cast_fusion.cu @@ -602,7 +602,7 @@ __launch_bounds__(512, 1) __global__ static void group_row_col_rht_gemm_device_g mma.accumulate_ = UMMA::ScaleOut::Zero; - tmem_allocator.allocate(TmemAllocator::Sm100TmemCapacityColumns, + tmem_allocator.allocate(cute::TMEM::Sm100TmemCapacityColumns, &shared_storage.tmem_base_ptr); __syncwarp(); tmem_allocation_result_barrier.arrive(); @@ -651,7 +651,7 @@ __launch_bounds__(512, 1) __global__ static void group_row_col_rht_gemm_device_g } while (scheduler.is_valid()); tmem_allocator.release_allocation_lock(); accumulator_pipeline.producer_tail(accumulator_pipe_producer_state); - tmem_allocator.free(tmem_base_ptr, TmemAllocator::Sm100TmemCapacityColumns); + tmem_allocator.free(tmem_base_ptr, cute::TMEM::Sm100TmemCapacityColumns); } } else if (is_sched_warp) { // Scheduler warp manages tile assignment and pipeline progress for warps diff --git a/transformer_engine/common/hadamard_transform/group_hadamard_transform_cast_fusion.cu b/transformer_engine/common/hadamard_transform/group_hadamard_transform_cast_fusion.cu index 4b1435f9eb..b3d9cb9719 100644 --- a/transformer_engine/common/hadamard_transform/group_hadamard_transform_cast_fusion.cu +++ b/transformer_engine/common/hadamard_transform/group_hadamard_transform_cast_fusion.cu @@ -407,7 +407,7 @@ __global__ static void group_rht_gemm_device( } else if (is_mma_warp) { mma.accumulate_ = UMMA::ScaleOut::Zero; - tmem_allocator.allocate(TmemAllocator::Sm100TmemCapacityColumns, + tmem_allocator.allocate(cute::TMEM::Sm100TmemCapacityColumns, &shared_storage.tmem_base_ptr); __syncwarp(); tmem_allocation_result_barrier.arrive(); @@ -453,7 +453,7 @@ __global__ static void group_rht_gemm_device( } while (tile_idx_m < tiles_in_m && tile_idx_n < tiles_in_n); tmem_allocator.release_allocation_lock(); accumulator_pipeline.producer_tail(accumulator_pipe_producer_state); - tmem_allocator.free(tmem_base_ptr, TmemAllocator::Sm100TmemCapacityColumns); + tmem_allocator.free(tmem_base_ptr, cute::TMEM::Sm100TmemCapacityColumns); } else if (is_epilogue_warp) { static constexpr int FragmentSize = 256 / sizeof_bits_v; diff --git a/transformer_engine/common/hadamard_transform/group_row_cast_col_hadamard_transform_cast_fusion.cu b/transformer_engine/common/hadamard_transform/group_row_cast_col_hadamard_transform_cast_fusion.cu index 2e6d383ce1..1f3ec4c538 100644 --- a/transformer_engine/common/hadamard_transform/group_row_cast_col_hadamard_transform_cast_fusion.cu +++ b/transformer_engine/common/hadamard_transform/group_row_cast_col_hadamard_transform_cast_fusion.cu @@ -590,7 +590,7 @@ __launch_bounds__(512, 1) __global__ static void group_row_col_rht_gemm_device( mma.accumulate_ = UMMA::ScaleOut::Zero; - tmem_allocator.allocate(TmemAllocator::Sm100TmemCapacityColumns, + tmem_allocator.allocate(cute::TMEM::Sm100TmemCapacityColumns, &shared_storage.tmem_base_ptr); __syncwarp(); tmem_allocation_result_barrier.arrive(); @@ -639,7 +639,7 @@ __launch_bounds__(512, 1) __global__ static void group_row_col_rht_gemm_device( } while (scheduler.is_valid()); tmem_allocator.release_allocation_lock(); accumulator_pipeline.producer_tail(accumulator_pipe_producer_state); - tmem_allocator.free(tmem_base_ptr, TmemAllocator::Sm100TmemCapacityColumns); + tmem_allocator.free(tmem_base_ptr, cute::TMEM::Sm100TmemCapacityColumns); } } else if (is_sched_warp) { // Scheduler warp manages tile assignment and pipeline progress for warps diff --git a/transformer_engine/common/hadamard_transform/hadamard_transform_cast_fusion.cu b/transformer_engine/common/hadamard_transform/hadamard_transform_cast_fusion.cu index 433da1f0f0..8c88a6ec04 100644 --- a/transformer_engine/common/hadamard_transform/hadamard_transform_cast_fusion.cu +++ b/transformer_engine/common/hadamard_transform/hadamard_transform_cast_fusion.cu @@ -367,7 +367,7 @@ rht_gemm_device(MShape M, NShape N, KShape K, ClusterTileShape cluster_tile, } else if (is_mma_warp) { mma.accumulate_ = UMMA::ScaleOut::Zero; - tmem_allocator.allocate(TmemAllocator::Sm100TmemCapacityColumns, &shared_storage.tmem_base_ptr); + tmem_allocator.allocate(cute::TMEM::Sm100TmemCapacityColumns, &shared_storage.tmem_base_ptr); __syncwarp(); tmem_allocation_result_barrier.arrive(); uint32_t tmem_base_ptr = shared_storage.tmem_base_ptr; @@ -410,7 +410,7 @@ rht_gemm_device(MShape M, NShape N, KShape K, ClusterTileShape cluster_tile, } while (tile_idx_m < tiles_in_m && tile_idx_n < tiles_in_n); tmem_allocator.release_allocation_lock(); accumulator_pipeline.producer_tail(accumulator_pipe_producer_state); - tmem_allocator.free(tmem_base_ptr, TmemAllocator::Sm100TmemCapacityColumns); + tmem_allocator.free(tmem_base_ptr, cute::TMEM::Sm100TmemCapacityColumns); } else if (is_epilogue_warp) { const float global_amax_val = *global_amax; static constexpr int FragmentSize = 256 / sizeof_bits_v; diff --git a/transformer_engine/common/hadamard_transform/row_cast_col_hadamard_transform_cast_fusion.cu b/transformer_engine/common/hadamard_transform/row_cast_col_hadamard_transform_cast_fusion.cu index 8d8ab20165..e107a47113 100644 --- a/transformer_engine/common/hadamard_transform/row_cast_col_hadamard_transform_cast_fusion.cu +++ b/transformer_engine/common/hadamard_transform/row_cast_col_hadamard_transform_cast_fusion.cu @@ -585,7 +585,8 @@ __global__ static void row_col_rht_gemm_device( mma.accumulate_ = UMMA::ScaleOut::Zero; - tmem_allocator.allocate(TmemAllocator::Sm100TmemCapacityColumns, &shared_storage.tmem_base_ptr); + tmem_allocator.allocate(cute::TMEM::Sm100TmemCapacityColumns, + &shared_storage.tmem_base_ptr); __syncwarp(); tmem_allocation_result_barrier.arrive(); uint32_t tmem_base_ptr = shared_storage.tmem_base_ptr; @@ -633,7 +634,7 @@ __global__ static void row_col_rht_gemm_device( } while (scheduler.is_valid()); tmem_allocator.release_allocation_lock(); accumulator_pipeline.producer_tail(accumulator_pipe_producer_state); - tmem_allocator.free(tmem_base_ptr, TmemAllocator::Sm100TmemCapacityColumns); + tmem_allocator.free(tmem_base_ptr, cute::TMEM::Sm100TmemCapacityColumns); } } else if(is_sched_warp) { cutlass::arch::warpgroup_reg_dealloc<32>(); From 9aeae41dd63eb0a24854a5ee38a9e585a8b5e81a Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Tue, 1 Sep 2026 07:58:58 +0000 Subject: [PATCH 3/3] [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --- .../hadamard_transform/group_hadamard_transform_cast_fusion.cu | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/transformer_engine/common/hadamard_transform/group_hadamard_transform_cast_fusion.cu b/transformer_engine/common/hadamard_transform/group_hadamard_transform_cast_fusion.cu index b3d9cb9719..c99796adae 100644 --- a/transformer_engine/common/hadamard_transform/group_hadamard_transform_cast_fusion.cu +++ b/transformer_engine/common/hadamard_transform/group_hadamard_transform_cast_fusion.cu @@ -407,8 +407,7 @@ __global__ static void group_rht_gemm_device( } else if (is_mma_warp) { mma.accumulate_ = UMMA::ScaleOut::Zero; - tmem_allocator.allocate(cute::TMEM::Sm100TmemCapacityColumns, - &shared_storage.tmem_base_ptr); + tmem_allocator.allocate(cute::TMEM::Sm100TmemCapacityColumns, &shared_storage.tmem_base_ptr); __syncwarp(); tmem_allocation_result_barrier.arrive(); uint32_t tmem_base_ptr = shared_storage.tmem_base_ptr;