[CPU] Add conversion for unsupported BF16 ops via target-specific sta…

…ge (#27) * Remove unused code. Signed-off-by: Ilya Enkovich <[email protected]> * Fma is always allowed on CPU. Signed-off-by: Ilya Enkovich <[email protected]> * Add unsupported op conversions for BF16 type. Signed-off-by: Ilya Enkovich <[email protected]> --------- Signed-off-by: Ilya Enkovich <[email protected]>
triton-lang · Aug 28, 2024 · ddbb2c3 · ddbb2c3
1 parent 5ba8381
commit ddbb2c3
Show file tree

Hide file tree

Showing 35 changed files with 458 additions and 721 deletions.
diff --git a/bin/RegisterTritonDialects.h b/bin/RegisterTritonDialects.h
@@ -16,6 +16,7 @@
 #include "triton/Dialect/TritonNvidiaGPU/Transforms/Passes.h"
 
 #include "cpu/include/TritonCPUToLLVM/Passes.h"
+#include "cpu/include/TritonCPUTransforms/Passes.h"
 #include "cpu/include/TritonToTritonCPU/Passes.h"
 #include "nvidia/include/NVGPUToLLVM/Passes.h"
 #include "nvidia/include/TritonNVIDIAGPUToLLVM/Passes.h"
@@ -69,6 +70,7 @@ inline void registerTritonDialects(mlir::DialectRegistry &registry) {
   // CPU passes
   mlir::triton::cpu::registerTritonToTritonCPUPasses();
   mlir::triton::cpu::registerTritonToTritonCPUPipeline();
+  mlir::triton::cpu::registerTritonCPUTransformsPasses();
   mlir::triton::cpu::registerTritonCPUToLLVMPasses();
   mlir::triton::cpu::registerTritonCPUToLLVMPipeline();
 

diff --git a/include/triton/Dialect/TritonCPU/CMakeLists.txt b/include/triton/Dialect/TritonCPU/CMakeLists.txt
@@ -1,2 +1 @@
 add_subdirectory(IR)
-add_subdirectory(Transforms)
diff --git a/include/triton/Dialect/TritonCPU/Transforms/CMakeLists.txt b/include/triton/Dialect/TritonCPU/Transforms/CMakeLists.txt
diff --git a/include/triton/Dialect/TritonCPU/Transforms/Passes.h b/include/triton/Dialect/TritonCPU/Transforms/Passes.h
diff --git a/include/triton/Dialect/TritonCPU/Transforms/Passes.td b/include/triton/Dialect/TritonCPU/Transforms/Passes.td
diff --git a/include/triton/Dialect/TritonCPU/Transforms/TritonCPUConversion.h b/include/triton/Dialect/TritonCPU/Transforms/TritonCPUConversion.h
diff --git a/lib/Conversion/CMakeLists.txt b/lib/Conversion/CMakeLists.txt
@@ -1,5 +1,2 @@
-# TODO(minjang): I will remove these scratches soon.
-# add_subdirectory(TritonToTritonCPU)
 add_subdirectory(TritonToTritonGPU)
-# add_subdirectory(TritonCPUToLLVM)
 add_subdirectory(TritonGPUToLLVM)
diff --git a/lib/Conversion/TritonCPUToLLVM/CMakeLists.txt b/lib/Conversion/TritonCPUToLLVM/CMakeLists.txt
diff --git a/lib/Conversion/TritonCPUToLLVM/CPUTargetInfo.cpp b/lib/Conversion/TritonCPUToLLVM/CPUTargetInfo.cpp
diff --git a/lib/Conversion/TritonCPUToLLVM/ControlFlowOpToLLVM.cpp b/lib/Conversion/TritonCPUToLLVM/ControlFlowOpToLLVM.cpp
diff --git a/lib/Conversion/TritonCPUToLLVM/FuncOpToLLVM.cpp b/lib/Conversion/TritonCPUToLLVM/FuncOpToLLVM.cpp
diff --git a/lib/Conversion/TritonCPUToLLVM/PrintOpToLLVM.cpp b/lib/Conversion/TritonCPUToLLVM/PrintOpToLLVM.cpp
Original file line number	Diff line number	Diff line change
		@@ -1,2 +1 @@
		add_subdirectory(IR)
		add_subdirectory(Transforms)