Merge pull request PaddlePaddle#4 from sunnxinn/fp8_cutlass

fix cublaslt search
YZW-explorer · May 29, 2024 · 26eacf9 · 26eacf9
2 parents 59658cb + ed3f983
commit 26eacf9
Show file tree

Hide file tree

Showing 4 changed files with 4 additions and 4 deletions.
diff --git a/paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/cublaslt_gemm.h b/paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/cublaslt_gemm.h
@@ -361,7 +361,7 @@ class CublasLtAlgoCache {
                     sizeof(custom_option));
                 status = dyl::cublasLtMatmulAlgoConfigSetAttribute(
                     &algo, CUBLASLT_ALGO_CONFIG_CTA_SWIZZLING, &k, sizeof(k));
-                int split_k_val = 0;
+                int split_k_val = 1;
                 int reduction_scheme = CUBLASLT_REDUCTION_SCHEME_NONE;
                 status = dyl::cublasLtMatmulAlgoConfigSetAttribute(
                     &algo,

diff --git a/paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/fp8_fp8_bf16_gemm.cu b/paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/fp8_fp8_bf16_gemm.cu
@@ -14,7 +14,7 @@
 
 #include <iostream>
 
-#include "paddle/phi/kernels/fusion/fp8_gemm_with_cublasLt/cublaslt_gemm.h"
+#include "paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/cublaslt_gemm.h"
 
 #include "paddle/phi/backends/gpu/gpu_context.h"
 #include "paddle/phi/common/memory_utils.h"

diff --git a/paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/fp8_fp8_fp16_gemm.cu b/paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/fp8_fp8_fp16_gemm.cu
@@ -14,7 +14,7 @@
 
 #include <iostream>
 
-#include "paddle/phi/kernels/fusion/fp8_gemm_with_cublasLt/cublaslt_gemm.h"
+#include "paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/cublaslt_gemm.h"
 
 #include "paddle/phi/backends/gpu/gpu_context.h"
 #include "paddle/phi/common/memory_utils.h"

diff --git a/...usion/fp8_gemm/fp8_gemm_with_cutlass/cutlass_kernels/dual_gemm/thread/left_silu_and_mul.h b/...usion/fp8_gemm/fp8_gemm_with_cutlass/cutlass_kernels/dual_gemm/thread/left_silu_and_mul.h
@@ -145,7 +145,7 @@ class LeftSiLUAndMul {
     auto silu_lhs = silu(convert_lhs);
     // return ElementOutput(mul(silu_lhs, convert_rhs));
     auto tmp = mul(silu_lhs, convert_rhs);
-    return compute_to_output(mul(alpha_, tmp));
+    return ElementOutput(mul(alpha_, tmp));
   }
 };