From 492d525445e3af710574c71e99556de4a38e0313 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Tue, 23 Jul 2024 18:07:27 +0000 Subject: [PATCH 01/19] Add gc-gpu-legalize-module pass --- include/gc/Transforms/Passes.td | 9 +++ lib/gc/Transforms/GPU/CMakeLists.txt | 1 + lib/gc/Transforms/GPU/GPULegalizeModule.cpp | 67 +++++++++++++++++++++ 3 files changed, 77 insertions(+) create mode 100644 lib/gc/Transforms/GPU/GPULegalizeModule.cpp diff --git a/include/gc/Transforms/Passes.td b/include/gc/Transforms/Passes.td index 0d5274eb9..9f06e688a 100644 --- a/include/gc/Transforms/Passes.td +++ b/include/gc/Transforms/Passes.td @@ -61,4 +61,13 @@ def LinalgToXeGPU : Pass<"linalg-to-xegpu", "func::FuncOp"> { } #endif +def GpuLegalizeModule: Pass<"gc-gpu-legalize-module", ""> { + let summary = "Legalizes a GPU module for spirv conversion."; + let description = [{ + Update all the nested gpu modules with an appropriate spirv target + information that is used further down in the pipeline. + }]; + let dependentDialects = ["gpu::GPUDialect", "spirv::SPIRVDialect"]; +} + #endif // GC_DIALECT_GC_PASSES diff --git a/lib/gc/Transforms/GPU/CMakeLists.txt b/lib/gc/Transforms/GPU/CMakeLists.txt index 18a7434e2..240dfb971 100644 --- a/lib/gc/Transforms/GPU/CMakeLists.txt +++ b/lib/gc/Transforms/GPU/CMakeLists.txt @@ -1,5 +1,6 @@ gc_add_mlir_library(GcGpuPasses LinalgToXeGPU.cpp + GPULegalizeModule.cpp DEPENDS GraphCompilerPassIncGen diff --git a/lib/gc/Transforms/GPU/GPULegalizeModule.cpp b/lib/gc/Transforms/GPU/GPULegalizeModule.cpp new file mode 100644 index 000000000..f2e697f82 --- /dev/null +++ b/lib/gc/Transforms/GPU/GPULegalizeModule.cpp @@ -0,0 +1,67 @@ +//===- GPULegalizeModule.cpp - Legalize target for gpu module ---*- C++ -*-===// +// +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "gc/Transforms/Passes.h" + +#include "mlir/Dialect/GPU/IR/GPUDialect.h" +#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h" +#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" + +using namespace mlir; + +namespace mlir { +namespace gc { +#define GEN_PASS_DEF_GPULEGALIZEMODULE +#include "gc/Transforms/Passes.h.inc" +} // namespace gc +} // namespace mlir + +struct GpuLegalizeModule + : public gc::impl::GpuLegalizeModuleBase { + using GpuLegalizeModuleBase::GpuLegalizeModuleBase; + + void runOnOperation() override; +}; + +void GpuLegalizeModule::runOnOperation() { + OpBuilder builder(&getContext()); + using namespace mlir::spirv; + + auto version = Version::V_1_0; + SmallVector capabilities = { + Capability::Addresses, Capability::Int64, Capability::Kernel}; + SmallVector extensions{}; + + auto caps = ArrayRef(capabilities); + auto exts = ArrayRef(extensions); + VerCapExtAttr vce = VerCapExtAttr::get(version, caps, exts, &getContext()); + + auto limits = ResourceLimitsAttr::get( + &getContext(), /*max_compute_shared_memory_size=*/16384, + /*max_compute_workgroup_invocations=*/128, + /*max_compute_workgroup_size=*/builder.getI32ArrayAttr({128, 128, 64}), + /*subgroup_size=*/16, + /*min_subgroup_size=*/std::nullopt, + /*max_subgroup_size=*/std::nullopt, + /*cooperative_matrix_properties_khr=*/ArrayAttr{}, + /*cooperative_matrix_properties_nv=*/ArrayAttr{}); + + auto target = spirv::TargetEnvAttr::get( + vce, limits, ClientAPI::OpenCL, Vendor::Intel, DeviceType::DiscreteGPU, + TargetEnvAttr::kUnknownDeviceID); + + getOperation()->walk([&](gpu::GPUModuleOp gpuModule) { + SmallVector targets; + if (std::optional attrs = gpuModule.getTargets()) + targets.append(attrs->getValue().begin(), attrs->getValue().end()); + targets.push_back(target); + // Remove any duplicate targets. + targets.erase(llvm::unique(targets), targets.end()); + gpuModule.setTargetsAttr(builder.getArrayAttr(targets)); + }); +} From f8b665e0aa867894b403f80cbc11e6977cfc8336 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Wed, 24 Jul 2024 14:31:03 +0000 Subject: [PATCH 02/19] Add gc-gpu-signatures-to-llvm pass --- include/gc/Transforms/Passes.td | 5 + lib/gc/Transforms/GPU/CMakeLists.txt | 1 + .../GPU/ConvertGpuSignaturesToLLVM.cpp | 58 ++++++++ lib/gc/Transforms/GPU/GPUOpsLowering.h | 140 ++++++++++++++++++ src/gc-opt/gc-opt.cpp | 4 +- 5 files changed, 207 insertions(+), 1 deletion(-) create mode 100644 lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp create mode 100644 lib/gc/Transforms/GPU/GPUOpsLowering.h diff --git a/include/gc/Transforms/Passes.td b/include/gc/Transforms/Passes.td index 9f06e688a..0925697b1 100644 --- a/include/gc/Transforms/Passes.td +++ b/include/gc/Transforms/Passes.td @@ -70,4 +70,9 @@ def GpuLegalizeModule: Pass<"gc-gpu-legalize-module", ""> { let dependentDialects = ["gpu::GPUDialect", "spirv::SPIRVDialect"]; } +def ConvertGpuSignaturesToLLVM: Pass<"gc-gpu-signatures-to-llvm", "gpu::GPUModuleOp"> { + let summary = "Legalize GPU kernel signatures for runtime code conversion."; + let dependentDialects = ["gpu::GPUDialect", "memref::MemRefDialect"]; +} + #endif // GC_DIALECT_GC_PASSES diff --git a/lib/gc/Transforms/GPU/CMakeLists.txt b/lib/gc/Transforms/GPU/CMakeLists.txt index 240dfb971..ddca23285 100644 --- a/lib/gc/Transforms/GPU/CMakeLists.txt +++ b/lib/gc/Transforms/GPU/CMakeLists.txt @@ -1,6 +1,7 @@ gc_add_mlir_library(GcGpuPasses LinalgToXeGPU.cpp GPULegalizeModule.cpp + ConvertGpuSignaturesToLLVM.cpp DEPENDS GraphCompilerPassIncGen diff --git a/lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp b/lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp new file mode 100644 index 000000000..c2d129070 --- /dev/null +++ b/lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp @@ -0,0 +1,58 @@ +//===- ConvertGpuSignaturesToLLVM.cpp - Legalize signatures -----*- C++ -*-===// +// +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "gc/Transforms/Passes.h" + +#include "mlir/Conversion/LLVMCommon/ConversionTarget.h" +#include "mlir/Conversion/LLVMCommon/TypeConverter.h" +#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/Dialect/MemRef/IR/MemRef.h" + +// TODO: replace once upstream support signature conversion +#include "GPUOpsLowering.h" + +using namespace mlir; + +namespace mlir { +namespace gc { +#define GEN_PASS_DEF_CONVERTGPUSIGNATURESTOLLVM +#include "gc/Transforms/Passes.h.inc" +} // namespace gc +} // namespace mlir + +struct ConvertGpuSignaturesToLLVM + : public gc::impl::ConvertGpuSignaturesToLLVMBase< + ConvertGpuSignaturesToLLVM> { + using ConvertGpuSignaturesToLLVMBase::ConvertGpuSignaturesToLLVMBase; + + void runOnOperation() override; +}; + +void ConvertGpuSignaturesToLLVM::runOnOperation() { + gpu::GPUModuleOp gpuModule = getOperation(); + + for (auto func : gpuModule.getOps()) { + func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(), + UnitAttr::get(&getContext())); + } + + LLVMTypeConverter converter(gpuModule.getContext()); + RewritePatternSet patterns(gpuModule.getContext()); + LLVMConversionTarget target(getContext()); + + patterns.add(converter); + patterns.add( + converter, 0 /*local*/, 3 /*shared*/, + StringAttr::get(&converter.getContext(), "xe.kernel")); + + if (failed(applyPartialConversion(gpuModule, target, std::move(patterns)))) + signalPassFailure(); +} diff --git a/lib/gc/Transforms/GPU/GPUOpsLowering.h b/lib/gc/Transforms/GPU/GPUOpsLowering.h new file mode 100644 index 000000000..92e69badc --- /dev/null +++ b/lib/gc/Transforms/GPU/GPUOpsLowering.h @@ -0,0 +1,140 @@ +//===- GPUOpsLowering.h - GPU FuncOp / ReturnOp lowering -------*- C++ -*--===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +#ifndef MLIR_CONVERSION_GPUCOMMON_GPUOPSLOWERING_H_ +#define MLIR_CONVERSION_GPUCOMMON_GPUOPSLOWERING_H_ + +#include "mlir/Conversion/LLVMCommon/Pattern.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" + +namespace mlir { + +/// Lowering for gpu.dynamic.shared.memory to LLVM dialect. The pattern first +/// create a 0-sized global array symbol similar as LLVM expects. It constructs +/// a memref descriptor with these values and return it. +struct GPUDynamicSharedMemoryOpLowering + : public ConvertOpToLLVMPattern { + using ConvertOpToLLVMPattern< + gpu::DynamicSharedMemoryOp>::ConvertOpToLLVMPattern; + GPUDynamicSharedMemoryOpLowering(const LLVMTypeConverter &converter, + unsigned alignmentBit = 0) + : ConvertOpToLLVMPattern(converter), + alignmentBit(alignmentBit) {} + + LogicalResult + matchAndRewrite(gpu::DynamicSharedMemoryOp op, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override; + +private: + // Alignment bit + unsigned alignmentBit; +}; + +struct GPUFuncOpLowering : ConvertOpToLLVMPattern { + GPUFuncOpLowering( + const LLVMTypeConverter &converter, unsigned allocaAddrSpace, + unsigned workgroupAddrSpace, StringAttr kernelAttributeName, + std::optional kernelBlockSizeAttributeName = std::nullopt) + : ConvertOpToLLVMPattern(converter), + allocaAddrSpace(allocaAddrSpace), + workgroupAddrSpace(workgroupAddrSpace), + kernelAttributeName(kernelAttributeName), + kernelBlockSizeAttributeName(kernelBlockSizeAttributeName) {} + + LogicalResult + matchAndRewrite(gpu::GPUFuncOp gpuFuncOp, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override; + +private: + /// The address space to use for `alloca`s in private memory. + unsigned allocaAddrSpace; + /// The address space to use declaring workgroup memory. + unsigned workgroupAddrSpace; + + /// The attribute name to use instead of `gpu.kernel`. + StringAttr kernelAttributeName; + + /// The attribute name to to set block size + std::optional kernelBlockSizeAttributeName; +}; + +/// The lowering of gpu.printf to a call to HIP hostcalls +/// +/// Simplifies llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp, as we don't have +/// to deal with %s (even if there were first-class strings in MLIR, they're not +/// legal input to gpu.printf) or non-constant format strings +struct GPUPrintfOpToHIPLowering : public ConvertOpToLLVMPattern { + using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern; + + LogicalResult + matchAndRewrite(gpu::PrintfOp gpuPrintfOp, gpu::PrintfOpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override; +}; + +/// The lowering of gpu.printf to a call to an external printf() function +/// +/// This pass will add a declaration of printf() to the GPUModule if needed +/// and separate out the format strings into global constants. For some +/// runtimes, such as OpenCL on AMD, this is sufficient setup, as the compiler +/// will lower printf calls to appropriate device-side code +struct GPUPrintfOpToLLVMCallLowering + : public ConvertOpToLLVMPattern { + GPUPrintfOpToLLVMCallLowering(const LLVMTypeConverter &converter, + int addressSpace = 0) + : ConvertOpToLLVMPattern(converter), + addressSpace(addressSpace) {} + + LogicalResult + matchAndRewrite(gpu::PrintfOp gpuPrintfOp, gpu::PrintfOpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override; + +private: + int addressSpace; +}; + +/// Lowering of gpu.printf to a vprintf standard library. +struct GPUPrintfOpToVPrintfLowering + : public ConvertOpToLLVMPattern { + using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern; + + LogicalResult + matchAndRewrite(gpu::PrintfOp gpuPrintfOp, gpu::PrintfOpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override; +}; + +struct GPUReturnOpLowering : public ConvertOpToLLVMPattern { + using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern; + + LogicalResult + matchAndRewrite(gpu::ReturnOp op, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override; +}; + +namespace impl { +/// Unrolls op if it's operating on vectors. +LogicalResult scalarizeVectorOp(Operation *op, ValueRange operands, + ConversionPatternRewriter &rewriter, + const LLVMTypeConverter &converter); +} // namespace impl + +/// Rewriting that unrolls SourceOp to scalars if it's operating on vectors. +template +struct ScalarizeVectorOpLowering : public ConvertOpToLLVMPattern { +public: + using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern; + + LogicalResult + matchAndRewrite(SourceOp op, typename SourceOp::Adaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + return impl::scalarizeVectorOp(op, adaptor.getOperands(), rewriter, + *this->getTypeConverter()); + } +}; +} // namespace mlir + +#endif // MLIR_CONVERSION_GPUCOMMON_GPUOPSLOWERING_H_ diff --git a/src/gc-opt/gc-opt.cpp b/src/gc-opt/gc-opt.cpp index 7526106e4..8cba6ab32 100644 --- a/src/gc-opt/gc-opt.cpp +++ b/src/gc-opt/gc-opt.cpp @@ -25,6 +25,7 @@ #endif #include "gc/Transforms/Passes.h" #include "mlir/InitAllDialects.h" +#include "mlir/InitAllExtensions.h" #include "mlir/InitAllPasses.h" #include "mlir/Tools/mlir-opt/MlirOptMain.h" @@ -59,7 +60,8 @@ int main(int argc, char *argv[]) { registry.insert(); registry.insert(); mlir::registerAllDialects(registry); -#ifdef GC_USE_IMEX + mlir::registerAllExtensions(registry); +#ifdef GC_USE_GPU registry.insert<::imex::xetile::XeTileDialect, ::imex::gpux::GPUXDialect>(); #endif mlir::cpuruntime::registerConvertCPURuntimeToLLVMInterface(registry); From 7f8c61199094af514c4fecdcd856320c8ddbaf6a Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Mon, 29 Jul 2024 10:06:59 +0000 Subject: [PATCH 03/19] Add gen dialect to hold the gen target --- include/gc/Dialect/CMakeLists.txt | 1 + include/gc/Dialect/LLVMIR/CMakeLists.txt | 6 ++ include/gc/Dialect/LLVMIR/GENDialect.h | 22 ++++++ include/gc/Dialect/LLVMIR/GenOps.td | 57 +++++++++++++ include/gc/Target/LLVM/GEN/Target.h | 30 +++++++ include/gc/Transforms/Passes.td | 5 ++ lib/gc/CMakeLists.txt | 1 + lib/gc/Dialect/CMakeLists.txt | 1 + lib/gc/Dialect/LLVMIR/CMakeLists.txt | 22 ++++++ lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp | 43 ++++++++++ lib/gc/Target/LLVM/CMakeLists.txt | 16 ++++ lib/gc/Target/LLVM/GEN/Target.cpp | 79 +++++++++++++++++++ lib/gc/Transforms/GPU/CMakeLists.txt | 3 + lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp | 45 +++++++++++ .../GPU/Pipelines/LLVMSPVToBinaryPipeline.cpp | 0 lib/gc/Transforms/Pipeline.cpp | 15 ++++ src/gc-opt/gc-opt.cpp | 11 ++- 17 files changed, 355 insertions(+), 2 deletions(-) create mode 100644 include/gc/Dialect/LLVMIR/CMakeLists.txt create mode 100644 include/gc/Dialect/LLVMIR/GENDialect.h create mode 100644 include/gc/Dialect/LLVMIR/GenOps.td create mode 100644 include/gc/Target/LLVM/GEN/Target.h create mode 100644 lib/gc/Dialect/LLVMIR/CMakeLists.txt create mode 100644 lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp create mode 100644 lib/gc/Target/LLVM/CMakeLists.txt create mode 100644 lib/gc/Target/LLVM/GEN/Target.cpp create mode 100644 lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp create mode 100644 lib/gc/Transforms/GPU/Pipelines/LLVMSPVToBinaryPipeline.cpp diff --git a/include/gc/Dialect/CMakeLists.txt b/include/gc/Dialect/CMakeLists.txt index db17a6f99..c2fe46c4d 100644 --- a/include/gc/Dialect/CMakeLists.txt +++ b/include/gc/Dialect/CMakeLists.txt @@ -2,3 +2,4 @@ add_subdirectory(CPURuntime) add_subdirectory(OneDNNGraph) add_subdirectory(Microkernel) add_subdirectory(Linalgx) +add_subdirectory(LLVMIR) diff --git a/include/gc/Dialect/LLVMIR/CMakeLists.txt b/include/gc/Dialect/LLVMIR/CMakeLists.txt new file mode 100644 index 000000000..b36730481 --- /dev/null +++ b/include/gc/Dialect/LLVMIR/CMakeLists.txt @@ -0,0 +1,6 @@ +add_mlir_dialect(GenOps gen) +add_mlir_doc(GenOps GENDialect Dialects/ -gen-dialect-doc -dialect=gen) +set(LLVM_TARGET_DEFINITIONS GenOps.td) +mlir_tablegen(GenOpsAttributes.h.inc -gen-attrdef-decls -attrdefs-dialect=gen) +mlir_tablegen(GenOpsAttributes.cpp.inc -gen-attrdef-defs -attrdefs-dialect=gen) +add_public_tablegen_target(MLIRGENConversionsIncGen) \ No newline at end of file diff --git a/include/gc/Dialect/LLVMIR/GENDialect.h b/include/gc/Dialect/LLVMIR/GENDialect.h new file mode 100644 index 000000000..8745a830b --- /dev/null +++ b/include/gc/Dialect/LLVMIR/GENDialect.h @@ -0,0 +1,22 @@ +//===- GENDialect.h - MLIR GEN target definitions ---------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_DIALECT_LLVMIR_GENDIALECT_H_ +#define MLIR_DIALECT_LLVMIR_GENDIALECT_H_ + +#include "mlir/Bytecode/BytecodeOpInterface.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/IR/Dialect.h" +#include "mlir/IR/OpDefinition.h" + +#define GET_ATTRDEF_CLASSES +#include "gc/Dialect/LLVMIR/GenOpsAttributes.h.inc" + +#include "gc/Dialect/LLVMIR/GenOpsDialect.h.inc" + +#endif /* MLIR_DIALECT_LLVMIR_XEDEFS_H_ */ diff --git a/include/gc/Dialect/LLVMIR/GenOps.td b/include/gc/Dialect/LLVMIR/GenOps.td new file mode 100644 index 000000000..cbc071ee1 --- /dev/null +++ b/include/gc/Dialect/LLVMIR/GenOps.td @@ -0,0 +1,57 @@ +#ifndef GENIR_OPS +#define GENIR_OPS + +include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td" +include "mlir/Dialect/LLVMIR/LLVMOpBase.td" +include "mlir/Dialect/SPIRV/IR/SPIRVBase.td" +include "mlir/Interfaces/SideEffectInterfaces.td" + +def GEN_Dialect : Dialect { + let name = "gen"; + let cppNamespace = "::mlir::gen"; + let dependentDialects = ["LLVM::LLVMDialect"]; + let hasOperationAttrVerify = 1; + + let extraClassDeclaration = [{ + }]; + + let useDefaultAttributePrinterParser = 1; +} + +class GEN_Attr traits = []> + : AttrDef { + let mnemonic = attrMnemonic; +} + +def GEN_TargettAttr : GEN_Attr<"GenTarget", "target"> { + let description = [{ + GPU target attribute for controlling compilation of targets. All + parameters decay into default values if not present. + + Examples: + + 1. Target with default values. + ``` + gpu.module @mymodule [#gen.target] attributes {...} { + ... + } + ``` + }]; + let parameters = (ins + DefaultValuedParameter<"int", "2", "Optimization level to apply.">:$O, + StringRefParameter<"Target triple.", "\"spir64-unknown-unknown\"">:$triple + ); + let assemblyFormat = [{ + (`<` struct($O, $triple)^ `>`)? + }]; + let builders = [ + AttrBuilder<(ins CArg<"int", "2">:$optLevel, + CArg<"StringRef", "\"spir64-unknown-unknown\"">:$triple), [{ + return Base::get($_ctxt, optLevel, triple); + }]> + ]; + let skipDefaultBuilders = 1; + let genVerifyDecl = 1; +} + +#endif // GENIR_OPS diff --git a/include/gc/Target/LLVM/GEN/Target.h b/include/gc/Target/LLVM/GEN/Target.h new file mode 100644 index 000000000..6d3438cb7 --- /dev/null +++ b/include/gc/Target/LLVM/GEN/Target.h @@ -0,0 +1,30 @@ +//===- Target.h - MLIR Xe target registration -------------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This provides registration calls for attaching the Gen target interface. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_TARGET_GEN_TARGET_H +#define MLIR_TARGET_GEN_TARGET_H + +namespace mlir { +class DialectRegistry; +class MLIRContext; +namespace gen { +/// Registers the `TargetAttrInterface` for the `#gen.target` attribute in +/// the given registry. +void registerGenTargetInterfaceExternalModels(DialectRegistry ®istry); + +/// Registers the `TargetAttrInterface` for the `#gen.target` attribute in +/// the registry associated with the given context. +void registerGenTargetInterfaceExternalModels(MLIRContext &context); +} // namespace gen +} // namespace mlir + +#endif // MLIR_TARGET_GEN_TARGET_H diff --git a/include/gc/Transforms/Passes.td b/include/gc/Transforms/Passes.td index 0925697b1..8525ad53c 100644 --- a/include/gc/Transforms/Passes.td +++ b/include/gc/Transforms/Passes.td @@ -75,4 +75,9 @@ def ConvertGpuSignaturesToLLVM: Pass<"gc-gpu-signatures-to-llvm", "gpu::GPUModul let dependentDialects = ["gpu::GPUDialect", "memref::MemRefDialect"]; } +def GpuGenAttachTarget: Pass<"gc-attach-gen-target", ""> { + let summary = "Attaches Gen target to a GPU module."; + let dependentDialects = ["gpu::GPUDialect", "gen::GENDialect"]; +} + #endif // GC_DIALECT_GC_PASSES diff --git a/lib/gc/CMakeLists.txt b/lib/gc/CMakeLists.txt index 7e955ffe9..0ccb16e49 100644 --- a/lib/gc/CMakeLists.txt +++ b/lib/gc/CMakeLists.txt @@ -2,3 +2,4 @@ add_subdirectory(CAPI) add_subdirectory(Dialect) add_subdirectory(Transforms) add_subdirectory(ExecutionEngine) +add_subdirectory(Target/LLVM) diff --git a/lib/gc/Dialect/CMakeLists.txt b/lib/gc/Dialect/CMakeLists.txt index fe07dda0d..3d1dea1ab 100644 --- a/lib/gc/Dialect/CMakeLists.txt +++ b/lib/gc/Dialect/CMakeLists.txt @@ -2,3 +2,4 @@ add_subdirectory(CPURuntime) add_subdirectory(Linalgx) add_subdirectory(Microkernel) add_subdirectory(OneDNNGraph) +add_subdirectory(LLVMIR) diff --git a/lib/gc/Dialect/LLVMIR/CMakeLists.txt b/lib/gc/Dialect/LLVMIR/CMakeLists.txt new file mode 100644 index 000000000..9d61678e9 --- /dev/null +++ b/lib/gc/Dialect/LLVMIR/CMakeLists.txt @@ -0,0 +1,22 @@ +add_mlir_dialect_library(MLIRGENDialect + IR/GENDialect.cpp + + ADDITIONAL_HEADER_DIRS + ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/LLVMIR + ${PROJECT_SOURCE_DIR}/include/gc/Dialect/LLVMIR + + DEPENDS +# MLIRGPUCompilationAttrInterfacesIncGen +# MLIRGENOpsIncGen + MLIRGENConversionsIncGen + + LINK_COMPONENTS + AsmParser + Core + + LINK_LIBS PUBLIC + MLIRIR + MLIRLLVMDialect + MLIRSideEffectInterfaces + ) +set_property(GLOBAL APPEND PROPERTY GC_DIALECT_LIBS MLIRGENDialect) diff --git a/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp b/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp new file mode 100644 index 000000000..336f1e978 --- /dev/null +++ b/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp @@ -0,0 +1,43 @@ +#include "gc/Dialect/LLVMIR/GENDialect.h" + +#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/IR/DialectImplementation.h" +#include "llvm/ADT/TypeSwitch.h" + +using namespace mlir; +using namespace gen; + +#include "gc/Dialect/LLVMIR/GenOpsDialect.cpp.inc" + +LogicalResult +GenTargetAttr::verify(function_ref emitError, int O, + StringRef triple) { + if (O < 0 || O > 3) { + emitError() << "The optimization level must be a number between 0 and 3."; + return failure(); + } + if (triple.empty()) { + emitError() << "The target triple cannot be empty."; + return failure(); + } + return success(); +} + +LogicalResult GENDialect::verifyOperationAttribute(Operation *op, + NamedAttribute attr) { + return success(); +} + +void GENDialect::initialize() { + addAttributes< +#define GET_ATTRDEF_LIST +#include "gc/Dialect/LLVMIR/GenOpsAttributes.cpp.inc" + >(); + + allowUnknownOperations(); + declarePromisedInterface(); +} + +#define GET_ATTRDEF_CLASSES +#include "gc/Dialect/LLVMIR/GenOpsAttributes.cpp.inc" diff --git a/lib/gc/Target/LLVM/CMakeLists.txt b/lib/gc/Target/LLVM/CMakeLists.txt new file mode 100644 index 000000000..1ad4b0b40 --- /dev/null +++ b/lib/gc/Target/LLVM/CMakeLists.txt @@ -0,0 +1,16 @@ +add_mlir_dialect_library(MLIRGENTarget + GEN/Target.cpp + + OBJECT + + ADDITIONAL_HEADER_DIRS + ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/LLVMIR + ${PROJECT_SOURCE_DIR}/include/gc/Dialect/LLVMIR + + LINK_LIBS PUBLIC + MLIRIR + MLIRExecutionEngineUtils + MLIRSupport + MLIRGPUDialect + MLIRTargetLLVM + ) \ No newline at end of file diff --git a/lib/gc/Target/LLVM/GEN/Target.cpp b/lib/gc/Target/LLVM/GEN/Target.cpp new file mode 100644 index 000000000..7f15bd8eb --- /dev/null +++ b/lib/gc/Target/LLVM/GEN/Target.cpp @@ -0,0 +1,79 @@ +//===- Target.cpp - MLIR LLVM XE target compilation -------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file defines Xe target related functions including registration +// calls for the `#xe.target` compilation attribute. +// +//===----------------------------------------------------------------------===// + +#include "gc/Target/LLVM/GEN/Target.h" + +#include "gc/Dialect/LLVMIR/GENDialect.h" +#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/IR/ExtensibleDialect.h" + +using namespace mlir; +using namespace mlir::gen; + +namespace { + +// Xe implementation of the gpu:TargetAttrInterface. +class GenTargetAttrImpl + : public gpu::TargetAttrInterface::FallbackModel { +public: + std::optional> + serializeToObject(Attribute attribute, Operation *module, + const gpu::TargetOptions &options) const; + + Attribute createObject(Attribute attribute, + const SmallVector &object, + const gpu::TargetOptions &options) const; +}; +} // namespace + +void mlir::gen::registerGenTargetInterfaceExternalModels( + DialectRegistry ®istry) { + registry.addExtension(+[](MLIRContext *ctx, gen::GENDialect *dialect) { + GenTargetAttr::attachInterface(*ctx); + }); +} + +void mlir::gen::registerGenTargetInterfaceExternalModels(MLIRContext &context) { + DialectRegistry registry; + registerGenTargetInterfaceExternalModels(registry); + context.appendDialectRegistry(registry); +} + +std::optional> +GenTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module, + const gpu::TargetOptions &options) const { + if (!module) + return std::nullopt; + auto gpuMod = dyn_cast(module); + if (!gpuMod) { + module->emitError("expected to be a gpu.module op"); + return std::nullopt; + } + + // todo +} + +Attribute +GenTargetAttrImpl::createObject(Attribute attribute, + const SmallVector &object, + const gpu::TargetOptions &options) const { + gpu::CompilationTarget format = options.getCompilationTarget(); + DictionaryAttr objectProps; + Builder builder(attribute.getContext()); + return builder.getAttr( + attribute, format, + builder.getStringAttr(StringRef(object.data(), object.size())), + objectProps); +} diff --git a/lib/gc/Transforms/GPU/CMakeLists.txt b/lib/gc/Transforms/GPU/CMakeLists.txt index ddca23285..9b3bf5d77 100644 --- a/lib/gc/Transforms/GPU/CMakeLists.txt +++ b/lib/gc/Transforms/GPU/CMakeLists.txt @@ -2,11 +2,14 @@ gc_add_mlir_library(GcGpuPasses LinalgToXeGPU.cpp GPULegalizeModule.cpp ConvertGpuSignaturesToLLVM.cpp + GPUAttachGenTarget.cpp DEPENDS GraphCompilerPassIncGen LINK_LIBS PUBLIC + MLIRGENDialect + MLIRGENTarget MLIRGPUDialect MLIRXeGPUDialect MLIRGPUTransforms diff --git a/lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp b/lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp new file mode 100644 index 000000000..59d0843c4 --- /dev/null +++ b/lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp @@ -0,0 +1,45 @@ +//===- GPUAttachGenTarget.cpp - Attach Gen target to gpu module -*- C++ -*-===// +// +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "gc/Transforms/Passes.h" + +#include "gc/Dialect/LLVMIR/GENDialect.h" +#include "gc/Target/LLVM/GEN/Target.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" + +using namespace mlir; + +namespace mlir { +namespace gc { +#define GEN_PASS_DEF_GPUGENATTACHTARGET +#include "gc/Transforms/Passes.h.inc" +} // namespace gc +} // namespace mlir + +struct GpuGenAttachTarget + : public gc::impl::GpuGenAttachTargetBase { + using GpuGenAttachTargetBase::GpuGenAttachTargetBase; + + void runOnOperation() override; +}; + +void GpuGenAttachTarget::runOnOperation() { + OpBuilder builder(&getContext()); + auto target = + builder.getAttr(2, "spir64-unknown-unknown"); + getOperation()->walk([&](gpu::GPUModuleOp gpuModule) { + SmallVector targets; + if (std::optional attrs = gpuModule.getTargets()) + targets.append(attrs->getValue().begin(), attrs->getValue().end()); + targets.push_back(target); + // Remove any duplicate targets. + targets.erase(llvm::unique(targets), targets.end()); + gpuModule.setTargetsAttr(builder.getArrayAttr(targets)); + }); +} diff --git a/lib/gc/Transforms/GPU/Pipelines/LLVMSPVToBinaryPipeline.cpp b/lib/gc/Transforms/GPU/Pipelines/LLVMSPVToBinaryPipeline.cpp new file mode 100644 index 000000000..e69de29bb diff --git a/lib/gc/Transforms/Pipeline.cpp b/lib/gc/Transforms/Pipeline.cpp index 7d487f149..5d1fed278 100644 --- a/lib/gc/Transforms/Pipeline.cpp +++ b/lib/gc/Transforms/Pipeline.cpp @@ -149,6 +149,21 @@ void populateCPUPipeline(mlir::OpPassManager &pm) { populateLLVMPasses(pm); } +void populateGPUPipeline(mlir::OpPassManager &pm) { + pm.addPass(createLinalgGeneralizeNamedOpsPass()); + bufferization::OneShotBufferizationOptions options; + options.bufferizeFunctionBoundaries = true; + options.setFunctionBoundaryTypeConversion( + bufferization::LayoutMapOption::IdentityLayoutMap); + pm.addPass(bufferization::createOneShotBufferizePass(options)); +} + +void buildGPUBinaryGenerationPipeline(mlir::OpPassManager &pm) { + GpuModuleToBinaryPassOptions gpuModuleToBinaryPassOptions; + gpuModuleToBinaryPassOptions.compilationTarget = "genisa"; + pm.addPass(createGpuModuleToBinaryPass(gpuModuleToBinaryPassOptions)); +} + void registerCPUPipeline() { PassPipelineRegistration<>("gc-cpu-pipeline", "The CPU pipeline for Graph Compiler", diff --git a/src/gc-opt/gc-opt.cpp b/src/gc-opt/gc-opt.cpp index 8cba6ab32..1264fcd40 100644 --- a/src/gc-opt/gc-opt.cpp +++ b/src/gc-opt/gc-opt.cpp @@ -18,11 +18,16 @@ */ #include "gc/Dialect/CPURuntime/Transforms/CPURuntimePasses.h" +#include "gc/Dialect/LLVMIR/GENDialect.h" #include "gc/Dialect/Linalgx/LinalgxDialect.h" #include "gc/Dialect/Microkernel/MicrokernelDialect.h" #ifdef GC_HAS_ONEDNN_DIALECT #include "gc/Dialect/OneDNNGraph/OneDNNGraphDialect.h" +<<<<<<< HEAD #endif + ======= +#include "gc/Target/LLVM/GEN/Target.h" + >>>>>>> abeedf6 (Add gen dialect to hold the gen target) #include "gc/Transforms/Passes.h" #include "mlir/InitAllDialects.h" #include "mlir/InitAllExtensions.h" @@ -34,8 +39,8 @@ #include #endif -namespace mlir::gc { -void registerCPUPipeline(); + namespace mlir::gc { + void registerCPUPipeline(); } // namespace mlir::gc int main(int argc, char *argv[]) { @@ -59,8 +64,10 @@ int main(int argc, char *argv[]) { registry.insert(); registry.insert(); registry.insert(); + registry.insert(); mlir::registerAllDialects(registry); mlir::registerAllExtensions(registry); + mlir::gen::registerGenTargetInterfaceExternalModels(registry); #ifdef GC_USE_GPU registry.insert<::imex::xetile::XeTileDialect, ::imex::gpux::GPUXDialect>(); #endif From fa3ba4b571b896d62f39699b32b814c1de7c294e Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Tue, 30 Jul 2024 10:17:51 +0000 Subject: [PATCH 04/19] Add gen target and a lowering pipeline through gpu-module-to-binary --- include/gc/Dialect/LLVMIR/CMakeLists.txt | 2 +- include/gc/Dialect/LLVMIR/GenOps.td | 19 ++- include/gc/Target/LLVM/GEN/Utils.h | 53 +++++++ .../Dialect/GEN/GENToLLVMIRTranslation.h | 31 ++++ lib/gc/CMakeLists.txt | 2 +- lib/gc/Dialect/LLVMIR/CMakeLists.txt | 2 - lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp | 6 +- lib/gc/Target/CMakeLists.txt | 2 + lib/gc/Target/LLVM/CMakeLists.txt | 3 +- lib/gc/Target/LLVM/GEN/Target.cpp | 146 +++++++++++++++++- lib/gc/Target/LLVMIR/CMakeLists.txt | 1 + .../Target/LLVMIR/Dialect/GEN/CMakeLists.txt | 16 ++ .../Dialect/GEN/GENToLLVMIRTranslation.cpp | 76 +++++++++ .../GPU/ConvertGpuSignaturesToLLVM.cpp | 7 +- lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp | 7 +- .../GPU/Pipelines/LLVMSPVToBinaryPipeline.cpp | 0 lib/gc/Transforms/Pipeline.cpp | 15 -- src/gc-opt/gc-opt.cpp | 11 +- 18 files changed, 358 insertions(+), 41 deletions(-) create mode 100644 include/gc/Target/LLVM/GEN/Utils.h create mode 100644 include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h create mode 100644 lib/gc/Target/CMakeLists.txt create mode 100644 lib/gc/Target/LLVMIR/CMakeLists.txt create mode 100644 lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt create mode 100644 lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp delete mode 100644 lib/gc/Transforms/GPU/Pipelines/LLVMSPVToBinaryPipeline.cpp diff --git a/include/gc/Dialect/LLVMIR/CMakeLists.txt b/include/gc/Dialect/LLVMIR/CMakeLists.txt index b36730481..72d155166 100644 --- a/include/gc/Dialect/LLVMIR/CMakeLists.txt +++ b/include/gc/Dialect/LLVMIR/CMakeLists.txt @@ -3,4 +3,4 @@ add_mlir_doc(GenOps GENDialect Dialects/ -gen-dialect-doc -dialect=gen) set(LLVM_TARGET_DEFINITIONS GenOps.td) mlir_tablegen(GenOpsAttributes.h.inc -gen-attrdef-decls -attrdefs-dialect=gen) mlir_tablegen(GenOpsAttributes.cpp.inc -gen-attrdef-defs -attrdefs-dialect=gen) -add_public_tablegen_target(MLIRGENConversionsIncGen) \ No newline at end of file +add_public_tablegen_target(MLIRGENConversionsIncGen) diff --git a/include/gc/Dialect/LLVMIR/GenOps.td b/include/gc/Dialect/LLVMIR/GenOps.td index cbc071ee1..fab4cb3bd 100644 --- a/include/gc/Dialect/LLVMIR/GenOps.td +++ b/include/gc/Dialect/LLVMIR/GenOps.td @@ -13,6 +13,15 @@ def GEN_Dialect : Dialect { let hasOperationAttrVerify = 1; let extraClassDeclaration = [{ + /// Get the name of the attribute used to annotate external kernel + /// functions. + static StringRef getKernelFuncAttrName() { return "gen.kernel"; } + /// The address space value that represents global memory. + static constexpr unsigned kGlobalMemoryAddressSpace = 1; + /// The address space value that represents shared memory. + static constexpr unsigned kSharedMemoryAddressSpace = 3; + /// The address space value that represents private memory. + static constexpr unsigned kPrivateMemoryAddressSpace = 0; }]; let useDefaultAttributePrinterParser = 1; @@ -39,15 +48,17 @@ def GEN_TargettAttr : GEN_Attr<"GenTarget", "target"> { }]; let parameters = (ins DefaultValuedParameter<"int", "2", "Optimization level to apply.">:$O, - StringRefParameter<"Target triple.", "\"spir64-unknown-unknown\"">:$triple + StringRefParameter<"Target triple.", "\"spirv64-unknown-unknown\"">:$triple, + StringRefParameter<"Target chip.", "\"xe_1\"">:$chip ); let assemblyFormat = [{ - (`<` struct($O, $triple)^ `>`)? + (`<` struct($O, $triple, $chip)^ `>`)? }]; let builders = [ AttrBuilder<(ins CArg<"int", "2">:$optLevel, - CArg<"StringRef", "\"spir64-unknown-unknown\"">:$triple), [{ - return Base::get($_ctxt, optLevel, triple); + CArg<"StringRef", "\"spirv64-unknown-unknown\"">:$triple, + CArg<"StringRef", "\"xe_1\"">:$chip), [{ + return Base::get($_ctxt, optLevel, triple, chip); }]> ]; let skipDefaultBuilders = 1; diff --git a/include/gc/Target/LLVM/GEN/Utils.h b/include/gc/Target/LLVM/GEN/Utils.h new file mode 100644 index 000000000..22854b715 --- /dev/null +++ b/include/gc/Target/LLVM/GEN/Utils.h @@ -0,0 +1,53 @@ +//===- Utils.h - MLIR GEN target utils --------------------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This files declares GEN target related utility classes and functions. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_TARGET_LLVM_GEN_UTILS_H +#define MLIR_TARGET_LLVM_GEN_UTILS_H + +#include "gc/Dialect/LLVMIR/GENDialect.h" +#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h" +#include "mlir/Target/LLVM/ModuleToObject.h" + +namespace mlir { +namespace gen { + +StringRef getONEAPIToolkitPath(); + +/// Base class for all GEN serializations from GPU modules into binary strings. +/// By default this class serializes into LLVM bitcode. +class SerializeGPUModuleBase : public LLVM::ModuleToObject { +public: + /// Initializes the `toolkitPath` with the path in `targetOptions` or if empty + /// with the path in `getONEAPIToolkitPath`. + SerializeGPUModuleBase(Operation &module, GenTargetAttr target, + const gpu::TargetOptions &targetOptions = {}); + + // Initialize intermediate spirv target llvm backend + static void init(); + + /// Returns the target attribute. + GenTargetAttr getTarget() const; + + /// Returns the ONEAPI toolkit path. + StringRef getToolkitPath() const; + +protected: + /// GEN target attribute. + GenTargetAttr target; + + /// ONEAPI toolkit path. + std::string toolkitPath; +}; +} // namespace gen +} // namespace mlir + +#endif // MLIR_TARGET_LLVM_GEN_UTILS_H diff --git a/include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h b/include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h new file mode 100644 index 000000000..701199cf9 --- /dev/null +++ b/include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h @@ -0,0 +1,31 @@ +//===- GENToLLVMIRTranslation.h - GEN to LLVM IR ----------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This provides registration calls for GEN dialect to LLVM IR translation. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_TARGET_LLVMIR_DIALECT_GEN_GENTOLLVMIRTRANSLATION_H +#define MLIR_TARGET_LLVMIR_DIALECT_GEN_GENTOLLVMIRTRANSLATION_H + +namespace mlir { + +class DialectRegistry; +class MLIRContext; + +/// Register the GEN dialect and the translation from it to the LLVM IR in the +/// given registry; +void registerGENDialectTranslation(DialectRegistry ®istry); + +/// Register the GEN dialect and the translation from it in the registry +/// associated with the given context. +void registerGENDialectTranslation(MLIRContext &context); + +} // namespace mlir + +#endif // MLIR_TARGET_LLVMIR_DIALECT_GEN_GENTOLLVMIRTRANSLATION_H diff --git a/lib/gc/CMakeLists.txt b/lib/gc/CMakeLists.txt index 0ccb16e49..441d43b8c 100644 --- a/lib/gc/CMakeLists.txt +++ b/lib/gc/CMakeLists.txt @@ -2,4 +2,4 @@ add_subdirectory(CAPI) add_subdirectory(Dialect) add_subdirectory(Transforms) add_subdirectory(ExecutionEngine) -add_subdirectory(Target/LLVM) +add_subdirectory(Target) diff --git a/lib/gc/Dialect/LLVMIR/CMakeLists.txt b/lib/gc/Dialect/LLVMIR/CMakeLists.txt index 9d61678e9..05ec2078a 100644 --- a/lib/gc/Dialect/LLVMIR/CMakeLists.txt +++ b/lib/gc/Dialect/LLVMIR/CMakeLists.txt @@ -6,8 +6,6 @@ add_mlir_dialect_library(MLIRGENDialect ${PROJECT_SOURCE_DIR}/include/gc/Dialect/LLVMIR DEPENDS -# MLIRGPUCompilationAttrInterfacesIncGen -# MLIRGENOpsIncGen MLIRGENConversionsIncGen LINK_COMPONENTS diff --git a/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp b/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp index 336f1e978..93a713659 100644 --- a/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp +++ b/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp @@ -12,7 +12,7 @@ using namespace gen; LogicalResult GenTargetAttr::verify(function_ref emitError, int O, - StringRef triple) { + StringRef triple, StringRef chip) { if (O < 0 || O > 3) { emitError() << "The optimization level must be a number between 0 and 3."; return failure(); @@ -21,6 +21,10 @@ GenTargetAttr::verify(function_ref emitError, int O, emitError() << "The target triple cannot be empty."; return failure(); } + if (chip.empty()) { + emitError() << "The target chip cannot be empty."; + return failure(); + } return success(); } diff --git a/lib/gc/Target/CMakeLists.txt b/lib/gc/Target/CMakeLists.txt new file mode 100644 index 000000000..3a8c89369 --- /dev/null +++ b/lib/gc/Target/CMakeLists.txt @@ -0,0 +1,2 @@ +add_subdirectory(LLVMIR) +add_subdirectory(LLVM) diff --git a/lib/gc/Target/LLVM/CMakeLists.txt b/lib/gc/Target/LLVM/CMakeLists.txt index 1ad4b0b40..bd5475919 100644 --- a/lib/gc/Target/LLVM/CMakeLists.txt +++ b/lib/gc/Target/LLVM/CMakeLists.txt @@ -13,4 +13,5 @@ add_mlir_dialect_library(MLIRGENTarget MLIRSupport MLIRGPUDialect MLIRTargetLLVM - ) \ No newline at end of file + MLIRGENToLLVMIRTranslation + ) diff --git a/lib/gc/Target/LLVM/GEN/Target.cpp b/lib/gc/Target/LLVM/GEN/Target.cpp index 7f15bd8eb..a47b67cf1 100644 --- a/lib/gc/Target/LLVM/GEN/Target.cpp +++ b/lib/gc/Target/LLVM/GEN/Target.cpp @@ -1,4 +1,4 @@ -//===- Target.cpp - MLIR LLVM XE target compilation -------------*- C++ -*-===// +//===- Target.cpp - MLIR LLVM GEN target compilation ------------*- C++ -*-===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -6,25 +6,31 @@ // //===----------------------------------------------------------------------===// // -// This file defines Xe target related functions including registration -// calls for the `#xe.target` compilation attribute. +// This file defines GEN target related functions including registration +// calls for the `#gen.target` compilation attribute. // //===----------------------------------------------------------------------===// #include "gc/Target/LLVM/GEN/Target.h" #include "gc/Dialect/LLVMIR/GENDialect.h" +#include "gc/Target/LLVM/GEN/Utils.h" #include "mlir/Dialect/GPU/IR/CompilationInterfaces.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/IR/ExtensibleDialect.h" +#include "mlir/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.h" +#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h" + +#include "llvm/IR/LegacyPassManager.h" +#include "llvm/Support/TargetSelect.h" +#include "llvm/Target/TargetMachine.h" using namespace mlir; using namespace mlir::gen; namespace { - -// Xe implementation of the gpu:TargetAttrInterface. +// Gen implementation of the gpu:TargetAttrInterface. class GenTargetAttrImpl : public gpu::TargetAttrInterface::FallbackModel { public: @@ -51,6 +57,131 @@ void mlir::gen::registerGenTargetInterfaceExternalModels(MLIRContext &context) { context.appendDialectRegistry(registry); } +StringRef mlir::gen::getONEAPIToolkitPath() { + if (const char *var = std::getenv("ONEAPI_ROOT")) + return var; + return ""; +} + +SerializeGPUModuleBase::SerializeGPUModuleBase( + Operation &module, GenTargetAttr target, + const gpu::TargetOptions &targetOptions) + : ModuleToObject(module, target.getTriple(), target.getChip(), {}, + target.getO()), + target(target), toolkitPath(targetOptions.getToolkitPath()) { + if (toolkitPath.empty()) + toolkitPath = getONEAPIToolkitPath(); +} + +void SerializeGPUModuleBase::init() { + static llvm::once_flag initializeBackendOnce; + llvm::call_once(initializeBackendOnce, []() { +#if LLVM_HAS_SPIRV_TARGET + LLVMInitializeSPIRVTarget(); + LLVMInitializeSPIRVTargetInfo(); + LLVMInitializeSPIRVTargetMC(); + LLVMInitializeSPIRVAsmPrinter(); +#endif + }); +} + +GenTargetAttr SerializeGPUModuleBase::getTarget() const { return target; } + +StringRef SerializeGPUModuleBase::getToolkitPath() const { return toolkitPath; } + +namespace { +class GenSerializer : public SerializeGPUModuleBase { +public: + GenSerializer(Operation &module, GenTargetAttr target, + const gpu::TargetOptions &targetOptions); + + gpu::GPUModuleOp getOperation(); + + std::optional> + compileToBinary(const std::string &serializedISA); + + std::optional> + moduleToObject(llvm::Module &llvmModule) override; + +private: + std::optional + translateToSPIRVBinary(llvm::Module &llvmModule, + llvm::TargetMachine &targetMachine); + gpu::TargetOptions targetOptions; +}; +} // namespace + +GenSerializer::GenSerializer(Operation &module, GenTargetAttr target, + const gpu::TargetOptions &targetOptions) + : SerializeGPUModuleBase(module, target, targetOptions) {} + +gpu::GPUModuleOp GenSerializer::getOperation() { + return dyn_cast(&SerializeGPUModuleBase::getOperation()); +} + +std::optional> +GenSerializer::moduleToObject(llvm::Module &llvmModule) { + // Return LLVM IR if the compilation target is `offload`. + if (targetOptions.getCompilationTarget() == gpu::CompilationTarget::Offload) + return SerializeGPUModuleBase::moduleToObject(llvmModule); + +#if !LLVM_HAS_SPIRV_TARGET + getOperation()->emitError( + "The `SPIRV` target was not built. Please enable it when building LLVM."); + return std::nullopt; +#endif // LLVM_HAS_SPIRV_TARGET + + std::optional targetMachine = + getOrCreateTargetMachine(); + if (!targetMachine) { + getOperation().emitError() << "Target Machine unavailable for triple " + << triple << ", can't compile with LLVM\n"; + return std::nullopt; + } + + std::optional serializedISA = + translateToISA(llvmModule, **targetMachine); + if (!serializedISA) { + getOperation().emitError() << "Failed translating the module to ISA."; + return std::nullopt; + } + + // Return SPIRV if the compilation target is `assembly`. + if (targetOptions.getCompilationTarget() == + gpu::CompilationTarget::Assembly) { + // Make sure to include the null terminator. + StringRef bin(serializedISA->c_str(), serializedISA->size() + 1); + return SmallVector(bin.begin(), bin.end()); + } + + return compileToBinary(*serializedISA); +} + +std::optional> +GenSerializer::compileToBinary(const std::string &serializedSPV) { + // FIXME + return SmallVector(serializedSPV.begin(), serializedSPV.end()); +} + +std::optional +translateToSPIRVBinary(llvm::Module &llvmModule, + llvm::TargetMachine &targetMachine) { + std::string targetISA; + llvm::raw_string_ostream stream(targetISA); + + { // Drop pstream after this to prevent the ISA from being stuck buffering + llvm::buffer_ostream pstream(stream); + llvm::legacy::PassManager codegenPasses; + + if (targetMachine.addPassesToEmitFile(codegenPasses, pstream, nullptr, + llvm::CodeGenFileType::ObjectFile)) + return std::nullopt; + + codegenPasses.run(llvmModule); + } + return stream.str(); +} + std::optional> GenTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module, const gpu::TargetOptions &options) const { @@ -62,7 +193,10 @@ GenTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module, return std::nullopt; } - // todo + GenSerializer serializer(*module, cast(attribute), options); + serializer.init(); + + return serializer.run(); } Attribute diff --git a/lib/gc/Target/LLVMIR/CMakeLists.txt b/lib/gc/Target/LLVMIR/CMakeLists.txt new file mode 100644 index 000000000..51b942d57 --- /dev/null +++ b/lib/gc/Target/LLVMIR/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(Dialect/GEN) diff --git a/lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt b/lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt new file mode 100644 index 000000000..e84dc22b9 --- /dev/null +++ b/lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt @@ -0,0 +1,16 @@ +add_mlir_translation_library(MLIRGENToLLVMIRTranslation + GENToLLVMIRTranslation.cpp + + DEPENDS + MLIRGENConversionsIncGen + + LINK_COMPONENTS + Core + + LINK_LIBS PUBLIC + MLIRIR + MLIRLLVMDialect + MLIRGENDialect + MLIRSupport + MLIRTargetLLVMIRExport + ) diff --git a/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp b/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp new file mode 100644 index 000000000..ccf055ad6 --- /dev/null +++ b/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp @@ -0,0 +1,76 @@ +//===- GENToLLVMIRTranslation.cpp - Translate GEN to LLVM IR --------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file implements a translation between the MLIR GEN dialect and +// LLVM IR. +// +//===----------------------------------------------------------------------===// + +#include "gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h" +#include "gc/Dialect/LLVMIR/GENDialect.h" +#include "mlir/IR/BuiltinAttributes.h" +#include "mlir/IR/Operation.h" +#include "mlir/Target/LLVMIR/ModuleTranslation.h" + +#include "llvm/IR/ConstantRange.h" +#include "llvm/IR/IRBuilder.h" +#include "llvm/Support/raw_ostream.h" + +using namespace mlir; +using namespace mlir::LLVM; + +namespace { +/// Implementation of the dialect interface that converts operations belonging +/// to the GEN dialect to LLVM IR. +class GENDialectLLVMIRTranslationInterface + : public LLVMTranslationDialectInterface { +public: + using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface; + + /// Translates the given operation to LLVM IR using the provided IR builder + /// and saving the state in `moduleTranslation`. + LogicalResult + convertOperation(Operation *op, llvm::IRBuilderBase &builder, + LLVM::ModuleTranslation &moduleTranslation) const final { + // no operations, not supposed to be called + return failure(); + } + + /// Attaches module-level metadata for functions marked as kernels. + LogicalResult + amendOperation(Operation *op, ArrayRef instructions, + NamedAttribute attribute, + LLVM::ModuleTranslation &moduleTranslation) const final { + auto func = dyn_cast(op); + if (!func) + return failure(); + llvm::LLVMContext &llvmContext = moduleTranslation.getLLVMContext(); + llvm::Function *llvmFunc = moduleTranslation.lookupFunction(func.getName()); + + if (attribute.getName() == gen::GENDialect::getKernelFuncAttrName()) { + // todo; note: migth not need it as we'll have storage classes translated + // already + } + + return success(); + } +}; +} // namespace + +void mlir::registerGENDialectTranslation(DialectRegistry ®istry) { + registry.insert(); + registry.addExtension(+[](MLIRContext *ctx, gen::GENDialect *dialect) { + dialect->addInterfaces(); + }); +} + +void mlir::registerGENDialectTranslation(MLIRContext &context) { + DialectRegistry registry; + registerGENDialectTranslation(registry); + context.appendDialectRegistry(registry); +} diff --git a/lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp b/lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp index c2d129070..0bcdc22bf 100644 --- a/lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp +++ b/lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp @@ -8,6 +8,7 @@ #include "gc/Transforms/Passes.h" +#include "gc/Dialect/LLVMIR/GENDialect.h" #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/TypeConverter.h" #include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h" @@ -50,8 +51,10 @@ void ConvertGpuSignaturesToLLVM::runOnOperation() { patterns.add(converter); patterns.add( - converter, 0 /*local*/, 3 /*shared*/, - StringAttr::get(&converter.getContext(), "xe.kernel")); + converter, gen::GENDialect::kPrivateMemoryAddressSpace /*local*/, + gen::GENDialect::kSharedMemoryAddressSpace /*shared*/, + StringAttr::get(&converter.getContext(), + gen::GENDialect::getKernelFuncAttrName())); if (failed(applyPartialConversion(gpuModule, target, std::move(patterns)))) signalPassFailure(); diff --git a/lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp b/lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp index 59d0843c4..8e84e66ae 100644 --- a/lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp +++ b/lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp @@ -32,11 +32,12 @@ struct GpuGenAttachTarget void GpuGenAttachTarget::runOnOperation() { OpBuilder builder(&getContext()); auto target = - builder.getAttr(2, "spir64-unknown-unknown"); + builder.getAttr(2, "spirv64-unknown-unknown"); getOperation()->walk([&](gpu::GPUModuleOp gpuModule) { SmallVector targets; - if (std::optional attrs = gpuModule.getTargets()) - targets.append(attrs->getValue().begin(), attrs->getValue().end()); + // Temporary solution to avoid an attempt to create a spirv binary + // if (std::optional attrs = gpuModule.getTargets()) + // targets.append(attrs->getValue().begin(), attrs->getValue().end()); targets.push_back(target); // Remove any duplicate targets. targets.erase(llvm::unique(targets), targets.end()); diff --git a/lib/gc/Transforms/GPU/Pipelines/LLVMSPVToBinaryPipeline.cpp b/lib/gc/Transforms/GPU/Pipelines/LLVMSPVToBinaryPipeline.cpp deleted file mode 100644 index e69de29bb..000000000 diff --git a/lib/gc/Transforms/Pipeline.cpp b/lib/gc/Transforms/Pipeline.cpp index 5d1fed278..7d487f149 100644 --- a/lib/gc/Transforms/Pipeline.cpp +++ b/lib/gc/Transforms/Pipeline.cpp @@ -149,21 +149,6 @@ void populateCPUPipeline(mlir::OpPassManager &pm) { populateLLVMPasses(pm); } -void populateGPUPipeline(mlir::OpPassManager &pm) { - pm.addPass(createLinalgGeneralizeNamedOpsPass()); - bufferization::OneShotBufferizationOptions options; - options.bufferizeFunctionBoundaries = true; - options.setFunctionBoundaryTypeConversion( - bufferization::LayoutMapOption::IdentityLayoutMap); - pm.addPass(bufferization::createOneShotBufferizePass(options)); -} - -void buildGPUBinaryGenerationPipeline(mlir::OpPassManager &pm) { - GpuModuleToBinaryPassOptions gpuModuleToBinaryPassOptions; - gpuModuleToBinaryPassOptions.compilationTarget = "genisa"; - pm.addPass(createGpuModuleToBinaryPass(gpuModuleToBinaryPassOptions)); -} - void registerCPUPipeline() { PassPipelineRegistration<>("gc-cpu-pipeline", "The CPU pipeline for Graph Compiler", diff --git a/src/gc-opt/gc-opt.cpp b/src/gc-opt/gc-opt.cpp index 1264fcd40..df3e15603 100644 --- a/src/gc-opt/gc-opt.cpp +++ b/src/gc-opt/gc-opt.cpp @@ -23,15 +23,14 @@ #include "gc/Dialect/Microkernel/MicrokernelDialect.h" #ifdef GC_HAS_ONEDNN_DIALECT #include "gc/Dialect/OneDNNGraph/OneDNNGraphDialect.h" -<<<<<<< HEAD #endif - ======= #include "gc/Target/LLVM/GEN/Target.h" - >>>>>>> abeedf6 (Add gen dialect to hold the gen target) +#include "gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h" #include "gc/Transforms/Passes.h" #include "mlir/InitAllDialects.h" #include "mlir/InitAllExtensions.h" #include "mlir/InitAllPasses.h" +#include "mlir/Target/LLVMIR/Dialect/All.h" #include "mlir/Tools/mlir-opt/MlirOptMain.h" #ifdef GC_USE_IMEX @@ -39,8 +38,8 @@ #include #endif - namespace mlir::gc { - void registerCPUPipeline(); +namespace mlir::gc { +void registerCPUPipeline(); } // namespace mlir::gc int main(int argc, char *argv[]) { @@ -67,7 +66,9 @@ int main(int argc, char *argv[]) { registry.insert(); mlir::registerAllDialects(registry); mlir::registerAllExtensions(registry); + mlir::registerAllToLLVMIRTranslations(registry); mlir::gen::registerGenTargetInterfaceExternalModels(registry); + mlir::registerGENDialectTranslation(registry); #ifdef GC_USE_GPU registry.insert<::imex::xetile::XeTileDialect, ::imex::gpux::GPUXDialect>(); #endif From 59d586bb505e8633402832e69213c35e690e8edb Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Tue, 30 Jul 2024 05:41:54 -0700 Subject: [PATCH 05/19] Add ocloc integration --- include/gc/Dialect/LLVMIR/GenOps.td | 4 +- lib/gc/Target/LLVM/GEN/Target.cpp | 124 +++++++++++++++++++++++++--- 2 files changed, 113 insertions(+), 15 deletions(-) diff --git a/include/gc/Dialect/LLVMIR/GenOps.td b/include/gc/Dialect/LLVMIR/GenOps.td index fab4cb3bd..134edb2c3 100644 --- a/include/gc/Dialect/LLVMIR/GenOps.td +++ b/include/gc/Dialect/LLVMIR/GenOps.td @@ -49,7 +49,7 @@ def GEN_TargettAttr : GEN_Attr<"GenTarget", "target"> { let parameters = (ins DefaultValuedParameter<"int", "2", "Optimization level to apply.">:$O, StringRefParameter<"Target triple.", "\"spirv64-unknown-unknown\"">:$triple, - StringRefParameter<"Target chip.", "\"xe_1\"">:$chip + StringRefParameter<"Target chip.", "\"pvc\"">:$chip ); let assemblyFormat = [{ (`<` struct($O, $triple, $chip)^ `>`)? @@ -57,7 +57,7 @@ def GEN_TargettAttr : GEN_Attr<"GenTarget", "target"> { let builders = [ AttrBuilder<(ins CArg<"int", "2">:$optLevel, CArg<"StringRef", "\"spirv64-unknown-unknown\"">:$triple, - CArg<"StringRef", "\"xe_1\"">:$chip), [{ + CArg<"StringRef", "\"pvc\"">:$chip), [{ return Base::get($_ctxt, optLevel, triple, chip); }]> ]; diff --git a/lib/gc/Target/LLVM/GEN/Target.cpp b/lib/gc/Target/LLVM/GEN/Target.cpp index a47b67cf1..7d393c486 100644 --- a/lib/gc/Target/LLVM/GEN/Target.cpp +++ b/lib/gc/Target/LLVM/GEN/Target.cpp @@ -23,6 +23,13 @@ #include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h" #include "llvm/IR/LegacyPassManager.h" +#include "llvm/Support/FileSystem.h" +#include "llvm/Support/FileUtilities.h" +#include "llvm/Support/FormatVariadic.h" +#include "llvm/Support/MemoryBuffer.h" +#include "llvm/Support/Path.h" +#include "llvm/Support/Process.h" +#include "llvm/Support/Program.h" #include "llvm/Support/TargetSelect.h" #include "llvm/Target/TargetMachine.h" @@ -60,7 +67,7 @@ void mlir::gen::registerGenTargetInterfaceExternalModels(MLIRContext &context) { StringRef mlir::gen::getONEAPIToolkitPath() { if (const char *var = std::getenv("ONEAPI_ROOT")) return var; - return ""; + return "/usr/"; } SerializeGPUModuleBase::SerializeGPUModuleBase( @@ -103,7 +110,12 @@ class GenSerializer : public SerializeGPUModuleBase { std::optional> moduleToObject(llvm::Module &llvmModule) override; + std::optional findTool(StringRef tool); + private: + using TmpFile = std::pair, llvm::FileRemover>; + std::optional createTemp(StringRef name, StringRef suffix); + std::optional translateToSPIRVBinary(llvm::Module &llvmModule, llvm::TargetMachine &targetMachine); @@ -119,6 +131,31 @@ gpu::GPUModuleOp GenSerializer::getOperation() { return dyn_cast(&SerializeGPUModuleBase::getOperation()); } +std::optional +GenSerializer::createTemp(StringRef name, StringRef suffix) { + llvm::SmallString<128> filename; + std::error_code ec = + llvm::sys::fs::createTemporaryFile(name, suffix, filename); + if (ec) { + getOperation().emitError() << "Couldn't create the temp file: `" << filename + << "`, error message: " << ec.message(); + return std::nullopt; + } + return TmpFile(filename, llvm::FileRemover(filename.c_str())); +} + +std::optional GenSerializer::findTool(StringRef tool) { + if (std::optional toolPath = + llvm::sys::Process::FindInEnvPath("PATH", tool)) + return *toolPath; + getOperation().emitError() + << "Couldn't find the `" << tool + << "` binary. Please specify the toolkit " + "path, add the compiler to $PATH, or set one of the environment " + "variables in `gen::getGENToolkitPath()`."; + return std::nullopt; +} + std::optional> GenSerializer::moduleToObject(llvm::Module &llvmModule) { // Return LLVM IR if the compilation target is `offload`. @@ -139,33 +176,94 @@ GenSerializer::moduleToObject(llvm::Module &llvmModule) { return std::nullopt; } - std::optional serializedISA = - translateToISA(llvmModule, **targetMachine); - if (!serializedISA) { - getOperation().emitError() << "Failed translating the module to ISA."; - return std::nullopt; - } - // Return SPIRV if the compilation target is `assembly`. if (targetOptions.getCompilationTarget() == gpu::CompilationTarget::Assembly) { + std::optional serializedISA = + translateToISA(llvmModule, **targetMachine); + if (!serializedISA) { + getOperation().emitError() << "Failed translating the module to ISA."; + return std::nullopt; + } // Make sure to include the null terminator. StringRef bin(serializedISA->c_str(), serializedISA->size() + 1); return SmallVector(bin.begin(), bin.end()); } - return compileToBinary(*serializedISA); + std::optional serializedSPIRVBinary = + translateToSPIRVBinary(llvmModule, **targetMachine); + if (!serializedSPIRVBinary) { + getOperation().emitError() << "Failed translating the module to Binary."; + return std::nullopt; + } + + return compileToBinary(*serializedSPIRVBinary); } std::optional> GenSerializer::compileToBinary(const std::string &serializedSPV) { - // FIXME - return SmallVector(serializedSPV.begin(), serializedSPV.end()); + std::optional ocloc = findTool("ocloc"); + if (!ocloc) + return std::nullopt; + + std::string basename = + llvm::formatv("mlir-{0}-{1}-{2}", getOperation().getNameAttr().getValue(), + getTarget().getTriple(), getTarget().getChip()); + + std::optional spvFile = createTemp(basename, "spv"); + if (!spvFile) + return std::nullopt; + std::optional binaryFile = createTemp(basename, "bin"); + if (!binaryFile) + return std::nullopt; + + Location loc = getOperation().getLoc(); + std::error_code ec; + { + llvm::raw_fd_ostream spvStream(spvFile->first, ec); + if (ec) { + emitError(loc) << "Couldn't open the file: `" << spvFile->first + << "`, error message: " << ec.message(); + return std::nullopt; + } + spvStream << serializedSPV; + if (spvStream.has_error()) { + emitError(loc) << "An error occurred while writing the SPIRV to: `" + << spvFile->first << "`."; + return std::nullopt; + } + spvStream.flush(); + } + + SmallVector oclocArgs( + {StringRef("compile"), StringRef("-device"), getTarget().getChip(), + StringRef("-spirv_input"), StringRef("-file"), StringRef(spvFile->first), + StringRef("-o"), StringRef(binaryFile->first)}); + + std::string message; + if (llvm::sys::ExecuteAndWait(ocloc.value(), oclocArgs, + /*Env=*/std::nullopt, + /*Redirects=*/std::nullopt, + /*SecondsToWait=*/0, + /*MemoryLimit=*/0, + /*ErrMsg=*/&message)) { + emitError(loc) << " ocloc invocation failed. Message:\n" << message; + return std::nullopt; + } + llvm::ErrorOr> binaryBuffer = + llvm::MemoryBuffer::getFile(binaryFile->first); + if (!binaryBuffer) { + emitError(loc) << "Couldn't open the file: `" << binaryFile->first + << "`, error message: " << binaryBuffer.getError().message(); + return std::nullopt; + } + StringRef result = (*binaryBuffer)->getBuffer(); + return SmallVector(result.begin(), result.end()); } std::optional -translateToSPIRVBinary(llvm::Module &llvmModule, - llvm::TargetMachine &targetMachine) { +GenSerializer::translateToSPIRVBinary(llvm::Module &llvmModule, + llvm::TargetMachine &targetMachine) { std::string targetISA; llvm::raw_string_ostream stream(targetISA); From dbeb7048b6f199694a10b8730365ca31f4d3fe98 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Tue, 30 Jul 2024 06:59:21 -0700 Subject: [PATCH 06/19] Add gpu pipeline registration --- lib/gc/CAPI/Passes.cpp | 2 ++ lib/gc/Transforms/Pipeline.cpp | 37 ++++++++++++++++++++++++++++++++++ src/gc-opt/gc-opt.cpp | 2 ++ 3 files changed, 41 insertions(+) diff --git a/lib/gc/CAPI/Passes.cpp b/lib/gc/CAPI/Passes.cpp index 07ff402e5..977b8706b 100644 --- a/lib/gc/CAPI/Passes.cpp +++ b/lib/gc/CAPI/Passes.cpp @@ -18,6 +18,7 @@ using namespace mlir::cpuruntime; namespace mlir::gc { void registerCPUPipeline(); +void registerGPUPipeline(); } // namespace mlir::gc #ifdef __cplusplus @@ -29,6 +30,7 @@ extern "C" { MLIR_CAPI_EXPORTED void mlirRegisterAllGCPassesAndPipelines() { registerCPUPipeline(); + registerGPUPipeline(); mlirRegisterCPURuntimePasses(); mlirRegisterGraphCompilerPasses(); } diff --git a/lib/gc/Transforms/Pipeline.cpp b/lib/gc/Transforms/Pipeline.cpp index 7d487f149..999ab0d0c 100644 --- a/lib/gc/Transforms/Pipeline.cpp +++ b/lib/gc/Transforms/Pipeline.cpp @@ -10,6 +10,8 @@ #include "mlir/Dialect/Arith/Transforms/Passes.h" #include "mlir/Dialect/Bufferization/Transforms/OneShotAnalysis.h" #include "mlir/Dialect/Bufferization/Transforms/Passes.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" +#include "mlir/Dialect/GPU/Transforms/Passes.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/LLVMIR/Transforms/Passes.h" #include "mlir/Dialect/Linalg/Passes.h" @@ -149,10 +151,45 @@ void populateCPUPipeline(mlir::OpPassManager &pm) { populateLLVMPasses(pm); } +void populateGPUPipeline(mlir::OpPassManager &pm) { + pm.addNestedPass(createLinalgGeneralizeNamedOpsPass()); + bufferization::OneShotBufferizationOptions options; + options.bufferizeFunctionBoundaries = true; + options.setFunctionBoundaryTypeConversion( + bufferization::LayoutMapOption::IdentityLayoutMap); + pm.addPass(bufferization::createOneShotBufferizePass(options)); + pm.addPass(createCSEPass()); + pm.addNestedPass(createConvertLinalgToParallelLoopsPass()); + pm.addNestedPass(createGpuMapParallelLoopsPass()); + pm.addNestedPass(createParallelLoopToGpuPass()); + pm.addNestedPass(createLowerAffinePass()); + pm.addPass(memref::createNormalizeMemRefsPass()); + pm.addPass(createGpuKernelOutliningPass()); + pm.addPass(createGpuLegalizeModule()); + pm.addPass(memref::createFoldMemRefAliasOpsPass()); + ConvertIndexToLLVMPassOptions idxOptions; + idxOptions.indexBitwidth = 32; + pm.addNestedPass(createConvertIndexToLLVMPass(idxOptions)); + pm.addNestedPass(createConvertGpuOpsToLLVMSPVOps()); + pm.addPass(createCanonicalizerPass()); + pm.addNestedPass(createConvertGpuSignaturesToLLVM()); + pm.addPass(createGpuToLLVMConversionPass()); + pm.addPass(createCanonicalizerPass()); + pm.addPass(createGpuGenAttachTarget()); + GpuModuleToBinaryPassOptions gpuModuleToBinaryPassOptions; + pm.addPass(createGpuModuleToBinaryPass(gpuModuleToBinaryPassOptions)); +} + void registerCPUPipeline() { PassPipelineRegistration<>("gc-cpu-pipeline", "The CPU pipeline for Graph Compiler", populateCPUPipeline); } +void registerGPUPipeline() { + PassPipelineRegistration<>("gc-gpu-pipeline", + "The GPU pipeline for Graph Compiler", + populateGPUPipeline); +} + } // namespace mlir::gc diff --git a/src/gc-opt/gc-opt.cpp b/src/gc-opt/gc-opt.cpp index df3e15603..3ee222cfb 100644 --- a/src/gc-opt/gc-opt.cpp +++ b/src/gc-opt/gc-opt.cpp @@ -40,6 +40,7 @@ namespace mlir::gc { void registerCPUPipeline(); +void registerGPUPipeline(); } // namespace mlir::gc int main(int argc, char *argv[]) { @@ -54,6 +55,7 @@ int main(int argc, char *argv[]) { #endif mlir::registerAllPasses(); mlir::gc::registerCPUPipeline(); + mlir::gc::registerGPUPipeline(); mlir::gc::registerGraphCompilerPasses(); mlir::cpuruntime::registerCPURuntimePasses(); mlir::DialectRegistry registry; From 6ee3d0bc417511b7a9c327923786787582798825 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Tue, 30 Jul 2024 10:14:21 -0700 Subject: [PATCH 07/19] Fix typo --- include/gc/Target/LLVM/GEN/Target.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/gc/Target/LLVM/GEN/Target.h b/include/gc/Target/LLVM/GEN/Target.h index 6d3438cb7..594b34752 100644 --- a/include/gc/Target/LLVM/GEN/Target.h +++ b/include/gc/Target/LLVM/GEN/Target.h @@ -1,4 +1,4 @@ -//===- Target.h - MLIR Xe target registration -------------------*- C++ -*-===// +//===- Target.h - MLIR GEN target registration ------------------*- C++ -*-===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. From 92918553dcd05977dc0981c85a7789d42ed3d34d Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Tue, 30 Jul 2024 10:38:59 -0700 Subject: [PATCH 08/19] Fix static gc-opt build & add comments for components registration --- src/gc-opt/CMakeLists.txt | 3 +++ src/gc-opt/gc-opt.cpp | 4 ++++ 2 files changed, 7 insertions(+) diff --git a/src/gc-opt/CMakeLists.txt b/src/gc-opt/CMakeLists.txt index 96d5ae860..1e5f4e38f 100644 --- a/src/gc-opt/CMakeLists.txt +++ b/src/gc-opt/CMakeLists.txt @@ -26,12 +26,15 @@ if(GC_DEV_LINK_LLVM_DYLIB) ) get_property(dialect_libs GLOBAL PROPERTY GC_DIALECT_LIBS) get_property(conversion_libs GLOBAL PROPERTY GC_PASS_LIBS) + get_property(extension_libs GLOBAL PROPERTY GC_EXTENSION_LIBS) else() set(MLIR_LINK_COMPONENTS MLIROptLib + MLIRToLLVMIRTranslationRegistration ) get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS) get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS) + get_property(extension_libs GLOBAL PROPERTY MLIR_EXTENSION_LIBS) endif() add_llvm_executable(gc-opt gc-opt.cpp) diff --git a/src/gc-opt/gc-opt.cpp b/src/gc-opt/gc-opt.cpp index 3ee222cfb..5c196762c 100644 --- a/src/gc-opt/gc-opt.cpp +++ b/src/gc-opt/gc-opt.cpp @@ -67,7 +67,11 @@ int main(int argc, char *argv[]) { registry.insert(); registry.insert(); mlir::registerAllDialects(registry); + // covers lowerings for weird dialects like ub + // TODO: avoid `registerALL` to remove this mlir::registerAllExtensions(registry); + // Adds missing `LLVMTranslationDialectInterface` registration for dialect for + // gpu.module op mlir::registerAllToLLVMIRTranslations(registry); mlir::gen::registerGenTargetInterfaceExternalModels(registry); mlir::registerGENDialectTranslation(registry); From 3cbebc18ce06b4f713c0bcaddd567752d6fe3bd8 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Tue, 30 Jul 2024 12:03:23 -0700 Subject: [PATCH 09/19] Fix warnings --- lib/gc/Target/LLVM/GEN/Target.cpp | 2 +- .../Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp | 9 ++------- 2 files changed, 3 insertions(+), 8 deletions(-) diff --git a/lib/gc/Target/LLVM/GEN/Target.cpp b/lib/gc/Target/LLVM/GEN/Target.cpp index 7d393c486..93bcc331b 100644 --- a/lib/gc/Target/LLVM/GEN/Target.cpp +++ b/lib/gc/Target/LLVM/GEN/Target.cpp @@ -105,7 +105,7 @@ class GenSerializer : public SerializeGPUModuleBase { gpu::GPUModuleOp getOperation(); std::optional> - compileToBinary(const std::string &serializedISA); + compileToBinary(const std::string &serializedSPV); std::optional> moduleToObject(llvm::Module &llvmModule) override; diff --git a/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp b/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp index ccf055ad6..b84076385 100644 --- a/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp +++ b/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp @@ -49,13 +49,8 @@ class GENDialectLLVMIRTranslationInterface auto func = dyn_cast(op); if (!func) return failure(); - llvm::LLVMContext &llvmContext = moduleTranslation.getLLVMContext(); - llvm::Function *llvmFunc = moduleTranslation.lookupFunction(func.getName()); - - if (attribute.getName() == gen::GENDialect::getKernelFuncAttrName()) { - // todo; note: migth not need it as we'll have storage classes translated - // already - } + // todo; note: migth not need it as we'll have storage classes translated + // already return success(); } From 8193d6561f10269ae94709125d482fc04da8b203 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Thu, 1 Aug 2024 04:28:13 -0700 Subject: [PATCH 10/19] Disable clang-tidy on the attribute definition --- lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp b/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp index 93a713659..aa032e89b 100644 --- a/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp +++ b/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp @@ -34,10 +34,13 @@ LogicalResult GENDialect::verifyOperationAttribute(Operation *op, } void GENDialect::initialize() { + // clang-tidy is confused by the registration mechanism + // NOLINTBEGIN addAttributes< #define GET_ATTRDEF_LIST #include "gc/Dialect/LLVMIR/GenOpsAttributes.cpp.inc" >(); + // NOLINTEND allowUnknownOperations(); declarePromisedInterface(); From 643f363691a989738e24b1336f2130ecededd006 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Thu, 1 Aug 2024 04:50:05 -0700 Subject: [PATCH 11/19] Fix licences --- include/gc/Dialect/LLVMIR/GENDialect.h | 4 ++-- include/gc/Dialect/LLVMIR/GenOps.td | 7 +++++++ include/gc/Target/LLVM/GEN/Target.h | 4 ++-- include/gc/Target/LLVM/GEN/Utils.h | 4 ++-- .../gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h | 4 ++-- lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp | 7 +++++++ lib/gc/Target/LLVM/GEN/Target.cpp | 4 ++-- .../Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp | 4 ++-- lib/gc/Transforms/GPU/GPUOpsLowering.h | 4 ++-- 9 files changed, 28 insertions(+), 14 deletions(-) diff --git a/include/gc/Dialect/LLVMIR/GENDialect.h b/include/gc/Dialect/LLVMIR/GENDialect.h index 8745a830b..4c9149fab 100644 --- a/include/gc/Dialect/LLVMIR/GENDialect.h +++ b/include/gc/Dialect/LLVMIR/GENDialect.h @@ -1,6 +1,6 @@ -//===- GENDialect.h - MLIR GEN target definitions ---------------*- C++ -*-===// +//===-- GENDialect.h - MLIR GEN target definitions --------------*- C++ -*-===// // -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // diff --git a/include/gc/Dialect/LLVMIR/GenOps.td b/include/gc/Dialect/LLVMIR/GenOps.td index 134edb2c3..7139d217f 100644 --- a/include/gc/Dialect/LLVMIR/GenOps.td +++ b/include/gc/Dialect/LLVMIR/GenOps.td @@ -1,3 +1,10 @@ +//===-- GenOps.td - Gen dialect definition -----------------*- tablegen -*-===// +// +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// #ifndef GENIR_OPS #define GENIR_OPS diff --git a/include/gc/Target/LLVM/GEN/Target.h b/include/gc/Target/LLVM/GEN/Target.h index 594b34752..462bc78e5 100644 --- a/include/gc/Target/LLVM/GEN/Target.h +++ b/include/gc/Target/LLVM/GEN/Target.h @@ -1,6 +1,6 @@ -//===- Target.h - MLIR GEN target registration ------------------*- C++ -*-===// +//===-- Target.h - MLIR GEN target registration -----------------*- C++ -*-===// // -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // diff --git a/include/gc/Target/LLVM/GEN/Utils.h b/include/gc/Target/LLVM/GEN/Utils.h index 22854b715..5017dcf67 100644 --- a/include/gc/Target/LLVM/GEN/Utils.h +++ b/include/gc/Target/LLVM/GEN/Utils.h @@ -1,6 +1,6 @@ -//===- Utils.h - MLIR GEN target utils --------------------------*- C++ -*-===// +//===-- Utils.h - MLIR GEN target utils -------------------------*- C++ -*-===// // -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // diff --git a/include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h b/include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h index 701199cf9..ac0a1f134 100644 --- a/include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h +++ b/include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h @@ -1,6 +1,6 @@ -//===- GENToLLVMIRTranslation.h - GEN to LLVM IR ----------------*- C++ -*-===// +//===-- GENToLLVMIRTranslation.h - GEN to LLVM IR ---------------*- C++ -*-===// // -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // diff --git a/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp b/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp index aa032e89b..326d561e4 100644 --- a/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp +++ b/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp @@ -1,3 +1,10 @@ +//===-- GENDialect.cpp - GEN Attrs and dialect registration -----*- C++ -*-===// +// +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// #include "gc/Dialect/LLVMIR/GENDialect.h" #include "mlir/Dialect/GPU/IR/CompilationInterfaces.h" diff --git a/lib/gc/Target/LLVM/GEN/Target.cpp b/lib/gc/Target/LLVM/GEN/Target.cpp index 93bcc331b..afa172ac2 100644 --- a/lib/gc/Target/LLVM/GEN/Target.cpp +++ b/lib/gc/Target/LLVM/GEN/Target.cpp @@ -1,6 +1,6 @@ -//===- Target.cpp - MLIR LLVM GEN target compilation ------------*- C++ -*-===// +//===-- Target.cpp - MLIR LLVM GEN target compilation -----------*- C++ -*-===// // -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // diff --git a/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp b/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp index b84076385..6a5466455 100644 --- a/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp +++ b/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp @@ -1,6 +1,6 @@ -//===- GENToLLVMIRTranslation.cpp - Translate GEN to LLVM IR --------------===// +//===-- GENToLLVMIRTranslation.cpp - Translate GEN to LLVM IR ---*- C++ -*-===// // -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // diff --git a/lib/gc/Transforms/GPU/GPUOpsLowering.h b/lib/gc/Transforms/GPU/GPUOpsLowering.h index 92e69badc..23638981c 100644 --- a/lib/gc/Transforms/GPU/GPUOpsLowering.h +++ b/lib/gc/Transforms/GPU/GPUOpsLowering.h @@ -1,6 +1,6 @@ -//===- GPUOpsLowering.h - GPU FuncOp / ReturnOp lowering -------*- C++ -*--===// +//===- GPUOpsLowering.h - GPU FuncOp / ReturnOp lowering --------*- C++ -*-===// // -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// This file is licensed under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // From 201327057838952cda23488273da37c555761d07 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Fri, 2 Aug 2024 06:47:33 -0700 Subject: [PATCH 12/19] Move xegpu pass to imex-only build --- .gitignore | 1 + cmake/functions.cmake | 11 +++++++-- lib/gc/Dialect/LLVMIR/CMakeLists.txt | 1 + lib/gc/Target/LLVM/CMakeLists.txt | 3 ++- .../Target/LLVMIR/Dialect/GEN/CMakeLists.txt | 3 ++- lib/gc/Transforms/CMakeLists.txt | 4 +--- lib/gc/Transforms/GPU/CMakeLists.txt | 24 ++++++++++++++++++- src/dnnl/CMakeLists.txt | 3 +++ src/gc-opt/CMakeLists.txt | 14 ++++++----- 9 files changed, 50 insertions(+), 14 deletions(-) diff --git a/.gitignore b/.gitignore index e1fe789da..4edd7d80c 100644 --- a/.gitignore +++ b/.gitignore @@ -3,3 +3,4 @@ build/ externals/ compile_commands.json +install/ diff --git a/cmake/functions.cmake b/cmake/functions.cmake index cbd173e75..16cab3e15 100644 --- a/cmake/functions.cmake +++ b/cmake/functions.cmake @@ -111,10 +111,17 @@ endfunction() function(gc_add_mlir_dialect_library name) add_mlir_dialect_library(${ARGV}) - target_link_libraries(obj.${name} PUBLIC GcInterface) set_property(GLOBAL APPEND PROPERTY GC_DIALECT_LIBS ${name}) if(GcInterface IN_LIST ARGN) target_link_libraries(obj.${name} PUBLIC GcInterface) endif() -endfunction() \ No newline at end of file +endfunction() + +function(gc_add_mlir_translation_library name) + add_mlir_translation_library(${ARGV}) + set_property(GLOBAL APPEND PROPERTY GC_MLIR_LIBS ${name}) + if(GcInterface IN_LIST ARGN) + target_link_libraries(obj.${name} PUBLIC GcInterface) + endif() +endfunction() diff --git a/lib/gc/Dialect/LLVMIR/CMakeLists.txt b/lib/gc/Dialect/LLVMIR/CMakeLists.txt index 05ec2078a..6a8d55641 100644 --- a/lib/gc/Dialect/LLVMIR/CMakeLists.txt +++ b/lib/gc/Dialect/LLVMIR/CMakeLists.txt @@ -16,5 +16,6 @@ add_mlir_dialect_library(MLIRGENDialect MLIRIR MLIRLLVMDialect MLIRSideEffectInterfaces + GcInterface ) set_property(GLOBAL APPEND PROPERTY GC_DIALECT_LIBS MLIRGENDialect) diff --git a/lib/gc/Target/LLVM/CMakeLists.txt b/lib/gc/Target/LLVM/CMakeLists.txt index bd5475919..e24628a7a 100644 --- a/lib/gc/Target/LLVM/CMakeLists.txt +++ b/lib/gc/Target/LLVM/CMakeLists.txt @@ -1,4 +1,4 @@ -add_mlir_dialect_library(MLIRGENTarget +gc_add_mlir_dialect_library(MLIRGENTarget GEN/Target.cpp OBJECT @@ -14,4 +14,5 @@ add_mlir_dialect_library(MLIRGENTarget MLIRGPUDialect MLIRTargetLLVM MLIRGENToLLVMIRTranslation + GcInterface ) diff --git a/lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt b/lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt index e84dc22b9..40630306d 100644 --- a/lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt +++ b/lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt @@ -1,4 +1,4 @@ -add_mlir_translation_library(MLIRGENToLLVMIRTranslation +gc_add_mlir_translation_library(MLIRGENToLLVMIRTranslation GENToLLVMIRTranslation.cpp DEPENDS @@ -13,4 +13,5 @@ add_mlir_translation_library(MLIRGENToLLVMIRTranslation MLIRGENDialect MLIRSupport MLIRTargetLLVMIRExport + GcInterface ) diff --git a/lib/gc/Transforms/CMakeLists.txt b/lib/gc/Transforms/CMakeLists.txt index 08ae24143..ab86311fe 100644 --- a/lib/gc/Transforms/CMakeLists.txt +++ b/lib/gc/Transforms/CMakeLists.txt @@ -25,6 +25,4 @@ gc_add_mlir_library(GcPasses GcInterface ) -if(GC_ENABLE_IMEX) - add_subdirectory(GPU) -endif() +add_subdirectory(GPU) diff --git a/lib/gc/Transforms/GPU/CMakeLists.txt b/lib/gc/Transforms/GPU/CMakeLists.txt index 9b3bf5d77..454cd4c92 100644 --- a/lib/gc/Transforms/GPU/CMakeLists.txt +++ b/lib/gc/Transforms/GPU/CMakeLists.txt @@ -1,9 +1,9 @@ gc_add_mlir_library(GcGpuPasses - LinalgToXeGPU.cpp GPULegalizeModule.cpp ConvertGpuSignaturesToLLVM.cpp GPUAttachGenTarget.cpp + PARTIAL_SOURCES_INTENDED DEPENDS GraphCompilerPassIncGen @@ -23,3 +23,25 @@ gc_add_mlir_library(GcGpuPasses GcUtilsIR ) +if(GC_ENABLE_IMEX) +gc_add_mlir_library(GcIMEXPasses + LinalgToXeGPU.cpp + + PARTIAL_SOURCES_INTENDED + DEPENDS + GraphCompilerPassIncGen + + LINK_LIBS PUBLIC + MLIRGPUDialect + MLIRXeGPUDialect + MLIRGPUTransforms + MLIRGPUToSPIRV + MLIRSCFToGPU + MLIRSCFToSPIRV + MLIRMathToSPIRV + MLIRControlFlowToSPIRV + MLIRMemRefTransforms + GcInterface + GcUtilsIR +) +endif() diff --git a/src/dnnl/CMakeLists.txt b/src/dnnl/CMakeLists.txt index 69e8cb29f..cea142dc2 100644 --- a/src/dnnl/CMakeLists.txt +++ b/src/dnnl/CMakeLists.txt @@ -24,9 +24,12 @@ set(GC_DNNL_SOURCES JsonParser.cpp dnnl_graph_compiler.cpp ) +# todo: replace with a gpu rutnime library once we have an appropriate target +get_property(gc_passes_libs GLOBAL PROPERTY GC_PASS_LIBS) set(GC_DNNL_LINKED_LIBS GcJitWrapper GcCpuRuntime + ${gc_passes_libs} ) gc_add_mlir_library(GcDnnl SHARED ${GC_DNNL_SOURCES} diff --git a/src/gc-opt/CMakeLists.txt b/src/gc-opt/CMakeLists.txt index 1e5f4e38f..54d7bf9b5 100644 --- a/src/gc-opt/CMakeLists.txt +++ b/src/gc-opt/CMakeLists.txt @@ -20,23 +20,24 @@ if(NOT GC_ENABLE_OPT) return() endif() +# todo: this needs further cleanup if(GC_DEV_LINK_LLVM_DYLIB) set(MLIR_LINK_COMPONENTS MLIR ) get_property(dialect_libs GLOBAL PROPERTY GC_DIALECT_LIBS) - get_property(conversion_libs GLOBAL PROPERTY GC_PASS_LIBS) - get_property(extension_libs GLOBAL PROPERTY GC_EXTENSION_LIBS) else() set(MLIR_LINK_COMPONENTS MLIROptLib MLIRToLLVMIRTranslationRegistration ) get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS) - get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS) - get_property(extension_libs GLOBAL PROPERTY MLIR_EXTENSION_LIBS) endif() +get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS) +get_property(extension_libs GLOBAL PROPERTY MLIR_EXTENSION_LIBS) +get_property(gc_passes_libs GLOBAL PROPERTY GC_PASS_LIBS) + add_llvm_executable(gc-opt gc-opt.cpp) llvm_update_compile_flags(gc-opt) mlir_check_all_link_libraries(gc-opt) @@ -45,8 +46,9 @@ target_link_libraries(gc-opt PUBLIC GcInterface) target_link_libraries(gc-opt PRIVATE ${dialect_libs} ${conversion_libs} + ${extension_libs} ${MLIR_LINK_COMPONENTS} - GcPasses + ${gc_passes_libs} ) if(GC_ENABLE_IMEX) @@ -55,7 +57,7 @@ if(GC_ENABLE_IMEX) get_property(IMEX_INCLUDES GLOBAL PROPERTY IMEX_INCLUDES) target_include_directories(gc-opt PRIVATE ${IMEX_INCLUDES}) target_link_libraries(gc-opt PRIVATE - GcGpuPasses + GcGpuIMEXPasses IMEXGPUXDialect IMEXXeTileDialect IMEXRegionDialect From 3f7a71a01468964e01d27b0b9e699d2e333e392e Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Fri, 2 Aug 2024 09:01:00 -0700 Subject: [PATCH 13/19] Fix python CAPI linkage --- python/CMakeLists.txt | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/python/CMakeLists.txt b/python/CMakeLists.txt index 68d522d26..bf88854d0 100644 --- a/python/CMakeLists.txt +++ b/python/CMakeLists.txt @@ -99,7 +99,9 @@ add_mlir_python_common_capi_library(GcPythonCAPI MLIRPythonExtension.RegisterEverything MLIRPythonSources.Core ) -target_link_libraries(GcPythonCAPI PUBLIC GcInterface) +# todo: replace with a gpu rutnime library once we have an appropriate target +get_property(gc_passes_libs GLOBAL PROPERTY GC_PASS_LIBS) +target_link_libraries(GcPythonCAPI PUBLIC GcInterface ${gc_passes_libs}) ################################################################################ # Instantiation of all Python modules From ec17d13342043cbc9cf7e9e10df2ccb1fb13498a Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Fri, 2 Aug 2024 09:02:24 -0700 Subject: [PATCH 14/19] fixup! Fix python CAPI linkage --- .gitignore | 1 - 1 file changed, 1 deletion(-) diff --git a/.gitignore b/.gitignore index 4edd7d80c..e1fe789da 100644 --- a/.gitignore +++ b/.gitignore @@ -3,4 +3,3 @@ build/ externals/ compile_commands.json -install/ From ce9e66f60f1b588b1b8cdc7f5cbbd707013fc6f8 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Fri, 2 Aug 2024 09:09:59 -0700 Subject: [PATCH 15/19] Fix merge issues --- lib/gc/Dialect/LLVMIR/CMakeLists.txt | 5 ++--- src/gc-opt/gc-opt.cpp | 2 +- 2 files changed, 3 insertions(+), 4 deletions(-) diff --git a/lib/gc/Dialect/LLVMIR/CMakeLists.txt b/lib/gc/Dialect/LLVMIR/CMakeLists.txt index 6a8d55641..5c46f5d99 100644 --- a/lib/gc/Dialect/LLVMIR/CMakeLists.txt +++ b/lib/gc/Dialect/LLVMIR/CMakeLists.txt @@ -1,4 +1,4 @@ -add_mlir_dialect_library(MLIRGENDialect +gc_add_mlir_dialect_library(MLIRGENDialect IR/GENDialect.cpp ADDITIONAL_HEADER_DIRS @@ -17,5 +17,4 @@ add_mlir_dialect_library(MLIRGENDialect MLIRLLVMDialect MLIRSideEffectInterfaces GcInterface - ) -set_property(GLOBAL APPEND PROPERTY GC_DIALECT_LIBS MLIRGENDialect) +) diff --git a/src/gc-opt/gc-opt.cpp b/src/gc-opt/gc-opt.cpp index 5c196762c..4f046f4ae 100644 --- a/src/gc-opt/gc-opt.cpp +++ b/src/gc-opt/gc-opt.cpp @@ -75,7 +75,7 @@ int main(int argc, char *argv[]) { mlir::registerAllToLLVMIRTranslations(registry); mlir::gen::registerGenTargetInterfaceExternalModels(registry); mlir::registerGENDialectTranslation(registry); -#ifdef GC_USE_GPU +#ifdef GC_USE_IMEX registry.insert<::imex::xetile::XeTileDialect, ::imex::gpux::GPUXDialect>(); #endif mlir::cpuruntime::registerConvertCPURuntimeToLLVMInterface(registry); From 09e8321a1ac7f80e13dec72989250599ecfe08b1 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Fri, 2 Aug 2024 09:49:04 -0700 Subject: [PATCH 16/19] Fix GCExecutionEngineTests linkage --- test/mlir/unittests/ExecutionEngine/CMakeLists.txt | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/test/mlir/unittests/ExecutionEngine/CMakeLists.txt b/test/mlir/unittests/ExecutionEngine/CMakeLists.txt index 2cfe3f77e..a063bbb3b 100644 --- a/test/mlir/unittests/ExecutionEngine/CMakeLists.txt +++ b/test/mlir/unittests/ExecutionEngine/CMakeLists.txt @@ -1,7 +1,10 @@ add_mlir_unittest(GCExecutionEngineTests JitWrapper.cpp ) +# todo: remove once we have a gpu runtime library +get_property(gc_passes_libs GLOBAL PROPERTY GC_PASS_LIBS) target_link_libraries(GCExecutionEngineTests PRIVATE GcJitWrapper - GcCpuRuntime) + GcCpuRuntime + ${gc_passes_libs}) From c51053cad1d818fde8d59190278d0bf61ac22f61 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Mon, 5 Aug 2024 02:09:34 -0700 Subject: [PATCH 17/19] Add SPIRVCodeGen dependency to the gen target --- lib/gc/Target/LLVM/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/lib/gc/Target/LLVM/CMakeLists.txt b/lib/gc/Target/LLVM/CMakeLists.txt index e24628a7a..659898cb7 100644 --- a/lib/gc/Target/LLVM/CMakeLists.txt +++ b/lib/gc/Target/LLVM/CMakeLists.txt @@ -13,6 +13,7 @@ gc_add_mlir_dialect_library(MLIRGENTarget MLIRSupport MLIRGPUDialect MLIRTargetLLVM + SPIRVCodeGen MLIRGENToLLVMIRTranslation GcInterface ) From eefddb6eedc6e6a8beb49b7279fa1bf865ede374 Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Mon, 5 Aug 2024 02:42:43 -0700 Subject: [PATCH 18/19] Revert "Add SPIRVCodeGen dependency to the gen target" This reverts commit a31238eb3bd68757ce9fc99631f12935dbac0650. --- lib/gc/Target/LLVM/CMakeLists.txt | 1 - 1 file changed, 1 deletion(-) diff --git a/lib/gc/Target/LLVM/CMakeLists.txt b/lib/gc/Target/LLVM/CMakeLists.txt index 659898cb7..e24628a7a 100644 --- a/lib/gc/Target/LLVM/CMakeLists.txt +++ b/lib/gc/Target/LLVM/CMakeLists.txt @@ -13,7 +13,6 @@ gc_add_mlir_dialect_library(MLIRGENTarget MLIRSupport MLIRGPUDialect MLIRTargetLLVM - SPIRVCodeGen MLIRGENToLLVMIRTranslation GcInterface ) From c2b485e03d30dc07ace109617fbba704c50bcb2e Mon Sep 17 00:00:00 2001 From: Petr Kurapov Date: Mon, 5 Aug 2024 02:52:57 -0700 Subject: [PATCH 19/19] Add LLVMSPIRVCodeGen dependency --- lib/gc/Target/LLVM/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/lib/gc/Target/LLVM/CMakeLists.txt b/lib/gc/Target/LLVM/CMakeLists.txt index e24628a7a..fcc2a06b8 100644 --- a/lib/gc/Target/LLVM/CMakeLists.txt +++ b/lib/gc/Target/LLVM/CMakeLists.txt @@ -13,6 +13,7 @@ gc_add_mlir_dialect_library(MLIRGENTarget MLIRSupport MLIRGPUDialect MLIRTargetLLVM + LLVMSPIRVCodeGen MLIRGENToLLVMIRTranslation GcInterface )