-
Notifications
You must be signed in to change notification settings - Fork 2.8k
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[Mosaic GPU] Load LLVM lowering interfaces for all dialects
Apparently we were missing interface registration code for LLVM lowering, which the gpu-to-llvm pass gracefully ignores unless compiled with debug assertions enabled. But, simply adding the assertions in fact makes the pass _too powerful_ and makes it lower _all dialects to LLVM_, which is not what we want. That's why I've replaced it with a minimal version that is only repsponsible for handling the GPU dialect, making the lowering similar to the one prior to extra registrations. PiperOrigin-RevId: 641874183
- Loading branch information
Showing
7 changed files
with
217 additions
and
6 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,64 @@ | ||
| /* Copyright 2024 The JAX Authors. | ||
| Licensed under the Apache License, Version 2.0 (the "License"); | ||
| you may not use this file except in compliance with the License. | ||
| You may obtain a copy of the License at | ||
| http://www.apache.org/licenses/LICENSE-2.0 | ||
| Unless required by applicable law or agreed to in writing, software | ||
| distributed under the License is distributed on an "AS IS" BASIS, | ||
| WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| See the License for the specific language governing permissions and | ||
| limitations under the License. | ||
| ==============================================================================*/ | ||
|
|
||
| #ifndef JAXLIB_MOSAIC_GPU_PASS_BOILERPLATE_H_ | ||
| #define JAXLIB_MOSAIC_GPU_PASS_BOILERPLATE_H_ | ||
|
|
||
| #include "mlir/include/mlir/IR/DialectRegistry.h" | ||
| #include "mlir/include/mlir/Pass/Pass.h" | ||
| #include "mlir/include/mlir/Support/LLVM.h" | ||
| #include "mlir/include/mlir/Support/TypeID.h" | ||
| namespace mosaic { | ||
| namespace gpu { | ||
|
|
||
| template <typename Derived, typename Op = void> | ||
| class Pass : public ::mlir::OperationPass<Op> { | ||
| public: | ||
| Pass() : ::mlir::OperationPass<Op>(::mlir::TypeID::get<Derived>()) {} | ||
| Pass(const Pass &other) : ::mlir::OperationPass<Op>(other) {} | ||
| Pass &operator=(const Pass &) = delete; | ||
| Pass(Pass &&) = delete; | ||
| Pass &operator=(Pass &&) = delete; | ||
| ~Pass() = default; | ||
|
|
||
| static constexpr ::llvm::StringLiteral getArgumentName() { | ||
| return ::llvm::StringLiteral(Derived::kArgumentName); | ||
| } | ||
| ::llvm::StringRef getArgument() const override { return getArgumentName(); } | ||
| ::llvm::StringRef getDescription() const override { return ""; } | ||
| static constexpr ::llvm::StringLiteral getPassName() { | ||
| return ::llvm::StringLiteral(Derived::kPassName); | ||
| } | ||
| ::llvm::StringRef getName() const override { return getPassName(); } | ||
| static bool classof(const ::mlir::Pass *pass) { | ||
| return pass->getTypeID() == ::mlir::TypeID::get<Derived>(); | ||
| } | ||
| std::unique_ptr<::mlir::Pass> clonePass() const override { | ||
| return std::make_unique<Derived>(*static_cast<const Derived *>(this)); | ||
| } | ||
| void getDependentDialects(::mlir::DialectRegistry ®istry) const override {} | ||
|
|
||
| private: | ||
| using This = | ||
| Pass<Derived, Op>; // Can't have a comma in the macro instantiation | ||
|
|
||
| public: | ||
| MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(This) | ||
| }; | ||
|
|
||
| } // namespace gpu | ||
| } // namespace mosaic | ||
|
|
||
| #endif // JAXLIB_MOSAIC_GPU_PASS_BOILERPLATE_H_ |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,81 @@ | ||
| /* Copyright 2024 The JAX Authors. | ||
| Licensed under the Apache License, Version 2.0 (the "License"); | ||
| you may not use this file except in compliance with the License. | ||
| You may obtain a copy of the License at | ||
| http://www.apache.org/licenses/LICENSE-2.0 | ||
| Unless required by applicable law or agreed to in writing, software | ||
| distributed under the License is distributed on an "AS IS" BASIS, | ||
| WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| See the License for the specific language governing permissions and | ||
| limitations under the License. | ||
| ==============================================================================*/ | ||
|
|
||
| #include "jaxlib/mosaic/gpu/passes.h" | ||
| #include <memory> | ||
| #include <utility> | ||
| #include <vector> | ||
|
|
||
| #include "llvm/include/llvm/ADT/StringRef.h" | ||
| #include "llvm/include/llvm/Support/Debug.h" | ||
| #include "mlir/include/mlir/Conversion/GPUCommon/GPUCommonPass.h" | ||
| #include "mlir/include/mlir/Conversion/LLVMCommon/TypeConverter.h" | ||
| #include "mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h" | ||
| #include "mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h" | ||
| #include "mlir/include/mlir/IR/BuiltinOps.h" | ||
| #include "mlir/include/mlir/IR/SymbolTable.h" | ||
| #include "mlir/include/mlir/Pass/PassRegistry.h" | ||
| #include "mlir/include/mlir/Support/LLVM.h" | ||
| #include "mlir/include/mlir/Transforms/DialectConversion.h" | ||
| #include "jaxlib/mosaic/gpu/pass_boilerplate.h" | ||
|
|
||
| namespace mosaic { | ||
| namespace gpu { | ||
|
|
||
| namespace { | ||
|
|
||
| class ConvertGpuToLLVMPass | ||
| : public mosaic::gpu::Pass<ConvertGpuToLLVMPass, mlir::ModuleOp> { | ||
| public: | ||
| using mosaic::gpu::Pass<ConvertGpuToLLVMPass, mlir::ModuleOp>::Pass; | ||
| static constexpr llvm::StringLiteral kArgumentName = | ||
| "mosaic-convert-gpu-to-llvm"; | ||
| static constexpr llvm::StringLiteral kPassName = "ConvertGpuToLLVMPass"; | ||
|
|
||
| void runOnOperation() override { | ||
| llvm::DebugFlag = true; | ||
| mlir::MLIRContext *ctx = &getContext(); | ||
| mlir::RewritePatternSet patterns(ctx); | ||
| mlir::LLVMTypeConverter converter(ctx); | ||
| mlir::ConversionTarget target(*ctx); | ||
| target.addLegalDialect<mlir::LLVM::LLVMDialect>(); | ||
| target.addLegalOp<mlir::gpu::GPUModuleOp>(); | ||
| target.addDynamicallyLegalOp<mlir::gpu::LaunchFuncOp>( | ||
| [&](mlir::gpu::LaunchFuncOp op) -> bool { | ||
| return converter.isLegal(op->getOperandTypes()) && | ||
| converter.isLegal(op->getResultTypes()); | ||
| }); | ||
| auto symtab = mlir::SymbolTable(getOperation()); | ||
| mlir::populateGpuToLLVMConversionPatterns(converter, patterns, "gpu.binary", | ||
| false, &symtab); | ||
| if (mlir::applyPartialConversion(getOperation(), target, | ||
| std::move(patterns)) | ||
| .failed()) { | ||
| signalPassFailure(); | ||
| } | ||
| llvm::DebugFlag = false; | ||
| } | ||
| }; | ||
|
|
||
| } // namespace | ||
|
|
||
| void registerConvertGpuToLLVMPass() { | ||
| ::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> { | ||
| return std::make_unique<ConvertGpuToLLVMPass>(); | ||
| }); | ||
| } | ||
|
|
||
| } // namespace gpu | ||
| } // namespace mosaic |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,27 @@ | ||
| /* Copyright 2024 The JAX Authors. | ||
| Licensed under the Apache License, Version 2.0 (the "License"); | ||
| you may not use this file except in compliance with the License. | ||
| You may obtain a copy of the License at | ||
| http://www.apache.org/licenses/LICENSE-2.0 | ||
| Unless required by applicable law or agreed to in writing, software | ||
| distributed under the License is distributed on an "AS IS" BASIS, | ||
| WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| See the License for the specific language governing permissions and | ||
| limitations under the License. | ||
| ==============================================================================*/ | ||
|
|
||
| #ifndef JAXLIB_MOSAIC_GPU_CONVERT_GPU_TO_LLVM_H_ | ||
| #define JAXLIB_MOSAIC_GPU_CONVERT_GPU_TO_LLVM_H_ | ||
|
|
||
| namespace mosaic { | ||
| namespace gpu { | ||
|
|
||
| void registerConvertGpuToLLVMPass(); | ||
|
|
||
| } // namespace gpu | ||
| } // namespace mosaic | ||
|
|
||
| #endif // JAXLIB_MOSAIC_GPU_CONVERT_GPU_TO_LLVM_H_ |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters