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 d5dbd0bd82..4405ab84ba 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 @@ -746,7 +746,7 @@ __launch_bounds__(512, 1) __global__ static void group_row_col_rht_gemm_device_g cutlass::arch::NamedBarrier::sync(NumEpilogueColQuantThreadCount, cutlass::arch::ReservedNamedBarriers::EpilogueBarrier); - // Aligning with TensorEngine's recipe to generate scale factors // {$nv-internal-release} + // Aligning with TensorEngine's recipe to generate scale factors static constexpr float fp4_max = 6.0f; static constexpr float fp8_max = 448.0f; static constexpr float fp4_max_inv = 1.0f / fp4_max; @@ -1003,7 +1003,7 @@ __launch_bounds__(512, 1) __global__ static void group_row_col_rht_gemm_device_g shape_rep, num_tensors, (scheduler.tile_n_base() * size<1>(epilogue_tiler)) * M, packed_N, M, offsets); float a_global_amax_val = shared_storage.global_a_amax[group_idx]; - // Aligning with TensorEngine's recipe to generate scale factors // {$nv-internal-release} + // Aligning with TensorEngine's recipe to generate scale factors static constexpr float fp4_max = 6.0f; static constexpr float fp8_max = 448.0f; static constexpr float fp4_max_inv = 1.0f / fp4_max; 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 cab0b38589..fbc3d740f5 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 @@ -728,7 +728,7 @@ __launch_bounds__(512, 1) __global__ static void group_row_col_rht_gemm_device( cutlass::arch::NamedBarrier::sync(NumEpilogueColQuantThreadCount, cutlass::arch::ReservedNamedBarriers::EpilogueBarrier); - // Aligning with TensorEngine's recipe to generate scale factors // {$nv-internal-release} + // Aligning with TensorEngine's recipe to generate scale factors static constexpr float fp4_max = 6.0f; static constexpr float fp8_max = 448.0f; static constexpr float fp4_max_inv = 1.0f / fp4_max; @@ -980,7 +980,7 @@ __launch_bounds__(512, 1) __global__ static void group_row_col_rht_gemm_device( int group_idx = GetGroupIdx(&args, scheduler.tile_n_base() * size<1>(epilogue_tiler)); float a_global_amax_val = shared_storage.global_a_amax[group_idx]; - // Aligning with TensorEngine's recipe to generate scale factors // {$nv-internal-release} + // Aligning with TensorEngine's recipe to generate scale factors static constexpr float fp4_max = 6.0f; static constexpr float fp8_max = 448.0f; static constexpr float fp4_max_inv = 1.0f / fp4_max; 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 479922a9bf..9daaadfcf4 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 @@ -709,7 +709,7 @@ __global__ static void row_col_rht_gemm_device( auto thr_t2r = tiled_t2r.get_slice(local_thread_idx); auto thr_r2g = tiled_r2g.get_slice(local_thread_idx); - // Aligning with TensorEngine's recipe to generate scale factors // {$nv-internal-release} + // Aligning with TensorEngine's recipe to generate scale factors static constexpr float fp4_max = 6.0f; static constexpr float fp8_max = 448.0f; float const fp4_max_inv = 1.0f / fp4_max; @@ -905,7 +905,7 @@ __global__ static void row_col_rht_gemm_device( cute::Tensor tQArSFA = make_tensor_like(tQAgSFA(_, _, _, _0{}, _0{})); cute::Tensor tQApSFA = thr_s2r.partition_D(pSFA_mn); - // Aligning with TensorEngine's recipe to generate scale factors // {$nv-internal-release} + // Aligning with TensorEngine's recipe to generate scale factors static constexpr float fp4_max = 6.0f; static constexpr float fp8_max = 448.0f; float const fp4_max_inv = 1.0f / fp4_max;