softmax with EVT (draft)

codeplaysoftware · Jan 22, 2025 · a84227b · a84227b
1 parent ea70d14
commit a84227b
Show file tree

Hide file tree

Showing 10 changed files with 995 additions and 24 deletions.
diff --git a/examples/sycl/pvc/CMakeLists.txt b/examples/sycl/pvc/CMakeLists.txt
@@ -42,6 +42,11 @@ cutlass_example_add_executable(
   pvc_gemm_with_epilogue_gelu.cpp
 )
 
+cutlass_example_add_executable(
+  pvc_gemm_with_epilogue_softmax
+  pvc_gemm_with_epilogue_softmax.cpp
+)
+
 cutlass_example_add_executable(
   pvc_gemm_with_epilogue_lincombdeeltact
   pvc_gemm_with_epilogue_lincombdeeltact.cpp

diff --git a/examples/sycl/pvc/flash_attention_v2/online_softmax.hpp b/examples/sycl/pvc/flash_attention_v2/online_softmax.hpp
@@ -37,24 +37,24 @@
 #include <sycl/sycl.hpp>
 #include "cutlass/cutlass.h"
 
-#ifdef __SYCL_DEVICE_ONLY__
-#define SYCL_DEVICE_OCL_FMA(x, y) SYCL_EXTERNAL x
-#else
-#define SYCL_DEVICE_OCL_FMA(x, y) \
-  inline x                        \
-  {                               \
-    assert(false);                \
-    return (y)0;                  \
-  }
-#endif
+// #ifdef __SYCL_DEVICE_ONLY__
+// #define SYCL_DEVICE_OCL_FMA(x, y) SYCL_EXTERNAL x
+// #else
+// #define SYCL_DEVICE_OCL_FMA(x, y) \
+//   inline x                        \
+//   {                               \
+//     assert(false);                \
+//     return (y)0;                  \
+//   }
+// #endif
 
 #define EXP sycl::native::exp
 #define DIV sycl::native::divide
 
-SYCL_DEVICE_OCL_FMA(float sub_group_reduce_add(float i), float);
-SYCL_DEVICE_OCL_FMA(float sub_group_reduce_max(float i), float);
+// SYCL_DEVICE_OCL_FMA(float sub_group_reduce_add(float i), float);
+// SYCL_DEVICE_OCL_FMA(float sub_group_reduce_max(float i), float);
 
-#undef SYCL_DEVICE_OCL_FMA
+// #undef SYCL_DEVICE_OCL_FMA
 
 namespace flash
 {