From d929350fbbfa16f4c2b4a36cf1ccf90e54553c3d Mon Sep 17 00:00:00 2001 From: Yuanwei Fang Date: Thu, 7 Nov 2024 18:02:03 -0800 Subject: [PATCH 1/5] added basic biolerplate code for proton dialect --- .../proton/dialect/proton/CMakeLists.txt | 7 ++ .../dialect/proton/include/CMakeLists.txt | 1 + .../proton/include/Dialect/CMakeLists.txt | 1 + .../include/Dialect/Proton/CMakeLists.txt | 1 + .../include/Dialect/Proton/IR/CMakeLists.txt | 16 ++++ .../include/Dialect/Proton/IR/Dialect.h | 47 +++++++++++ .../Dialect/Proton/IR/ProtonAttrDefs.td | 34 ++++++++ .../Dialect/Proton/IR/ProtonDialect.td | 41 ++++++++++ .../include/Dialect/Proton/IR/ProtonOps.td | 79 +++++++++++++++++++ .../proton/dialect/proton/lib/CMakeLists.txt | 1 + .../dialect/proton/lib/Dialect/CMakeLists.txt | 1 + .../proton/lib/Dialect/Proton/CMakeLists.txt | 1 + .../lib/Dialect/Proton/IR/CMakeLists.txt | 13 +++ .../proton/lib/Dialect/Proton/IR/Dialect.cpp | 52 ++++++++++++ .../proton/lib/Dialect/Proton/IR/Ops.cpp | 58 ++++++++++++++ .../proton/dialect/proton/triton_proton.cc | 15 ++++ 16 files changed, 368 insertions(+) create mode 100644 third_party/proton/dialect/proton/CMakeLists.txt create mode 100644 third_party/proton/dialect/proton/include/CMakeLists.txt create mode 100644 third_party/proton/dialect/proton/include/Dialect/CMakeLists.txt create mode 100644 third_party/proton/dialect/proton/include/Dialect/Proton/CMakeLists.txt create mode 100644 third_party/proton/dialect/proton/include/Dialect/Proton/IR/CMakeLists.txt create mode 100644 third_party/proton/dialect/proton/include/Dialect/Proton/IR/Dialect.h create mode 100644 third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonAttrDefs.td create mode 100644 third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonDialect.td create mode 100644 third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonOps.td create mode 100644 third_party/proton/dialect/proton/lib/CMakeLists.txt create mode 100644 third_party/proton/dialect/proton/lib/Dialect/CMakeLists.txt create mode 100644 third_party/proton/dialect/proton/lib/Dialect/Proton/CMakeLists.txt create mode 100644 third_party/proton/dialect/proton/lib/Dialect/Proton/IR/CMakeLists.txt create mode 100644 third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Dialect.cpp create mode 100644 third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Ops.cpp create mode 100644 third_party/proton/dialect/proton/triton_proton.cc diff --git a/third_party/proton/dialect/proton/CMakeLists.txt b/third_party/proton/dialect/proton/CMakeLists.txt new file mode 100644 index 000000000000..e8a2e8c1b495 --- /dev/null +++ b/third_party/proton/dialect/proton/CMakeLists.txt @@ -0,0 +1,7 @@ +include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include) +include_directories(${CMAKE_CURRENT_BINARY_DIR}/include) +add_subdirectory(include) +add_subdirectory(lib) +if(TRITON_BUILD_PYTHON_MODULE) + add_triton_plugin(TritonProton ${CMAKE_CURRENT_SOURCE_DIR}/python/triton_proton.cc) +endif() diff --git a/third_party/proton/dialect/proton/include/CMakeLists.txt b/third_party/proton/dialect/proton/include/CMakeLists.txt new file mode 100644 index 000000000000..0ca0f41c5af4 --- /dev/null +++ b/third_party/proton/dialect/proton/include/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(Dialect) diff --git a/third_party/proton/dialect/proton/include/Dialect/CMakeLists.txt b/third_party/proton/dialect/proton/include/Dialect/CMakeLists.txt new file mode 100644 index 000000000000..f18c30ba1a6d --- /dev/null +++ b/third_party/proton/dialect/proton/include/Dialect/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(Proton) diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/CMakeLists.txt b/third_party/proton/dialect/proton/include/Dialect/Proton/CMakeLists.txt new file mode 100644 index 000000000000..f33061b2d87c --- /dev/null +++ b/third_party/proton/dialect/proton/include/Dialect/Proton/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(IR) diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/CMakeLists.txt b/third_party/proton/dialect/proton/include/Dialect/Proton/IR/CMakeLists.txt new file mode 100644 index 000000000000..f25e45686907 --- /dev/null +++ b/third_party/proton/dialect/proton/include/Dialect/Proton/IR/CMakeLists.txt @@ -0,0 +1,16 @@ +set(MLIR_BINARY_DIR ${CMAKE_BINARY_DIR}) + +set(LLVM_TARGET_DEFINITIONS ProtonOps.td) +mlir_tablegen(Dialect.h.inc -gen-dialect-decls -dialect=proton) +mlir_tablegen(Dialect.cpp.inc -gen-dialect-defs -dialect=proton) +mlir_tablegen(OpsConversions.inc -gen-llvmir-conversions) +mlir_tablegen(Ops.h.inc -gen-op-decls) +mlir_tablegen(Ops.cpp.inc -gen-op-defs) +add_mlir_doc(ProtonDialect ProtonDialect dialects/ -gen-dialect-doc) +add_mlir_doc(ProtonOps ProtonOps dialects/ -gen-op-doc) +add_public_tablegen_target(ProtonTableGen) + +set(LLVM_TARGET_DEFINITIONS ProtonAttrDefs.td) +mlir_tablegen(ProtonAttrDefs.h.inc -gen-attrdef-decls) +mlir_tablegen(ProtonAttrDefs.cpp.inc -gen-attrdef-defs) +add_public_tablegen_target(ProtonAttrDefsIncGen) diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/Dialect.h b/third_party/proton/dialect/proton/include/Dialect/Proton/IR/Dialect.h new file mode 100644 index 000000000000..7a1ca249003f --- /dev/null +++ b/third_party/proton/dialect/proton/include/Dialect/Proton/IR/Dialect.h @@ -0,0 +1,47 @@ +/* + * Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights + * reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining + * a copy of this software and associated documentation files + * (the "Software"), to deal in the Software without restriction, + * including without limitation the rights to use, copy, modify, merge, + * publish, distribute, sublicense, and/or sell copies of the Software, + * and to permit persons to whom the Software is furnished to do so, + * subject to the following conditions: + * + * The above copyright notice and this permission notice shall be + * included in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, + * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF + * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. + * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY + * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, + * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE + * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + */ + +#ifndef TRITON_DIALECT_PROTON_IR_DIALECT_H_ +#define TRITON_DIALECT_PROTON_IR_DIALECT_H_ + +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/Dialect/Tensor/IR/Tensor.h" +#include "mlir/IR/BuiltinOps.h" +#include "mlir/IR/Dialect.h" +#include "mlir/IR/PatternMatch.h" +#include "proton/include/Dialect/Proton/IR/Dialect.h.inc" + +#define GET_ATTRDEF_CLASSES +#include "proton/include/Dialect/Proton/IR/ProtonAttrDefs.h.inc" + +#define GET_OP_CLASSES +#include "proton/include/Dialect/Proton/IR/Ops.h.inc" + +namespace mlir { +namespace triton { +namespace proton {} // namespace proton +} // namespace triton +} // namespace mlir + +#endif // TRITON_DIALECT_PROTON_IR_DIALECT_H_ diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonAttrDefs.td b/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonAttrDefs.td new file mode 100644 index 000000000000..8b4a82dcc89b --- /dev/null +++ b/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonAttrDefs.td @@ -0,0 +1,34 @@ +// Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining +// a copy of this software and associated documentation files +// (the "Software"), to deal in the Software without restriction, +// including without limitation the rights to use, copy, modify, merge, +// publish, distribute, sublicense, and/or sell copies of the Software, +// and to permit persons to whom the Software is furnished to do so, +// subject to the following conditions: +// +// The above copyright notice and this permission notice shall be +// included in all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, +// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF +// MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. +// IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY +// CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, +// TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE +// SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + + +#ifndef PROTON_ATTRDEFS +#define PROTON_ATTRDEFS + +include "mlir/IR/AttrTypeBase.td" +include "ProtonDialect.td" + +class Proton_Attr traits = [], + string baseCppClass = "::mlir::Attribute"> + : AttrDef { +} + +#endif // PROTON_ATTRDEFS diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonDialect.td b/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonDialect.td new file mode 100644 index 000000000000..63fe0bb51209 --- /dev/null +++ b/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonDialect.td @@ -0,0 +1,41 @@ +// Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining +// a copy of this software and associated documentation files +// (the "Software"), to deal in the Software without restriction, +// including without limitation the rights to use, copy, modify, merge, +// publish, distribute, sublicense, and/or sell copies of the Software, +// and to permit persons to whom the Software is furnished to do so, +// subject to the following conditions: +// +// The above copyright notice and this permission notice shall be +// included in all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, +// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF +// MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. +// IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY +// CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, +// TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE +// SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + +#ifndef PROTON_DIALECT +#define PROTON_DIALECT + +include "mlir/IR/OpBase.td" + +def Proton_Dialect : Dialect { + let name = "proton"; + let cppNamespace = "::mlir::triton::proton"; + + let description = [{ + Proton Dialect provides core ops for building third-party compiler-based + performance profiling and analysis tools. + }]; + + let dependentDialects = [ + "mlir::LLVM::LLVMDialect" + ]; +} + +#endif diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonOps.td b/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonOps.td new file mode 100644 index 000000000000..7e58e27fb3e9 --- /dev/null +++ b/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonOps.td @@ -0,0 +1,79 @@ +// Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining +// a copy of this software and associated documentation files +// (the "Software"), to deal in the Software without restriction, +// including without limitation the rights to use, copy, modify, merge, +// publish, distribute, sublicense, and/or sell copies of the Software, +// and to permit persons to whom the Software is furnished to do so, +// subject to the following conditions: +// +// The above copyright notice and this permission notice shall be +// included in all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, +// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF +// MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. +// IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY +// CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, +// TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE +// SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + + +#ifndef PROTON_OPS +#define PROTON_OPS + +include "mlir/IR/OpBase.td" +include "triton/Dialect/Triton/IR/TritonAttrDefs.td" +include "mlir/IR/EnumAttr.td" +include "triton/Dialect/Triton/IR/TritonTypes.td" +include "mlir/Dialect/LLVMIR/LLVMOpBase.td" +include "mlir/Interfaces/InferTypeOpInterface.td" +include "mlir/Interfaces/SideEffectInterfaces.td" +include "triton/Dialect/Triton/IR/TritonInterfaces.td" +include "ProtonDialect.td" +include "ProtonAttrDefs.td" + +class TT_Proton_Op traits = []> : + Op { +} + +// Proton profiling metric. +def MetricAttr : I32EnumAttr< + "Metric", "", + [ + I32EnumAttrCase<"CYCLE", 0, "cycle">, + I32EnumAttrCase<"INVALID", 1, "invalid">, + ]> { + let cppNamespace = "::mlir::triton::proton"; +} + +// Proton profiling granularity. +def GranularityAttr : I32EnumAttr< + "Granularity", "", + [ + I32EnumAttrCase<"WARPGROUP", 0, "warpgroup">, + I32EnumAttrCase<"WARP", 1, "warp">, + ]> { + let cppNamespace = "::mlir::triton::proton"; +} + +def TT_RecordOp : TT_Proton_Op<"record", [DeclareOpInterfaceMethods]> { + let summary = "Record a GPU hardware event"; + + let description = [{ + The operator records GPU events from performance counters. + Currently only cycle counter is supported. + }]; + let arguments = ( + ins BoolAttr: $isStart, + I32Attr: $regionId, + DefaultValuedAttr:$metric, + DefaultValuedAttr:$granularity + ); + + let hasCustomAssemblyFormat = 0; + let hasVerifier = 1; +} + +#endif // PROTON_OPS diff --git a/third_party/proton/dialect/proton/lib/CMakeLists.txt b/third_party/proton/dialect/proton/lib/CMakeLists.txt new file mode 100644 index 000000000000..0ca0f41c5af4 --- /dev/null +++ b/third_party/proton/dialect/proton/lib/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(Dialect) diff --git a/third_party/proton/dialect/proton/lib/Dialect/CMakeLists.txt b/third_party/proton/dialect/proton/lib/Dialect/CMakeLists.txt new file mode 100644 index 000000000000..f18c30ba1a6d --- /dev/null +++ b/third_party/proton/dialect/proton/lib/Dialect/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(Proton) diff --git a/third_party/proton/dialect/proton/lib/Dialect/Proton/CMakeLists.txt b/third_party/proton/dialect/proton/lib/Dialect/Proton/CMakeLists.txt new file mode 100644 index 000000000000..f33061b2d87c --- /dev/null +++ b/third_party/proton/dialect/proton/lib/Dialect/Proton/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(IR) diff --git a/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/CMakeLists.txt b/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/CMakeLists.txt new file mode 100644 index 000000000000..5eea5cb3cf9e --- /dev/null +++ b/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/CMakeLists.txt @@ -0,0 +1,13 @@ +add_triton_library(ProtonIR + Dialect.cpp + Ops.cpp + + DEPENDS + ProtonTableGen + ProtonAttrDefsIncGen + + LINK_LIBS PUBLIC + MLIRLLVMDialect + TritonIR + TritonGPUIR +) diff --git a/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Dialect.cpp b/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Dialect.cpp new file mode 100644 index 000000000000..3d6f935cca2b --- /dev/null +++ b/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Dialect.cpp @@ -0,0 +1,52 @@ +/* + * Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights + * reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining + * a copy of this software and associated documentation files + * (the "Software"), to deal in the Software without restriction, + * including without limitation the rights to use, copy, modify, merge, + * publish, distribute, sublicense, and/or sell copies of the Software, + * and to permit persons to whom the Software is furnished to do so, + * subject to the following conditions: + * + * The above copyright notice and this permission notice shall be + * included in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, + * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF + * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. + * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY + * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, + * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE + * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + */ + +#include "mlir/IR/DialectImplementation.h" +#include "mlir/IR/OpImplementation.h" + +// clang-format off +#include "Dialect/Proton/IR/Dialect.h" +#include "Dialect/Proton/IR/Dialect.cpp.inc" +// clang-format on + +using namespace mlir; +using namespace mlir::triton::proton; + +void mlir::triton::proton::ProtonDialect::initialize() { + addAttributes< +#define GET_ATTRDEF_LIST +#include "Dialect/Proton/IR/ProtonAttrDefs.cpp.inc" + >(); + + addOperations< +#define GET_OP_LIST +#include "Dialect/Proton/IR/Ops.cpp.inc" + >(); +} + +#define GET_ATTRDEF_CLASSES +#include "Dialect/Proton/IR/ProtonAttrDefs.cpp.inc" + +#define GET_OP_CLASSES +#include "Dialect/Proton/IR/Ops.cpp.inc" diff --git a/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Ops.cpp b/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Ops.cpp new file mode 100644 index 000000000000..6627eb7f2c8e --- /dev/null +++ b/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Ops.cpp @@ -0,0 +1,58 @@ +/* + * Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights + * reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining + * a copy of this software and associated documentation files + * (the "Software"), to deal in the Software without restriction, + * including without limitation the rights to use, copy, modify, merge, + * publish, distribute, sublicense, and/or sell copies of the Software, + * and to permit persons to whom the Software is furnished to do so, + * subject to the following conditions: + * + * The above copyright notice and this permission notice shall be + * included in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, + * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF + * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. + * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY + * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, + * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE + * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + */ + +#include "mlir/IR/Builders.h" +#include "mlir/IR/BuiltinAttributes.h" +#include "mlir/IR/BuiltinTypes.h" +#include "mlir/IR/OperationSupport.h" +#include "mlir/Interfaces/FunctionImplementation.h" +#include "mlir/Interfaces/FunctionInterfaces.h" +#include "mlir/Support/LLVM.h" +#include "proton/Dialect/Proton/IR/Dialect.h" +#include "triton/Dialect/Triton/IR/Dialect.h" +#include "triton/Dialect/Triton/IR/Types.h" +#include "triton/Dialect/Triton/IR/Utility.h" + +#define GET_OP_CLASSES +#include "proton/Dialect/Proton/IR/Ops.cpp.inc" + +namespace mlir { +namespace triton { +namespace proton { + +// -- RecordOp -- +void RecordOp::getEffects( + SmallVectorImpl> + &effects) { + effects.emplace_back(MemoryEffects::Write::get(), + SideEffects::DefaultResource::get()); + effects.emplace_back(MemoryEffects::Read::get(), + SideEffects::DefaultResource::get()); +} + +LogicalResult RecordOp::verify() { return success(); } + +} // namespace proton +} // namespace triton +} // namespace mlir diff --git a/third_party/proton/dialect/proton/triton_proton.cc b/third_party/proton/dialect/proton/triton_proton.cc new file mode 100644 index 000000000000..573526f11227 --- /dev/null +++ b/third_party/proton/dialect/proton/triton_proton.cc @@ -0,0 +1,15 @@ +#include "Dialect/Proton/IR/Dialect.h" +#include "mlir/Pass/PassManager.h" +#include "passes.h" +#include "llvm/IR/Constants.h" +#include "llvm/Support/raw_ostream.h" +#include +#include +#include + +namespace py = pybind11; + +void init_triton_proton(py::module &&m) { + auto passes = m.def_submodule("passes"); + llvm::outs() << "init_triton_proton!\n"; +} From e7697097eaf836bd546fd677bedd0249f8ffd4ec Mon Sep 17 00:00:00 2001 From: Yuanwei Fang Date: Mon, 11 Nov 2024 14:26:03 -0800 Subject: [PATCH 2/5] integrate proton dialect in the build system --- CMakeLists.txt | 4 + bin/RegisterTritonDialects.h | 2 + test/Proton/ops.mlir | 15 ++++ .../dialect/{proton => }/CMakeLists.txt | 2 +- .../proton/dialect/include/CMakeLists.txt | 1 + .../include => include/proton}/CMakeLists.txt | 0 .../proton}/Dialect/CMakeLists.txt | 0 .../proton}/Dialect/Proton/CMakeLists.txt | 0 .../proton}/Dialect/Proton/IR/CMakeLists.txt | 2 + .../proton}/Dialect/Proton/IR/Dialect.h | 8 +- .../Dialect/Proton/IR/ProtonAttrDefs.td | 0 .../Dialect/Proton/IR/ProtonDialect.td | 0 .../proton}/Dialect/Proton/IR/ProtonOps.td | 7 +- third_party/proton/dialect/lib/CMakeLists.txt | 1 + .../{proton/lib => lib/proton}/CMakeLists.txt | 0 .../lib => lib/proton}/Dialect/CMakeLists.txt | 0 .../proton}/Dialect/Proton/CMakeLists.txt | 0 .../proton}/Dialect/Proton/IR/CMakeLists.txt | 0 .../proton}/Dialect/Proton/IR/Dialect.cpp | 13 ++- .../proton}/Dialect/Proton/IR/Ops.cpp | 89 +++++++++++++++++++ .../proton/dialect/proton/triton_proton.cc | 15 ---- third_party/proton/dialect/triton_proton.cc | 20 +++++ 22 files changed, 147 insertions(+), 32 deletions(-) create mode 100644 test/Proton/ops.mlir rename third_party/proton/dialect/{proton => }/CMakeLists.txt (67%) create mode 100644 third_party/proton/dialect/include/CMakeLists.txt rename third_party/proton/dialect/{proton/include => include/proton}/CMakeLists.txt (100%) rename third_party/proton/dialect/{proton/include => include/proton}/Dialect/CMakeLists.txt (100%) rename third_party/proton/dialect/{proton/include => include/proton}/Dialect/Proton/CMakeLists.txt (100%) rename third_party/proton/dialect/{proton/include => include/proton}/Dialect/Proton/IR/CMakeLists.txt (88%) rename third_party/proton/dialect/{proton/include => include/proton}/Dialect/Proton/IR/Dialect.h (84%) rename third_party/proton/dialect/{proton/include => include/proton}/Dialect/Proton/IR/ProtonAttrDefs.td (100%) rename third_party/proton/dialect/{proton/include => include/proton}/Dialect/Proton/IR/ProtonDialect.td (100%) rename third_party/proton/dialect/{proton/include => include/proton}/Dialect/Proton/IR/ProtonOps.td (91%) create mode 100644 third_party/proton/dialect/lib/CMakeLists.txt rename third_party/proton/dialect/{proton/lib => lib/proton}/CMakeLists.txt (100%) rename third_party/proton/dialect/{proton/lib => lib/proton}/Dialect/CMakeLists.txt (100%) rename third_party/proton/dialect/{proton/lib => lib/proton}/Dialect/Proton/CMakeLists.txt (100%) rename third_party/proton/dialect/{proton/lib => lib/proton}/Dialect/Proton/IR/CMakeLists.txt (100%) rename third_party/proton/dialect/{proton/lib => lib/proton}/Dialect/Proton/IR/Dialect.cpp (84%) rename third_party/proton/dialect/{proton/lib => lib/proton}/Dialect/Proton/IR/Ops.cpp (50%) delete mode 100644 third_party/proton/dialect/proton/triton_proton.cc create mode 100644 third_party/proton/dialect/triton_proton.cc diff --git a/CMakeLists.txt b/CMakeLists.txt index 6ee47748602f..ac08849e1f0c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -199,6 +199,9 @@ if(TRITON_BUILD_PYTHON_MODULE) if (TRITON_BUILD_PROTON) add_subdirectory(third_party/proton) endif() + # We always build proton dialect + list(APPEND TRITON_PLUGIN_NAMES "proton") + add_subdirectory(third_party/proton/dialect) get_property(triton_libs GLOBAL PROPERTY TRITON_LIBS) get_property(triton_plugins GLOBAL PROPERTY TRITON_PLUGINS) @@ -304,6 +307,7 @@ if(NOT TRITON_BUILD_PYTHON_MODULE) foreach(CODEGEN_BACKEND ${TRITON_CODEGEN_BACKENDS}) add_subdirectory(third_party/${CODEGEN_BACKEND}) endforeach() + add_subdirectory(third_party/proton/dialect) endif() add_subdirectory(third_party/f2reduce) diff --git a/bin/RegisterTritonDialects.h b/bin/RegisterTritonDialects.h index c69e46792c3d..1e3400d6bd38 100644 --- a/bin/RegisterTritonDialects.h +++ b/bin/RegisterTritonDialects.h @@ -2,6 +2,7 @@ #include "amd/include/Dialect/TritonAMDGPU/IR/Dialect.h" #include "amd/include/TritonAMDGPUTransforms/Passes.h" #include "third_party/nvidia/include/Dialect/NVGPU/IR/Dialect.h" +#include "third_party/proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h" #include "triton/Dialect/Triton/IR/Dialect.h" #include "triton/Dialect/TritonGPU/IR/Dialect.h" #include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h" @@ -75,5 +76,6 @@ inline void registerTritonDialects(mlir::DialectRegistry ®istry) { mlir::gpu::GPUDialect, mlir::LLVM::LLVMDialect, mlir::NVVM::NVVMDialect, mlir::triton::nvgpu::NVGPUDialect, mlir::triton::amdgpu::TritonAMDGPUDialect, + mlir::triton::proton::ProtonDialect, mlir::ROCDL::ROCDLDialect>(); } diff --git a/test/Proton/ops.mlir b/test/Proton/ops.mlir new file mode 100644 index 000000000000..aeda8336bb43 --- /dev/null +++ b/test/Proton/ops.mlir @@ -0,0 +1,15 @@ +// RUN: triton-opt --split-input-file %s -cse -canonicalize | FileCheck %s + +module { + // CHECK-LABEL: proton_record + tt.func @proton_record() { + // CHECK: proton.record <1, "start", "cycle", "warpgroup"> + // CHECK-NEXT: proton.record <1, "end", "cycle", "warpgroup"> + // CHECK-NEXT: tt.return + proton.record <1, "start", "cycle", "warpgroup"> + proton.record <1, "end", "cycle", "warpgroup"> + tt.return + } +} // end module + +// ----- diff --git a/third_party/proton/dialect/proton/CMakeLists.txt b/third_party/proton/dialect/CMakeLists.txt similarity index 67% rename from third_party/proton/dialect/proton/CMakeLists.txt rename to third_party/proton/dialect/CMakeLists.txt index e8a2e8c1b495..c7b5413a0e15 100644 --- a/third_party/proton/dialect/proton/CMakeLists.txt +++ b/third_party/proton/dialect/CMakeLists.txt @@ -3,5 +3,5 @@ include_directories(${CMAKE_CURRENT_BINARY_DIR}/include) add_subdirectory(include) add_subdirectory(lib) if(TRITON_BUILD_PYTHON_MODULE) - add_triton_plugin(TritonProton ${CMAKE_CURRENT_SOURCE_DIR}/python/triton_proton.cc) + add_triton_plugin(TritonProton ${CMAKE_CURRENT_SOURCE_DIR}/triton_proton.cc LINK_LIBS ProtonIR) endif() diff --git a/third_party/proton/dialect/include/CMakeLists.txt b/third_party/proton/dialect/include/CMakeLists.txt new file mode 100644 index 000000000000..b2488278419a --- /dev/null +++ b/third_party/proton/dialect/include/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(proton) diff --git a/third_party/proton/dialect/proton/include/CMakeLists.txt b/third_party/proton/dialect/include/proton/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/proton/include/CMakeLists.txt rename to third_party/proton/dialect/include/proton/CMakeLists.txt diff --git a/third_party/proton/dialect/proton/include/Dialect/CMakeLists.txt b/third_party/proton/dialect/include/proton/Dialect/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/proton/include/Dialect/CMakeLists.txt rename to third_party/proton/dialect/include/proton/Dialect/CMakeLists.txt diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/CMakeLists.txt b/third_party/proton/dialect/include/proton/Dialect/Proton/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/proton/include/Dialect/Proton/CMakeLists.txt rename to third_party/proton/dialect/include/proton/Dialect/Proton/CMakeLists.txt diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/CMakeLists.txt b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/CMakeLists.txt similarity index 88% rename from third_party/proton/dialect/proton/include/Dialect/Proton/IR/CMakeLists.txt rename to third_party/proton/dialect/include/proton/Dialect/Proton/IR/CMakeLists.txt index f25e45686907..4645b0ebcd5a 100644 --- a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/CMakeLists.txt +++ b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/CMakeLists.txt @@ -6,6 +6,8 @@ mlir_tablegen(Dialect.cpp.inc -gen-dialect-defs -dialect=proton) mlir_tablegen(OpsConversions.inc -gen-llvmir-conversions) mlir_tablegen(Ops.h.inc -gen-op-decls) mlir_tablegen(Ops.cpp.inc -gen-op-defs) +mlir_tablegen(OpsEnums.h.inc -gen-enum-decls) +mlir_tablegen(OpsEnums.cpp.inc -gen-enum-defs) add_mlir_doc(ProtonDialect ProtonDialect dialects/ -gen-dialect-doc) add_mlir_doc(ProtonOps ProtonOps dialects/ -gen-op-doc) add_public_tablegen_target(ProtonTableGen) diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/Dialect.h b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h similarity index 84% rename from third_party/proton/dialect/proton/include/Dialect/Proton/IR/Dialect.h rename to third_party/proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h index 7a1ca249003f..68de4a3f3bbb 100644 --- a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/Dialect.h +++ b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h @@ -26,17 +26,17 @@ #define TRITON_DIALECT_PROTON_IR_DIALECT_H_ #include "mlir/Dialect/LLVMIR/LLVMDialect.h" -#include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/Dialect.h" #include "mlir/IR/PatternMatch.h" -#include "proton/include/Dialect/Proton/IR/Dialect.h.inc" +#include "proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h.inc" +#include "proton/dialect/include/proton/Dialect/Proton/IR/OpsEnums.h.inc" #define GET_ATTRDEF_CLASSES -#include "proton/include/Dialect/Proton/IR/ProtonAttrDefs.h.inc" +#include "proton/dialect/include/proton/Dialect/Proton/IR/ProtonAttrDefs.h.inc" #define GET_OP_CLASSES -#include "proton/include/Dialect/Proton/IR/Ops.h.inc" +#include "proton/dialect/include/proton/Dialect/Proton/IR/Ops.h.inc" namespace mlir { namespace triton { diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonAttrDefs.td b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonAttrDefs.td similarity index 100% rename from third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonAttrDefs.td rename to third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonAttrDefs.td diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonDialect.td b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonDialect.td similarity index 100% rename from third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonDialect.td rename to third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonDialect.td diff --git a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonOps.td b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonOps.td similarity index 91% rename from third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonOps.td rename to third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonOps.td index 7e58e27fb3e9..702258591013 100644 --- a/third_party/proton/dialect/proton/include/Dialect/Proton/IR/ProtonOps.td +++ b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonOps.td @@ -24,7 +24,6 @@ #define PROTON_OPS include "mlir/IR/OpBase.td" -include "triton/Dialect/Triton/IR/TritonAttrDefs.td" include "mlir/IR/EnumAttr.td" include "triton/Dialect/Triton/IR/TritonTypes.td" include "mlir/Dialect/LLVMIR/LLVMOpBase.td" @@ -68,11 +67,11 @@ def TT_RecordOp : TT_Proton_Op<"record", [DeclareOpInterfaceMethods:$metric, - DefaultValuedAttr:$granularity + DefaultValuedAttr:$metric, + DefaultValuedAttr:$granularity ); - let hasCustomAssemblyFormat = 0; + let hasCustomAssemblyFormat = 1; let hasVerifier = 1; } diff --git a/third_party/proton/dialect/lib/CMakeLists.txt b/third_party/proton/dialect/lib/CMakeLists.txt new file mode 100644 index 000000000000..b2488278419a --- /dev/null +++ b/third_party/proton/dialect/lib/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(proton) diff --git a/third_party/proton/dialect/proton/lib/CMakeLists.txt b/third_party/proton/dialect/lib/proton/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/proton/lib/CMakeLists.txt rename to third_party/proton/dialect/lib/proton/CMakeLists.txt diff --git a/third_party/proton/dialect/proton/lib/Dialect/CMakeLists.txt b/third_party/proton/dialect/lib/proton/Dialect/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/proton/lib/Dialect/CMakeLists.txt rename to third_party/proton/dialect/lib/proton/Dialect/CMakeLists.txt diff --git a/third_party/proton/dialect/proton/lib/Dialect/Proton/CMakeLists.txt b/third_party/proton/dialect/lib/proton/Dialect/Proton/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/proton/lib/Dialect/Proton/CMakeLists.txt rename to third_party/proton/dialect/lib/proton/Dialect/Proton/CMakeLists.txt diff --git a/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/CMakeLists.txt b/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/proton/lib/Dialect/Proton/IR/CMakeLists.txt rename to third_party/proton/dialect/lib/proton/Dialect/Proton/IR/CMakeLists.txt diff --git a/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Dialect.cpp b/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Dialect.cpp similarity index 84% rename from third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Dialect.cpp rename to third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Dialect.cpp index 3d6f935cca2b..9bb201851c68 100644 --- a/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Dialect.cpp +++ b/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Dialect.cpp @@ -26,8 +26,8 @@ #include "mlir/IR/OpImplementation.h" // clang-format off -#include "Dialect/Proton/IR/Dialect.h" -#include "Dialect/Proton/IR/Dialect.cpp.inc" +#include "proton/Dialect/Proton/IR/Dialect.h" +#include "proton/Dialect/Proton/IR/Dialect.cpp.inc" // clang-format on using namespace mlir; @@ -36,17 +36,14 @@ using namespace mlir::triton::proton; void mlir::triton::proton::ProtonDialect::initialize() { addAttributes< #define GET_ATTRDEF_LIST -#include "Dialect/Proton/IR/ProtonAttrDefs.cpp.inc" +#include "proton/Dialect/Proton/IR/ProtonAttrDefs.cpp.inc" >(); addOperations< #define GET_OP_LIST -#include "Dialect/Proton/IR/Ops.cpp.inc" +#include "proton/Dialect/Proton/IR/Ops.cpp.inc" >(); } #define GET_ATTRDEF_CLASSES -#include "Dialect/Proton/IR/ProtonAttrDefs.cpp.inc" - -#define GET_OP_CLASSES -#include "Dialect/Proton/IR/Ops.cpp.inc" +#include "proton/Dialect/Proton/IR/ProtonAttrDefs.cpp.inc" diff --git a/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Ops.cpp b/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp similarity index 50% rename from third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Ops.cpp rename to third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp index 6627eb7f2c8e..301dc489851d 100644 --- a/third_party/proton/dialect/proton/lib/Dialect/Proton/IR/Ops.cpp +++ b/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp @@ -36,6 +36,7 @@ #define GET_OP_CLASSES #include "proton/Dialect/Proton/IR/Ops.cpp.inc" +#include "proton/Dialect/Proton/IR/OpsEnums.cpp.inc" namespace mlir { namespace triton { @@ -53,6 +54,94 @@ void RecordOp::getEffects( LogicalResult RecordOp::verify() { return success(); } +ParseResult RecordOp::parse(OpAsmParser &parser, OperationState &result) { + MLIRContext *ctx = parser.getContext(); + int idx = 0; + std::string tag, metric, granularity; + mlir::IntegerAttr rid; + + auto parseAttr = [&]() { + switch (idx++) { + case 0: + return parser.parseAttribute(rid, mlir::IntegerType::get(ctx, 32)); + case 1: + return parser.parseString(&tag); + case 2: + return parser.parseOptionalString(&metric); + case 3: + return parser.parseOptionalString(&granularity); + } + return ParseResult(failure()); + }; + + if (parser.parseLess() || parser.parseCommaSeparatedList(parseAttr) || + parser.parseGreater()) + return failure(); + + bool isStart; + if (tag == "start") + isStart = true; + else if (tag == "end") + isStart = false; + else + return failure(); + result.addAttribute("isStart", BoolAttr::get(ctx, isStart)); + + result.addAttribute("regionId", rid); + + MetricAttr metricAttr; + if (!metric.empty()) { + if (metric == "cycle") + metricAttr = MetricAttr::get(ctx, Metric::CYCLE); + else if (metric == "invalid") + metricAttr = MetricAttr::get(ctx, Metric::INVALID); + else + return failure(); + result.addAttribute("metric", metricAttr); + } + + GranularityAttr granularityAttr; + if (!granularity.empty()) { + if (granularity == "warpgroup") + granularityAttr = + GranularityAttr::get(ctx, Granularity::WARPGROUP); + else if (granularity == "warp") + granularityAttr = + GranularityAttr::get(ctx, Granularity::WARP); + else + return failure(); + result.addAttribute("granularity", granularityAttr); + } + + return success(); +} + +void RecordOp::print(OpAsmPrinter &printer) { + Operation *op = getOperation(); + printer << " <"; + printer << getRegionId() << ", "; + + if (getIsStart()) { + printer << "\"start\", "; + } else { + printer << "\"end\", "; + } + + if (getMetric() == Metric::CYCLE) { + printer << "\"cycle\", "; + } else { + printer << "\"invalid\", "; + } + + if (getGranularity() == Granularity::WARP) { + printer << "\"warp\""; + } else { + printer << "\"warpgroup\""; + } + + printer << ">"; +} + } // namespace proton } // namespace triton } // namespace mlir diff --git a/third_party/proton/dialect/proton/triton_proton.cc b/third_party/proton/dialect/proton/triton_proton.cc deleted file mode 100644 index 573526f11227..000000000000 --- a/third_party/proton/dialect/proton/triton_proton.cc +++ /dev/null @@ -1,15 +0,0 @@ -#include "Dialect/Proton/IR/Dialect.h" -#include "mlir/Pass/PassManager.h" -#include "passes.h" -#include "llvm/IR/Constants.h" -#include "llvm/Support/raw_ostream.h" -#include -#include -#include - -namespace py = pybind11; - -void init_triton_proton(py::module &&m) { - auto passes = m.def_submodule("passes"); - llvm::outs() << "init_triton_proton!\n"; -} diff --git a/third_party/proton/dialect/triton_proton.cc b/third_party/proton/dialect/triton_proton.cc new file mode 100644 index 000000000000..7f084707b7cb --- /dev/null +++ b/third_party/proton/dialect/triton_proton.cc @@ -0,0 +1,20 @@ +#include "proton/Dialect/Proton/IR/Dialect.h" +#include "mlir/Pass/PassManager.h" +#include "passes.h" +#include +#include +#include + +namespace py = pybind11; + +void init_triton_proton(py::module &&m) { + auto passes = m.def_submodule("passes"); + + // load dialects + m.def("load_dialects", [](mlir::MLIRContext &context) { + mlir::DialectRegistry registry; + registry.insert(); + context.appendDialectRegistry(registry); + context.loadAllAvailableDialects(); + }); +} From 9bde345c149d33810f4068677d873c6b29c554c1 Mon Sep 17 00:00:00 2001 From: Yuanwei Fang Date: Mon, 11 Nov 2024 15:08:57 -0800 Subject: [PATCH 3/5] fix pre-commit lint --- bin/RegisterTritonDialects.h | 18 +++++++++--------- .../lib/proton/Dialect/Proton/IR/Ops.cpp | 6 ++---- third_party/proton/dialect/triton_proton.cc | 2 +- 3 files changed, 12 insertions(+), 14 deletions(-) diff --git a/bin/RegisterTritonDialects.h b/bin/RegisterTritonDialects.h index 1e3400d6bd38..b4864fb87563 100644 --- a/bin/RegisterTritonDialects.h +++ b/bin/RegisterTritonDialects.h @@ -69,13 +69,13 @@ inline void registerTritonDialects(mlir::DialectRegistry ®istry) { mlir::triton::registerTritonAMDGPULowerInstructionSchedHints(); // TODO: register Triton & TritonGPU passes - registry.insert(); + registry + .insert(); } diff --git a/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp b/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp index 301dc489851d..53842dc6f4ce 100644 --- a/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp +++ b/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp @@ -103,11 +103,9 @@ ParseResult RecordOp::parse(OpAsmParser &parser, OperationState &result) { GranularityAttr granularityAttr; if (!granularity.empty()) { if (granularity == "warpgroup") - granularityAttr = - GranularityAttr::get(ctx, Granularity::WARPGROUP); + granularityAttr = GranularityAttr::get(ctx, Granularity::WARPGROUP); else if (granularity == "warp") - granularityAttr = - GranularityAttr::get(ctx, Granularity::WARP); + granularityAttr = GranularityAttr::get(ctx, Granularity::WARP); else return failure(); result.addAttribute("granularity", granularityAttr); diff --git a/third_party/proton/dialect/triton_proton.cc b/third_party/proton/dialect/triton_proton.cc index 7f084707b7cb..00f2d21e6b7c 100644 --- a/third_party/proton/dialect/triton_proton.cc +++ b/third_party/proton/dialect/triton_proton.cc @@ -1,6 +1,6 @@ -#include "proton/Dialect/Proton/IR/Dialect.h" #include "mlir/Pass/PassManager.h" #include "passes.h" +#include "proton/Dialect/Proton/IR/Dialect.h" #include #include #include From 5192af2223d25baa3d5b14050cc5027be68804d8 Mon Sep 17 00:00:00 2001 From: Yuanwei Fang Date: Mon, 18 Nov 2024 13:04:13 -0800 Subject: [PATCH 4/5] remove license and revise based on the feedbacks regarding the dialect path and asm format --- bin/RegisterTritonDialects.h | 2 +- test/Proton/ops.mlir | 8 +- .../proton/dialect/include/CMakeLists.txt | 2 +- .../{proton => }/Dialect/CMakeLists.txt | 0 .../Dialect/Proton/CMakeLists.txt | 0 .../Dialect/Proton/IR/CMakeLists.txt | 0 .../include/Dialect/Proton/IR/Dialect.h | 23 +++ .../Dialect/Proton/IR/ProtonAttrDefs.td | 12 ++ .../Dialect/Proton/IR/ProtonDialect.td | 20 +++ .../Dialect/Proton/IR/ProtonOps.td | 41 ++--- .../dialect/include/proton/CMakeLists.txt | 1 - .../proton/Dialect/Proton/IR/Dialect.h | 47 ------ .../Dialect/Proton/IR/ProtonAttrDefs.td | 34 ---- .../proton/Dialect/Proton/IR/ProtonDialect.td | 41 ----- third_party/proton/dialect/lib/CMakeLists.txt | 2 +- .../lib/{proton => }/Dialect/CMakeLists.txt | 0 .../Dialect/Proton/CMakeLists.txt | 0 .../Dialect/Proton/IR/CMakeLists.txt | 0 .../dialect/lib/Dialect/Proton/IR/Dialect.cpp | 25 +++ .../dialect/lib/Dialect/Proton/IR/Ops.cpp | 33 ++++ .../proton/dialect/lib/proton/CMakeLists.txt | 1 - .../lib/proton/Dialect/Proton/IR/Dialect.cpp | 49 ------ .../lib/proton/Dialect/Proton/IR/Ops.cpp | 145 ------------------ third_party/proton/dialect/triton_proton.cc | 2 +- 24 files changed, 135 insertions(+), 353 deletions(-) rename third_party/proton/dialect/include/{proton => }/Dialect/CMakeLists.txt (100%) rename third_party/proton/dialect/include/{proton => }/Dialect/Proton/CMakeLists.txt (100%) rename third_party/proton/dialect/include/{proton => }/Dialect/Proton/IR/CMakeLists.txt (100%) create mode 100644 third_party/proton/dialect/include/Dialect/Proton/IR/Dialect.h create mode 100644 third_party/proton/dialect/include/Dialect/Proton/IR/ProtonAttrDefs.td create mode 100644 third_party/proton/dialect/include/Dialect/Proton/IR/ProtonDialect.td rename third_party/proton/dialect/include/{proton => }/Dialect/Proton/IR/ProtonOps.td (53%) delete mode 100644 third_party/proton/dialect/include/proton/CMakeLists.txt delete mode 100644 third_party/proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h delete mode 100644 third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonAttrDefs.td delete mode 100644 third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonDialect.td rename third_party/proton/dialect/lib/{proton => }/Dialect/CMakeLists.txt (100%) rename third_party/proton/dialect/lib/{proton => }/Dialect/Proton/CMakeLists.txt (100%) rename third_party/proton/dialect/lib/{proton => }/Dialect/Proton/IR/CMakeLists.txt (100%) create mode 100644 third_party/proton/dialect/lib/Dialect/Proton/IR/Dialect.cpp create mode 100644 third_party/proton/dialect/lib/Dialect/Proton/IR/Ops.cpp delete mode 100644 third_party/proton/dialect/lib/proton/CMakeLists.txt delete mode 100644 third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Dialect.cpp delete mode 100644 third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp diff --git a/bin/RegisterTritonDialects.h b/bin/RegisterTritonDialects.h index b4864fb87563..71d75b35dbf0 100644 --- a/bin/RegisterTritonDialects.h +++ b/bin/RegisterTritonDialects.h @@ -2,7 +2,7 @@ #include "amd/include/Dialect/TritonAMDGPU/IR/Dialect.h" #include "amd/include/TritonAMDGPUTransforms/Passes.h" #include "third_party/nvidia/include/Dialect/NVGPU/IR/Dialect.h" -#include "third_party/proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h" +#include "third_party/proton/dialect/include/Dialect/Proton/IR/Dialect.h" #include "triton/Dialect/Triton/IR/Dialect.h" #include "triton/Dialect/TritonGPU/IR/Dialect.h" #include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h" diff --git a/test/Proton/ops.mlir b/test/Proton/ops.mlir index aeda8336bb43..22a17e3f0f58 100644 --- a/test/Proton/ops.mlir +++ b/test/Proton/ops.mlir @@ -3,11 +3,11 @@ module { // CHECK-LABEL: proton_record tt.func @proton_record() { - // CHECK: proton.record <1, "start", "cycle", "warpgroup"> - // CHECK-NEXT: proton.record <1, "end", "cycle", "warpgroup"> + // CHECK: proton.record() {isStart = true, regionId = 1 : i32} + // CHECK-NEXT: proton.record() {isStart = false, regionId = 1 : i32} // CHECK-NEXT: tt.return - proton.record <1, "start", "cycle", "warpgroup"> - proton.record <1, "end", "cycle", "warpgroup"> + proton.record() {isStart = true, regionId = 1 : i32} + proton.record() {isStart = false, regionId = 1 : i32} tt.return } } // end module diff --git a/third_party/proton/dialect/include/CMakeLists.txt b/third_party/proton/dialect/include/CMakeLists.txt index b2488278419a..0ca0f41c5af4 100644 --- a/third_party/proton/dialect/include/CMakeLists.txt +++ b/third_party/proton/dialect/include/CMakeLists.txt @@ -1 +1 @@ -add_subdirectory(proton) +add_subdirectory(Dialect) diff --git a/third_party/proton/dialect/include/proton/Dialect/CMakeLists.txt b/third_party/proton/dialect/include/Dialect/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/include/proton/Dialect/CMakeLists.txt rename to third_party/proton/dialect/include/Dialect/CMakeLists.txt diff --git a/third_party/proton/dialect/include/proton/Dialect/Proton/CMakeLists.txt b/third_party/proton/dialect/include/Dialect/Proton/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/include/proton/Dialect/Proton/CMakeLists.txt rename to third_party/proton/dialect/include/Dialect/Proton/CMakeLists.txt diff --git a/third_party/proton/dialect/include/proton/Dialect/Proton/IR/CMakeLists.txt b/third_party/proton/dialect/include/Dialect/Proton/IR/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/include/proton/Dialect/Proton/IR/CMakeLists.txt rename to third_party/proton/dialect/include/Dialect/Proton/IR/CMakeLists.txt diff --git a/third_party/proton/dialect/include/Dialect/Proton/IR/Dialect.h b/third_party/proton/dialect/include/Dialect/Proton/IR/Dialect.h new file mode 100644 index 000000000000..680a205f08f1 --- /dev/null +++ b/third_party/proton/dialect/include/Dialect/Proton/IR/Dialect.h @@ -0,0 +1,23 @@ +#ifndef TRITON_DIALECT_PROTON_IR_DIALECT_H_ +#define TRITON_DIALECT_PROTON_IR_DIALECT_H_ + +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/IR/BuiltinOps.h" +#include "mlir/IR/Dialect.h" +#include "mlir/IR/PatternMatch.h" +#include "proton/dialect/include/Dialect/Proton/IR/Dialect.h.inc" +#include "proton/dialect/include/Dialect/Proton/IR/OpsEnums.h.inc" + +#define GET_ATTRDEF_CLASSES +#include "proton/dialect/include/Dialect/Proton/IR/ProtonAttrDefs.h.inc" + +#define GET_OP_CLASSES +#include "proton/dialect/include/Dialect/Proton/IR/Ops.h.inc" + +namespace mlir { +namespace triton { +namespace proton {} // namespace proton +} // namespace triton +} // namespace mlir + +#endif // TRITON_DIALECT_PROTON_IR_DIALECT_H_ diff --git a/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonAttrDefs.td b/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonAttrDefs.td new file mode 100644 index 000000000000..d469fbb35f6b --- /dev/null +++ b/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonAttrDefs.td @@ -0,0 +1,12 @@ +#ifndef PROTON_ATTRDEFS +#define PROTON_ATTRDEFS + +include "mlir/IR/AttrTypeBase.td" +include "ProtonDialect.td" + +class Proton_Attr traits = [], + string baseCppClass = "::mlir::Attribute"> + : AttrDef { +} + +#endif // PROTON_ATTRDEFS diff --git a/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonDialect.td b/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonDialect.td new file mode 100644 index 000000000000..71641ca8c2d6 --- /dev/null +++ b/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonDialect.td @@ -0,0 +1,20 @@ +#ifndef PROTON_DIALECT +#define PROTON_DIALECT + +include "mlir/IR/OpBase.td" + +def Proton_Dialect : Dialect { + let name = "proton"; + let cppNamespace = "::mlir::triton::proton"; + + let description = [{ + Proton Dialect provides core ops for building third-party compiler-based + performance profiling and analysis tools. + }]; + + let dependentDialects = [ + "mlir::LLVM::LLVMDialect" + ]; +} + +#endif diff --git a/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonOps.td b/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonOps.td similarity index 53% rename from third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonOps.td rename to third_party/proton/dialect/include/Dialect/Proton/IR/ProtonOps.td index 702258591013..d18a48d5d1a0 100644 --- a/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonOps.td +++ b/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonOps.td @@ -1,25 +1,3 @@ -// Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights reserved. -// -// Permission is hereby granted, free of charge, to any person obtaining -// a copy of this software and associated documentation files -// (the "Software"), to deal in the Software without restriction, -// including without limitation the rights to use, copy, modify, merge, -// publish, distribute, sublicense, and/or sell copies of the Software, -// and to permit persons to whom the Software is furnished to do so, -// subject to the following conditions: -// -// The above copyright notice and this permission notice shall be -// included in all copies or substantial portions of the Software. -// -// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, -// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF -// MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. -// IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY -// CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, -// TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE -// SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. - - #ifndef PROTON_OPS #define PROTON_OPS @@ -42,7 +20,6 @@ def MetricAttr : I32EnumAttr< "Metric", "", [ I32EnumAttrCase<"CYCLE", 0, "cycle">, - I32EnumAttrCase<"INVALID", 1, "invalid">, ]> { let cppNamespace = "::mlir::triton::proton"; } @@ -63,16 +40,26 @@ def TT_RecordOp : TT_Proton_Op<"record", [DeclareOpInterfaceMethods:$regionId, DefaultValuedAttr:$metric, DefaultValuedAttr:$granularity ); - - let hasCustomAssemblyFormat = 1; - let hasVerifier = 1; + let assemblyFormat = " `(` operands `)` attr-dict"; } #endif // PROTON_OPS diff --git a/third_party/proton/dialect/include/proton/CMakeLists.txt b/third_party/proton/dialect/include/proton/CMakeLists.txt deleted file mode 100644 index 0ca0f41c5af4..000000000000 --- a/third_party/proton/dialect/include/proton/CMakeLists.txt +++ /dev/null @@ -1 +0,0 @@ -add_subdirectory(Dialect) diff --git a/third_party/proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h deleted file mode 100644 index 68de4a3f3bbb..000000000000 --- a/third_party/proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h +++ /dev/null @@ -1,47 +0,0 @@ -/* - * Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights - * reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining - * a copy of this software and associated documentation files - * (the "Software"), to deal in the Software without restriction, - * including without limitation the rights to use, copy, modify, merge, - * publish, distribute, sublicense, and/or sell copies of the Software, - * and to permit persons to whom the Software is furnished to do so, - * subject to the following conditions: - * - * The above copyright notice and this permission notice shall be - * included in all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, - * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF - * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. - * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY - * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, - * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE - * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. - */ - -#ifndef TRITON_DIALECT_PROTON_IR_DIALECT_H_ -#define TRITON_DIALECT_PROTON_IR_DIALECT_H_ - -#include "mlir/Dialect/LLVMIR/LLVMDialect.h" -#include "mlir/IR/BuiltinOps.h" -#include "mlir/IR/Dialect.h" -#include "mlir/IR/PatternMatch.h" -#include "proton/dialect/include/proton/Dialect/Proton/IR/Dialect.h.inc" -#include "proton/dialect/include/proton/Dialect/Proton/IR/OpsEnums.h.inc" - -#define GET_ATTRDEF_CLASSES -#include "proton/dialect/include/proton/Dialect/Proton/IR/ProtonAttrDefs.h.inc" - -#define GET_OP_CLASSES -#include "proton/dialect/include/proton/Dialect/Proton/IR/Ops.h.inc" - -namespace mlir { -namespace triton { -namespace proton {} // namespace proton -} // namespace triton -} // namespace mlir - -#endif // TRITON_DIALECT_PROTON_IR_DIALECT_H_ diff --git a/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonAttrDefs.td b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonAttrDefs.td deleted file mode 100644 index 8b4a82dcc89b..000000000000 --- a/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonAttrDefs.td +++ /dev/null @@ -1,34 +0,0 @@ -// Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights reserved. -// -// Permission is hereby granted, free of charge, to any person obtaining -// a copy of this software and associated documentation files -// (the "Software"), to deal in the Software without restriction, -// including without limitation the rights to use, copy, modify, merge, -// publish, distribute, sublicense, and/or sell copies of the Software, -// and to permit persons to whom the Software is furnished to do so, -// subject to the following conditions: -// -// The above copyright notice and this permission notice shall be -// included in all copies or substantial portions of the Software. -// -// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, -// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF -// MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. -// IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY -// CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, -// TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE -// SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. - - -#ifndef PROTON_ATTRDEFS -#define PROTON_ATTRDEFS - -include "mlir/IR/AttrTypeBase.td" -include "ProtonDialect.td" - -class Proton_Attr traits = [], - string baseCppClass = "::mlir::Attribute"> - : AttrDef { -} - -#endif // PROTON_ATTRDEFS diff --git a/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonDialect.td b/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonDialect.td deleted file mode 100644 index 63fe0bb51209..000000000000 --- a/third_party/proton/dialect/include/proton/Dialect/Proton/IR/ProtonDialect.td +++ /dev/null @@ -1,41 +0,0 @@ -// Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights reserved. -// -// Permission is hereby granted, free of charge, to any person obtaining -// a copy of this software and associated documentation files -// (the "Software"), to deal in the Software without restriction, -// including without limitation the rights to use, copy, modify, merge, -// publish, distribute, sublicense, and/or sell copies of the Software, -// and to permit persons to whom the Software is furnished to do so, -// subject to the following conditions: -// -// The above copyright notice and this permission notice shall be -// included in all copies or substantial portions of the Software. -// -// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, -// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF -// MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. -// IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY -// CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, -// TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE -// SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. - -#ifndef PROTON_DIALECT -#define PROTON_DIALECT - -include "mlir/IR/OpBase.td" - -def Proton_Dialect : Dialect { - let name = "proton"; - let cppNamespace = "::mlir::triton::proton"; - - let description = [{ - Proton Dialect provides core ops for building third-party compiler-based - performance profiling and analysis tools. - }]; - - let dependentDialects = [ - "mlir::LLVM::LLVMDialect" - ]; -} - -#endif diff --git a/third_party/proton/dialect/lib/CMakeLists.txt b/third_party/proton/dialect/lib/CMakeLists.txt index b2488278419a..0ca0f41c5af4 100644 --- a/third_party/proton/dialect/lib/CMakeLists.txt +++ b/third_party/proton/dialect/lib/CMakeLists.txt @@ -1 +1 @@ -add_subdirectory(proton) +add_subdirectory(Dialect) diff --git a/third_party/proton/dialect/lib/proton/Dialect/CMakeLists.txt b/third_party/proton/dialect/lib/Dialect/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/lib/proton/Dialect/CMakeLists.txt rename to third_party/proton/dialect/lib/Dialect/CMakeLists.txt diff --git a/third_party/proton/dialect/lib/proton/Dialect/Proton/CMakeLists.txt b/third_party/proton/dialect/lib/Dialect/Proton/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/lib/proton/Dialect/Proton/CMakeLists.txt rename to third_party/proton/dialect/lib/Dialect/Proton/CMakeLists.txt diff --git a/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/CMakeLists.txt b/third_party/proton/dialect/lib/Dialect/Proton/IR/CMakeLists.txt similarity index 100% rename from third_party/proton/dialect/lib/proton/Dialect/Proton/IR/CMakeLists.txt rename to third_party/proton/dialect/lib/Dialect/Proton/IR/CMakeLists.txt diff --git a/third_party/proton/dialect/lib/Dialect/Proton/IR/Dialect.cpp b/third_party/proton/dialect/lib/Dialect/Proton/IR/Dialect.cpp new file mode 100644 index 000000000000..60c2852654db --- /dev/null +++ b/third_party/proton/dialect/lib/Dialect/Proton/IR/Dialect.cpp @@ -0,0 +1,25 @@ +#include "mlir/IR/DialectImplementation.h" +#include "mlir/IR/OpImplementation.h" + +// clang-format off +#include "Dialect/Proton/IR/Dialect.h" +#include "Dialect/Proton/IR/Dialect.cpp.inc" +// clang-format on + +using namespace mlir; +using namespace mlir::triton::proton; + +void mlir::triton::proton::ProtonDialect::initialize() { + addAttributes< +#define GET_ATTRDEF_LIST +#include "Dialect/Proton/IR/ProtonAttrDefs.cpp.inc" + >(); + + addOperations< +#define GET_OP_LIST +#include "Dialect/Proton/IR/Ops.cpp.inc" + >(); +} + +#define GET_ATTRDEF_CLASSES +#include "Dialect/Proton/IR/ProtonAttrDefs.cpp.inc" diff --git a/third_party/proton/dialect/lib/Dialect/Proton/IR/Ops.cpp b/third_party/proton/dialect/lib/Dialect/Proton/IR/Ops.cpp new file mode 100644 index 000000000000..1a0799aea127 --- /dev/null +++ b/third_party/proton/dialect/lib/Dialect/Proton/IR/Ops.cpp @@ -0,0 +1,33 @@ +#include "Dialect/Proton/IR/Dialect.h" +#include "mlir/IR/Builders.h" +#include "mlir/IR/BuiltinAttributes.h" +#include "mlir/IR/BuiltinTypes.h" +#include "mlir/IR/OperationSupport.h" +#include "mlir/Interfaces/FunctionImplementation.h" +#include "mlir/Interfaces/FunctionInterfaces.h" +#include "mlir/Support/LLVM.h" +#include "triton/Dialect/Triton/IR/Dialect.h" +#include "triton/Dialect/Triton/IR/Types.h" +#include "triton/Dialect/Triton/IR/Utility.h" + +#define GET_OP_CLASSES +#include "Dialect/Proton/IR/Ops.cpp.inc" +#include "Dialect/Proton/IR/OpsEnums.cpp.inc" + +namespace mlir { +namespace triton { +namespace proton { + +// -- RecordOp -- +void RecordOp::getEffects( + SmallVectorImpl> + &effects) { + effects.emplace_back(MemoryEffects::Write::get(), + SideEffects::DefaultResource::get()); + effects.emplace_back(MemoryEffects::Read::get(), + SideEffects::DefaultResource::get()); +} + +} // namespace proton +} // namespace triton +} // namespace mlir diff --git a/third_party/proton/dialect/lib/proton/CMakeLists.txt b/third_party/proton/dialect/lib/proton/CMakeLists.txt deleted file mode 100644 index 0ca0f41c5af4..000000000000 --- a/third_party/proton/dialect/lib/proton/CMakeLists.txt +++ /dev/null @@ -1 +0,0 @@ -add_subdirectory(Dialect) diff --git a/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Dialect.cpp b/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Dialect.cpp deleted file mode 100644 index 9bb201851c68..000000000000 --- a/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Dialect.cpp +++ /dev/null @@ -1,49 +0,0 @@ -/* - * Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights - * reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining - * a copy of this software and associated documentation files - * (the "Software"), to deal in the Software without restriction, - * including without limitation the rights to use, copy, modify, merge, - * publish, distribute, sublicense, and/or sell copies of the Software, - * and to permit persons to whom the Software is furnished to do so, - * subject to the following conditions: - * - * The above copyright notice and this permission notice shall be - * included in all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, - * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF - * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. - * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY - * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, - * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE - * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. - */ - -#include "mlir/IR/DialectImplementation.h" -#include "mlir/IR/OpImplementation.h" - -// clang-format off -#include "proton/Dialect/Proton/IR/Dialect.h" -#include "proton/Dialect/Proton/IR/Dialect.cpp.inc" -// clang-format on - -using namespace mlir; -using namespace mlir::triton::proton; - -void mlir::triton::proton::ProtonDialect::initialize() { - addAttributes< -#define GET_ATTRDEF_LIST -#include "proton/Dialect/Proton/IR/ProtonAttrDefs.cpp.inc" - >(); - - addOperations< -#define GET_OP_LIST -#include "proton/Dialect/Proton/IR/Ops.cpp.inc" - >(); -} - -#define GET_ATTRDEF_CLASSES -#include "proton/Dialect/Proton/IR/ProtonAttrDefs.cpp.inc" diff --git a/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp b/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp deleted file mode 100644 index 53842dc6f4ce..000000000000 --- a/third_party/proton/dialect/lib/proton/Dialect/Proton/IR/Ops.cpp +++ /dev/null @@ -1,145 +0,0 @@ -/* - * Copyright (c) 2024 Meta Platforms Corporation & Affiliates. All rights - * reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining - * a copy of this software and associated documentation files - * (the "Software"), to deal in the Software without restriction, - * including without limitation the rights to use, copy, modify, merge, - * publish, distribute, sublicense, and/or sell copies of the Software, - * and to permit persons to whom the Software is furnished to do so, - * subject to the following conditions: - * - * The above copyright notice and this permission notice shall be - * included in all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, - * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF - * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. - * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY - * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, - * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE - * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. - */ - -#include "mlir/IR/Builders.h" -#include "mlir/IR/BuiltinAttributes.h" -#include "mlir/IR/BuiltinTypes.h" -#include "mlir/IR/OperationSupport.h" -#include "mlir/Interfaces/FunctionImplementation.h" -#include "mlir/Interfaces/FunctionInterfaces.h" -#include "mlir/Support/LLVM.h" -#include "proton/Dialect/Proton/IR/Dialect.h" -#include "triton/Dialect/Triton/IR/Dialect.h" -#include "triton/Dialect/Triton/IR/Types.h" -#include "triton/Dialect/Triton/IR/Utility.h" - -#define GET_OP_CLASSES -#include "proton/Dialect/Proton/IR/Ops.cpp.inc" -#include "proton/Dialect/Proton/IR/OpsEnums.cpp.inc" - -namespace mlir { -namespace triton { -namespace proton { - -// -- RecordOp -- -void RecordOp::getEffects( - SmallVectorImpl> - &effects) { - effects.emplace_back(MemoryEffects::Write::get(), - SideEffects::DefaultResource::get()); - effects.emplace_back(MemoryEffects::Read::get(), - SideEffects::DefaultResource::get()); -} - -LogicalResult RecordOp::verify() { return success(); } - -ParseResult RecordOp::parse(OpAsmParser &parser, OperationState &result) { - MLIRContext *ctx = parser.getContext(); - int idx = 0; - std::string tag, metric, granularity; - mlir::IntegerAttr rid; - - auto parseAttr = [&]() { - switch (idx++) { - case 0: - return parser.parseAttribute(rid, mlir::IntegerType::get(ctx, 32)); - case 1: - return parser.parseString(&tag); - case 2: - return parser.parseOptionalString(&metric); - case 3: - return parser.parseOptionalString(&granularity); - } - return ParseResult(failure()); - }; - - if (parser.parseLess() || parser.parseCommaSeparatedList(parseAttr) || - parser.parseGreater()) - return failure(); - - bool isStart; - if (tag == "start") - isStart = true; - else if (tag == "end") - isStart = false; - else - return failure(); - result.addAttribute("isStart", BoolAttr::get(ctx, isStart)); - - result.addAttribute("regionId", rid); - - MetricAttr metricAttr; - if (!metric.empty()) { - if (metric == "cycle") - metricAttr = MetricAttr::get(ctx, Metric::CYCLE); - else if (metric == "invalid") - metricAttr = MetricAttr::get(ctx, Metric::INVALID); - else - return failure(); - result.addAttribute("metric", metricAttr); - } - - GranularityAttr granularityAttr; - if (!granularity.empty()) { - if (granularity == "warpgroup") - granularityAttr = GranularityAttr::get(ctx, Granularity::WARPGROUP); - else if (granularity == "warp") - granularityAttr = GranularityAttr::get(ctx, Granularity::WARP); - else - return failure(); - result.addAttribute("granularity", granularityAttr); - } - - return success(); -} - -void RecordOp::print(OpAsmPrinter &printer) { - Operation *op = getOperation(); - printer << " <"; - printer << getRegionId() << ", "; - - if (getIsStart()) { - printer << "\"start\", "; - } else { - printer << "\"end\", "; - } - - if (getMetric() == Metric::CYCLE) { - printer << "\"cycle\", "; - } else { - printer << "\"invalid\", "; - } - - if (getGranularity() == Granularity::WARP) { - printer << "\"warp\""; - } else { - printer << "\"warpgroup\""; - } - - printer << ">"; -} - -} // namespace proton -} // namespace triton -} // namespace mlir diff --git a/third_party/proton/dialect/triton_proton.cc b/third_party/proton/dialect/triton_proton.cc index 00f2d21e6b7c..8046539794e1 100644 --- a/third_party/proton/dialect/triton_proton.cc +++ b/third_party/proton/dialect/triton_proton.cc @@ -1,6 +1,6 @@ +#include "Dialect/Proton/IR/Dialect.h" #include "mlir/Pass/PassManager.h" #include "passes.h" -#include "proton/Dialect/Proton/IR/Dialect.h" #include #include #include From b7b9319d0cfddbbcf836b230d2caedbfebc075aa Mon Sep 17 00:00:00 2001 From: Yuanwei Fang Date: Wed, 20 Nov 2024 14:43:25 -0800 Subject: [PATCH 5/5] drop the dependency of llvm dialect. Will revisit this when needed. --- .../proton/dialect/include/Dialect/Proton/IR/ProtonDialect.td | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonDialect.td b/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonDialect.td index 71641ca8c2d6..245f2e09a2ec 100644 --- a/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonDialect.td +++ b/third_party/proton/dialect/include/Dialect/Proton/IR/ProtonDialect.td @@ -12,9 +12,7 @@ def Proton_Dialect : Dialect { performance profiling and analysis tools. }]; - let dependentDialects = [ - "mlir::LLVM::LLVMDialect" - ]; + let dependentDialects = []; } #endif