remove comments

microsoft · Jan 12, 2024 · d449971 · d449971
1 parent 1b66da8
commit d449971
Show file tree

Hide file tree

Showing 2 changed files with 14 additions and 21 deletions.
diff --git a/dskernels/ft_gemm/gemm_variants/cutlass_extensions/gemm/kernel/moe_cutlass_kernel.h b/dskernels/ft_gemm/gemm_variants/cutlass_extensions/gemm/kernel/moe_cutlass_kernel.h
@@ -518,10 +518,8 @@ struct MoeFCGemm {
 #elif defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 800) && (__CUDA_ARCH__ < 900)
         static constexpr bool compile_needed = platform::is_same<KernelArch, arch::Sm80>::value;
         KernelRunner<compile_needed>::run_kernel(params, shared_storage);#elif defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
-        // TODO Update the arch to Sm90 once CUTLASS hopper specialisations are available
         static constexpr bool compile_needed = platform::is_same<KernelArch, arch::Sm80>::value;
 #elif defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
-        // TODO Update the arch to Sm90 once CUTLASS hopper specialisations are available
         static constexpr bool compile_needed = platform::is_same<KernelArch, arch::Sm80>::value;
         KernelRunner<compile_needed>::run_kernel(params, shared_storage);
 #else

diff --git a/dskernels/ft_gemm/gemm_variants/moe_gemm/moe_gemm_kernels_template.h b/dskernels/ft_gemm/gemm_variants/moe_gemm/moe_gemm_kernels_template.h
@@ -778,25 +778,20 @@ void MoeGemmRunner<T, V>::dispatch_to_arch<EpilogueTag>(const T*          A,
                                                                                       occupancy);
     }
     else if (sm_ >=90) {
-
-        dispatch_moe_gemm_to_cutlass<T, WeightType, cutlass::arch::Sm80, EpilogueTag>(A, B, weight_scales, biases, C,
-            total_rows_before_expert, total_rows, gemm_n, gemm_k, num_experts, gemm_config, sm_, multi_processor_count_,
-            stream, occupancy);
-        //dispatch_moe_gemm_to_cutlass<T, WeightType, cutlass::arch::Sm90, EpilogueTag>(A,
-        //                                                                              B,
-        //                                                                              weight_scales,
-        //                                                                              biases,
-        //                                                                              C,
-        //                                                                              total_rows_before_expert,
-        //                                                                              total_rows,
-        //                                                                              gemm_n,
-        //                                                                              gemm_k,
-        //                                                                              num_experts,
-        //                                                                              gemm_config,
-        //                                                                              sm_,
-        //                                                                              multi_processor_count_,
-        //                                                                              stream,
-        //                                                                              occupancy);
+        dispatch_moe_gemm_to_cutlass<T, WeightType, cutlass::arch::Sm80, EpilogueTag>(A, 
+                                                                                      B, 
+                                                                                      weight_scales, 
+                                                                                      biases, 
+                                                                                      C,
+                                                                                      total_rows_before_expert, 
+                                                                                      total_rows, gemm_n, 
+                                                                                      gemm_k, 
+                                                                                      num_experts, 
+                                                                                      gemm_config, 
+                                                                                      sm_, 
+                                                                                      multi_processor_count_,
+                                                                                      stream, 
+                                                                                      occupancy);
     }
     else {
         throw std::runtime_error("[FT Error][MoE][GEMM Dispatch] Arch unsupported for MoE GEMM");