[Mlir-commits] [clang] [flang] [mlir] [mlir][OpenACC] Decouple frontend support from GPU operations (PR #233103)
Mehdi Amini
llvmlistbot at llvm.org
Sat Oct 17 08:43:25 PDT 2026
https://github.com/joker-eph created https://github.com/llvm/llvm-project/pull/233103
Allow clients such as CIR to use OpenACC frontend IR without linking the full GPU operation implementation and its Math dependencies.
Provide MLIRGPUBase for GPU representations and expose OpenACC's GPU, LLVM and MemRef external interface models through separate extension libraries. Keep GPU type and attribute ownership and parsing in the GPU dialect, and retain all OpenACC operations in MLIROpenACCDialect.
Register the extensions through MLIR's default extension registration, OpenACC passes and Flang. Preserve the C API's automatic registration and CIR's pointer models, and explicitly load Arith for CIR's OpenACC lowering.
Save >4MB in clang+CIR build.
Assisted-by: Codex
>From 456126e12399c8bb3ebdcd8e403ad7e134e0240e Mon Sep 17 00:00:00 2001
From: Mehdi Amini <joker.eph at gmail.com>
Date: Sat, 17 Oct 2026 03:03:35 -0700
Subject: [PATCH] [mlir][OpenACC] Decouple frontend support from GPU operations
Allow clients such as CIR to use OpenACC frontend IR without linking the
full GPU operation implementation and its Math dependencies.
Provide MLIRGPUBase for GPU representations and expose OpenACC's GPU,
LLVM and MemRef external interface models through separate extension
libraries. Keep GPU type and attribute ownership and parsing in the GPU
dialect, and retain all OpenACC operations in MLIROpenACCDialect.
Register the extensions through MLIR's default extension registration,
OpenACC passes and Flang. Preserve the C API's automatic registration
and CIR's pointer models, and explicitly load Arith for CIR's OpenACC
lowering.
Assisted-by: Codex
---
clang/lib/CIR/CMakeLists.txt | 1 +
clang/lib/CIR/CodeGen/CIRGenerator.cpp | 5 +-
clang/lib/CIR/RegisterAllDialects.cpp | 8 +-
.../Optimizer/OpenACC/Support/CMakeLists.txt | 2 +
.../Support/RegisterOpenACCExtensions.cpp | 2 +
mlir/docs/Dialects/OpenACCDialect.md | 31 ++
mlir/include/mlir/Dialect/GPU/IR/GPUBase.h | 209 +++++++
mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h | 180 +------
.../OpenACC/Extensions/AllExtensions.h | 22 +
.../OpenACC/Extensions/GPUInterfaces.h | 20 +
.../OpenACC/Extensions/LLVMInterfaces.h | 21 +
.../OpenACC/Extensions/MemRefInterfaces.h | 21 +
mlir/include/mlir/Dialect/OpenACC/OpenACC.h | 4 +-
.../mlir/Dialect/OpenACC/OpenACCBase.td | 3 -
mlir/lib/CAPI/Dialect/CMakeLists.txt | 2 +
mlir/lib/CAPI/Dialect/OpenACC.cpp | 36 +-
mlir/lib/Dialect/GPU/CMakeLists.txt | 24 +
mlir/lib/Dialect/GPU/IR/GPUBase.cpp | 509 ++++++++++++++++++
mlir/lib/Dialect/GPU/IR/GPUDialect.cpp | 500 +----------------
mlir/lib/Dialect/OpenACC/CMakeLists.txt | 1 +
.../OpenACC/Extensions/AllExtensions.cpp | 19 +
.../Dialect/OpenACC/Extensions/CMakeLists.txt | 41 ++
.../OpenACC/Extensions/GPUInterfaces.cpp | 32 ++
.../OpenACC/Extensions/LLVMInterfaces.cpp | 98 ++++
.../OpenACC/Extensions/MemRefInterfaces.cpp | 328 +++++++++++
mlir/lib/Dialect/OpenACC/IR/CMakeLists.txt | 4 +-
mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp | 394 --------------
mlir/lib/Dialect/OpenACC/IR/OpenACCCG.cpp | 7 +-
.../OpenACC/Transforms/ACCBindRoutine.cpp | 9 +-
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 7 +-
.../OpenACC/Transforms/ACCComputeLowering.cpp | 8 +-
.../ACCDeclareGPUModuleInsertion.cpp | 9 +-
.../OpenACC/Transforms/ACCEmitRemarksData.cpp | 8 +-
.../OpenACC/Transforms/ACCEmitRemarksLoop.cpp | 8 +-
.../Transforms/ACCEmitRemarksPrivate.cpp | 9 +-
.../Transforms/ACCIfClauseLowering.cpp | 9 +-
.../OpenACC/Transforms/ACCImplicitData.cpp | 9 +-
.../OpenACC/Transforms/ACCImplicitDeclare.cpp | 8 +-
.../OpenACC/Transforms/ACCImplicitRoutine.cpp | 8 +-
.../OpenACC/Transforms/ACCLegalizeSerial.cpp | 7 +-
.../OpenACC/Transforms/ACCLoopTiling.cpp | 7 +-
.../Transforms/ACCRecipeMaterialization.cpp | 9 +-
.../OpenACC/Transforms/ACCRoutineLowering.cpp | 8 +-
.../Transforms/ACCRoutineToGPUFunc.cpp | 9 +-
.../Transforms/ACCSpecializeForDevice.cpp | 9 +-
.../Transforms/ACCSpecializeForHost.cpp | 8 +-
.../Dialect/OpenACC/Transforms/CMakeLists.txt | 1 +
.../OpenACC/Transforms/LegalizeDataValues.cpp | 9 +-
.../OffloadLiveInValueCanonicalization.cpp | 9 +-
.../Transforms/OffloadTargetVerifier.cpp | 8 +-
mlir/lib/RegisterAllExtensions.cpp | 2 +
mlir/test/CAPI/CMakeLists.txt | 8 +
mlir/test/CAPI/openacc.c | 55 ++
mlir/test/CMakeLists.txt | 1 +
mlir/unittests/Dialect/OpenACC/CMakeLists.txt | 5 +
.../Dialect/OpenACC/OpenACCCGOpsTest.cpp | 4 +
.../Dialect/OpenACC/OpenACCExtensionsTest.cpp | 96 ++++
.../OpenACC/OpenACCOpsInterfacesTest.cpp | 4 +
.../Dialect/OpenACC/OpenACCOpsTest.cpp | 4 +
.../OpenACC/OpenACCTypeInterfacesTest.cpp | 6 +-
.../Dialect/OpenACC/OpenACCUtilsCGTest.cpp | 4 +
.../Dialect/OpenACC/OpenACCUtilsGPUTest.cpp | 4 +
.../Dialect/OpenACC/OpenACCUtilsLoopTest.cpp | 4 +
.../OpenACC/OpenACCUtilsReductionTest.cpp | 4 +
.../Dialect/OpenACC/OpenACCUtilsTest.cpp | 4 +
.../OpenACC/OpenACCUtilsTilingTest.cpp | 4 +
.../Dialect/OpenACC/OpenACCUtilsTypeTest.cpp | 4 +
.../Dialect/OpenACC/Standalone/CMakeLists.txt | 7 +
.../Standalone/OpenACCStandaloneTest.cpp | 36 ++
69 files changed, 1842 insertions(+), 1114 deletions(-)
create mode 100644 mlir/include/mlir/Dialect/GPU/IR/GPUBase.h
create mode 100644 mlir/include/mlir/Dialect/OpenACC/Extensions/AllExtensions.h
create mode 100644 mlir/include/mlir/Dialect/OpenACC/Extensions/GPUInterfaces.h
create mode 100644 mlir/include/mlir/Dialect/OpenACC/Extensions/LLVMInterfaces.h
create mode 100644 mlir/include/mlir/Dialect/OpenACC/Extensions/MemRefInterfaces.h
create mode 100644 mlir/lib/Dialect/GPU/IR/GPUBase.cpp
create mode 100644 mlir/lib/Dialect/OpenACC/Extensions/AllExtensions.cpp
create mode 100644 mlir/lib/Dialect/OpenACC/Extensions/CMakeLists.txt
create mode 100644 mlir/lib/Dialect/OpenACC/Extensions/GPUInterfaces.cpp
create mode 100644 mlir/lib/Dialect/OpenACC/Extensions/LLVMInterfaces.cpp
create mode 100644 mlir/lib/Dialect/OpenACC/Extensions/MemRefInterfaces.cpp
create mode 100644 mlir/test/CAPI/openacc.c
create mode 100644 mlir/unittests/Dialect/OpenACC/OpenACCExtensionsTest.cpp
create mode 100644 mlir/unittests/Dialect/OpenACC/Standalone/CMakeLists.txt
create mode 100644 mlir/unittests/Dialect/OpenACC/Standalone/OpenACCStandaloneTest.cpp
diff --git a/clang/lib/CIR/CMakeLists.txt b/clang/lib/CIR/CMakeLists.txt
index 6f9f4bb941aadb..c63c84afc8fa31 100644
--- a/clang/lib/CIR/CMakeLists.txt
+++ b/clang/lib/CIR/CMakeLists.txt
@@ -14,6 +14,7 @@ add_clang_library(CIRRegisterAllDialects
MLIRCIR
CIROpenACCSupport
CIROpenMPSupport
+ MLIRArithDialect
MLIRDLTIDialect
MLIRLLVMDialect
)
diff --git a/clang/lib/CIR/CodeGen/CIRGenerator.cpp b/clang/lib/CIR/CodeGen/CIRGenerator.cpp
index bc0aed0b01ff03..d2f1b54e43325c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenerator.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenerator.cpp
@@ -12,6 +12,7 @@
#include "CIRGenModule.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
@@ -54,8 +55,10 @@ void CIRGenerator::Initialize(ASTContext &astContext) {
mlirContext->loadDialect<mlir::DLTIDialect, mlir::LLVM::LLVMDialect,
cir::CIRDialect>();
const LangOptions &langOpts = astContext.getLangOpts();
+ // OpenACC lowering emits Arith constants for bounds and clause operands.
if (langOpts.OpenACC)
- mlirContext->getOrLoadDialect<mlir::acc::OpenACCDialect>();
+ mlirContext
+ ->loadDialect<mlir::acc::OpenACCDialect, mlir::arith::ArithDialect>();
if (langOpts.OpenMP)
mlirContext->getOrLoadDialect<mlir::omp::OpenMPDialect>();
diff --git a/clang/lib/CIR/RegisterAllDialects.cpp b/clang/lib/CIR/RegisterAllDialects.cpp
index 73effa12b31921..706635049afbe4 100644
--- a/clang/lib/CIR/RegisterAllDialects.cpp
+++ b/clang/lib/CIR/RegisterAllDialects.cpp
@@ -8,6 +8,7 @@
#include "clang/CIR/InitAllDialects.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
@@ -24,9 +25,10 @@ namespace cir {
void registerAllDialects(mlir::DialectRegistry ®istry) {
// The LLVM dialect is needed to parse CIR: data layout specs use LLVM
// pointer types as keys.
- registry.insert<mlir::BuiltinDialect, cir::CIRDialect, mlir::DLTIDialect,
- mlir::LLVM::LLVMDialect, mlir::omp::OpenMPDialect,
- mlir::acc::OpenACCDialect>();
+ registry
+ .insert<mlir::BuiltinDialect, cir::CIRDialect, mlir::arith::ArithDialect,
+ mlir::DLTIDialect, mlir::LLVM::LLVMDialect,
+ mlir::omp::OpenMPDialect, mlir::acc::OpenACCDialect>();
// Register extensions to integrate CIR types with OpenACC and OpenMP.
cir::omp::registerOpenMPExtensions(registry);
cir::acc::registerOpenACCExtensions(registry);
diff --git a/flang/lib/Optimizer/OpenACC/Support/CMakeLists.txt b/flang/lib/Optimizer/OpenACC/Support/CMakeLists.txt
index 9ff46c70db3cc0..94f7fdf872a860 100644
--- a/flang/lib/Optimizer/OpenACC/Support/CMakeLists.txt
+++ b/flang/lib/Optimizer/OpenACC/Support/CMakeLists.txt
@@ -27,10 +27,12 @@ add_flang_library(FIROpenACCSupport
HLFIRDialect
MLIR_DEPS
+ MLIROpenACCAllExtensions
MLIROpenACCDialect
MLIROpenACCUtils
MLIR_LIBS
+ MLIROpenACCAllExtensions
MLIROpenACCDialect
MLIROpenACCUtils
)
diff --git a/flang/lib/Optimizer/OpenACC/Support/RegisterOpenACCExtensions.cpp b/flang/lib/Optimizer/OpenACC/Support/RegisterOpenACCExtensions.cpp
index 0ffa43ac695858..262c39a72a5b1f 100644
--- a/flang/lib/Optimizer/OpenACC/Support/RegisterOpenACCExtensions.cpp
+++ b/flang/lib/Optimizer/OpenACC/Support/RegisterOpenACCExtensions.cpp
@@ -21,9 +21,11 @@
#include "flang/Optimizer/HLFIR/HLFIROps.h"
#include "flang/Optimizer/OpenACC/Support/FIROpenACCOpsInterfaces.h"
#include "flang/Optimizer/OpenACC/Support/FIROpenACCTypeInterfaces.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
namespace fir::acc {
void registerOpenACCExtensions(mlir::DialectRegistry ®istry) {
+ mlir::acc::registerAllExtensions(registry);
registry.addExtension(+[](mlir::MLIRContext *ctx,
fir::FIROpsDialect *dialect) {
fir::BoxType::attachInterface<OpenACCMappableModel<fir::BaseBoxType>>(*ctx);
diff --git a/mlir/docs/Dialects/OpenACCDialect.md b/mlir/docs/Dialects/OpenACCDialect.md
index 9e85d66ba340fd..f8864f28d508c9 100755
--- a/mlir/docs/Dialects/OpenACCDialect.md
+++ b/mlir/docs/Dialects/OpenACCDialect.md
@@ -19,6 +19,37 @@ operations, types, and attributes.
[TOC]
+## Registration and library dependencies
+
+`MLIROpenACCDialect` provides all OpenACC language and code-generation
+operations, types, and attributes.
+
+OpenACC interfaces for other dialects are provided by extension libraries:
+
+| Library | Registration API | Models |
+| --- | --- | --- |
+| `MLIROpenACCMemRefInterfaces` | `acc::registerMemRefInterfaces` | MemRef pointer and global interfaces |
+| `MLIROpenACCLLVMInterfaces` | `acc::registerLLVMInterfaces` | LLVM pointer and address-of interfaces |
+| `MLIROpenACCGPUInterfaces` | `acc::registerGPUInterfaces` | GPU launch offload-region interface |
+
+To use these models, link the corresponding libraries and register the
+extensions with a `DialectRegistry &`. Their headers are in
+`mlir/Dialect/OpenACC/Extensions`. `acc::registerAllExtensions`, declared in
+`Extensions/AllExtensions.h` and provided by `MLIROpenACCAllExtensions`,
+registers all three extensions.
+
+The MemRef and LLVM extensions load their target dialects and Arith when
+OpenACC is loaded to support IR generation. The MemRef extension attaches
+pointer models to builtin MemRef types. The GPU extension attaches the
+offload-region model to `gpu.launch` when both OpenACC and GPU are loaded,
+regardless of loading order. Extensions can also be registered by appending
+a registry to a context whose dialects are already loaded.
+
+`mlir::registerAllExtensions`, OpenACC passes, and Flang's OpenACC registration
+include these extensions. CIR registers its own pointer models.
+`MLIRCAPIOpenACC` provides a C API dialect handle that registers these
+extensions and loads GPU with OpenACC.
+
## Dialect Design Goals
* Needs to have complete representation of the OpenACC language.
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUBase.h b/mlir/include/mlir/Dialect/GPU/IR/GPUBase.h
new file mode 100644
index 00000000000000..b8ae0c18b8f7d5
--- /dev/null
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUBase.h
@@ -0,0 +1,209 @@
+//===- GPUBase.h - GPU types, attributes, and enums -----------*- 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 declares GPU types, attributes, and enums independently of GPU
+// operations.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_GPU_IR_GPUBASE_H
+#define MLIR_DIALECT_GPU_IR_GPUBASE_H
+
+// These representations can be inspected without linking GPU operations.
+// Constructing or parsing GPU types and attributes still requires loading
+// GPUDialect, which owns and registers them as before.
+// Attribute methods that build IR are provided by MLIRGPUDialect.
+#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
+#include "mlir/Dialect/GPU/IR/GPUDialectDecl.h"
+#include "mlir/Dialect/SCF/IR/DeviceMappingInterface.h"
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/Interfaces/FunctionInterfaces.h"
+
+namespace mlir::gpu {
+namespace detail {
+/// Register storage whose definitions are private to MLIRGPUBase. Attribute
+/// metadata and interfaces are registered by GPUDialect::initialize().
+void registerAttributeStorage(MLIRContext *context);
+} // namespace detail
+
+class AsyncTokenType
+ : public Type::TypeBase<AsyncTokenType, Type, TypeStorage> {
+public:
+ // Used for generic hooks in TypeBase.
+ using Base::Base;
+
+ static constexpr StringLiteral name = "gpu.async_token";
+};
+
+class NamedBarrierType
+ : public Type::TypeBase<NamedBarrierType, Type, TypeStorage> {
+public:
+ using Base::Base;
+
+ static constexpr StringLiteral name = "gpu.named_barrier";
+};
+
+/// MMAMatrixType storage and uniquing. Array is uniqued based on its shape
+/// and type.
+struct MMAMatrixStorageType : public TypeStorage {
+ MMAMatrixStorageType(unsigned numDims, const int64_t *dimShapes,
+ Type elementType, StringRef operand)
+ : dimShapes(dimShapes), numDims(numDims), elementType(elementType),
+ operand(operand) {}
+
+ /// The hash key for uniquing.
+ using KeyTy = std::tuple<ArrayRef<int64_t>, Type, StringRef>;
+ bool operator==(const KeyTy &key) const {
+ return key == KeyTy(getShape(), elementType, operand);
+ }
+
+ /// Construction.
+ static MMAMatrixStorageType *construct(TypeStorageAllocator &allocator,
+ const KeyTy &key) {
+ ArrayRef<int64_t> shape = allocator.copyInto(std::get<0>(key));
+ StringRef operand = allocator.copyInto(std::get<2>(key));
+
+ return new (allocator.allocate<MMAMatrixStorageType>())
+ MMAMatrixStorageType(shape.size(), shape.data(), std::get<1>(key),
+ operand);
+ }
+
+ ArrayRef<int64_t> getShape() const {
+ return ArrayRef<int64_t>(dimShapes, numDims);
+ }
+
+ StringRef getOperand() const { return operand; }
+
+ /// Reference to the shape of the MMA matrix.
+ const int64_t *dimShapes;
+
+ /// Number of dimensions in the MMA matrix.
+ unsigned numDims;
+
+ /// Element type of elements held in the MMA matrix.
+ Type elementType;
+
+ /// MMA operand that this MMAMatrix holds. The general form of operation this
+ /// type supports is given by the equation C += A*B. This field specifies
+ /// which operand in the given equation is held by this type. The valid values
+ /// are "AOp", "BOp" and "COp".
+ StringRef operand;
+};
+
+/// MMAMatrix represents a matrix held by a subgroup for matrix-matrix multiply
+/// accumulate operations. MMAMatrices are taken as direct operands by these
+/// operations and are also produced as results. These matrices are meant to
+/// reside in the registers. A limited number of pointwise operations can be
+/// performed on these matrices, i.e., operations which operate uniformly on
+/// all the elements in the matrix and do not change the order of matrix
+/// elements. The above conditions exist because the layout of matrix elements
+/// inside the matrix is opaque i.e., the elements may be present in the
+/// matrix in any order. The general usage of this type is shown as follows:-
+///
+/// %0 = gpu.subgroup_mma_load_matrix %arg0[%c0, %c0] {leadDimension = 16 :
+/// index} : memref<16x16xf16> -> !gpu.mma_matrix<16x16xf16, "AOp">
+///
+/// The MMAMatrixType describes the shape of the matrix being loaded and the
+/// operand being loaded too. The operand needs to be specified to aid the
+/// lowering of this type to dialects such as NVVM where each workitem may
+/// hold different amount of elements depending on the elementType of the
+/// matrix. For e.g., Each workitem holds 4 vector<2xf16>s for f16 data type
+/// and 8 f32s for f32 data type of MMAMatrix. Some other instances of usage
+/// are:-
+///
+/// %3 = gpu.subgroup_mma_compute %0, %1, %2 :
+/// !gpu.mma_matrix<16x16xf16, "AOp">, !gpu.mma_matrix<16x16xf16, "BOp">
+/// -> !gpu.mma_matrix<16x16xf32, "COp">
+///
+///
+/// gpu.subgroup_mma_store_matrix %3, %arg22[%c0, %c0] {leadDimension = 16
+/// : index}: !gpu.mma_matrix<16x16xf32, "COp">, memref<16x16xf32>
+///
+/// MMA matrices may be memref elements on targets that support storing them,
+/// for example SPIRV can represent arrays of MMA matrices.
+// TODO: consider moving this to ODS.
+class MMAMatrixType
+ : public Type::TypeBase<MMAMatrixType, Type, MMAMatrixStorageType,
+ MemRefElementTypeInterface::Trait> {
+public:
+ using Base::Base;
+
+ static constexpr StringLiteral name = "gpu.mma_matrix";
+
+ /// Get MMAMatrixType and verify construction Invariants.
+ static MMAMatrixType get(ArrayRef<int64_t> shape, Type elementType,
+ StringRef operand);
+
+ /// Get MMAMatrixType at a particular location and verify construction
+ /// Invariants.
+ static MMAMatrixType getChecked(function_ref<InFlightDiagnostic()> emitError,
+ ArrayRef<int64_t> shape, Type elementType,
+ StringRef operand);
+
+ /// Check if a type is valid a MMAMatrixType elementType.
+ static bool isValidElementType(Type elementType);
+
+ /// Verify that shape and elementType are actually allowed for the
+ /// MMAMatrixType.
+ static LogicalResult
+ verifyInvariants(function_ref<InFlightDiagnostic()> emitError,
+ ArrayRef<int64_t> shape, Type elementType,
+ StringRef operand);
+
+ /// Get number of dims.
+ unsigned getNumDims() const;
+
+ /// Get shape of the matrix.
+ ArrayRef<int64_t> getShape() const;
+
+ /// Get elementType of a single element.
+ Type getElementType() const;
+
+ /// The general form of operation this type supports is given by the equation
+ /// C += A*B. This function returns which operand in the given equation is
+ /// held by this type. String returned can be one of"AOp", "BOp" and "COp".
+ StringRef getOperand() const;
+};
+
+// Handle types for sparse.
+enum class SparseHandleKind { SpMat, DnTensor, SpGEMMOp };
+
+class SparseDnTensorHandleType
+ : public Type::TypeBase<SparseDnTensorHandleType, Type, TypeStorage> {
+public:
+ using Base =
+ Type::TypeBase<SparseDnTensorHandleType, Type, TypeStorage>::Base;
+ using Base::Base;
+
+ static constexpr StringLiteral name = "gpu.sparse.dntensor_handle";
+};
+
+class SparseSpMatHandleType
+ : public Type::TypeBase<SparseSpMatHandleType, Type, TypeStorage> {
+public:
+ using Base = Type::TypeBase<SparseSpMatHandleType, Type, TypeStorage>::Base;
+ using Base::Base;
+
+ static constexpr StringLiteral name = "gpu.sparse.spmat_handle";
+};
+
+class SparseSpGEMMOpHandleType
+ : public Type::TypeBase<SparseSpGEMMOpHandleType, Type, TypeStorage> {
+public:
+ using Base =
+ Type::TypeBase<SparseSpGEMMOpHandleType, Type, TypeStorage>::Base;
+ using Base::Base;
+
+ static constexpr StringLiteral name = "gpu.sparse.spgemmop_handle";
+};
+} // namespace mlir::gpu
+
+#define GET_ATTRDEF_CLASSES
+#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.h.inc"
+
+#endif // MLIR_DIALECT_GPU_IR_GPUBASE_H
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
index 34ff868db7d446..18b9d1dffb4274 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
@@ -17,7 +17,7 @@
#include "mlir/Bytecode/BytecodeOpInterface.h"
#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/DLTI/Traits.h"
-#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
+#include "mlir/Dialect/GPU/IR/GPUBase.h"
#include "mlir/Dialect/GPU/IR/GPUDialectDecl.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinTypes.h"
@@ -35,7 +35,6 @@
namespace mlir {
namespace gpu {
-
/// Utility class for the GPU dialect to represent triples of `Value`s
/// accessible through `.x`, `.y`, and `.z` similarly to CUDA notation.
struct KernelDim3 {
@@ -44,190 +43,13 @@ struct KernelDim3 {
Value z;
};
-class AsyncTokenType
- : public Type::TypeBase<AsyncTokenType, Type, TypeStorage> {
-public:
- // Used for generic hooks in TypeBase.
- using Base::Base;
-
- static constexpr StringLiteral name = "gpu.async_token";
-};
-
-class NamedBarrierType
- : public Type::TypeBase<NamedBarrierType, Type, TypeStorage> {
-public:
- using Base::Base;
-
- static constexpr StringLiteral name = "gpu.named_barrier";
-};
-
-/// MMAMatrixType storage and uniquing. Array is uniqued based on its shape
-/// and type.
-struct MMAMatrixStorageType : public TypeStorage {
- MMAMatrixStorageType(unsigned numDims, const int64_t *dimShapes,
- Type elementType, StringRef operand)
- : dimShapes(dimShapes), numDims(numDims), elementType(elementType),
- operand(operand) {}
-
- /// The hash key for uniquing.
- using KeyTy = std::tuple<ArrayRef<int64_t>, Type, StringRef>;
- bool operator==(const KeyTy &key) const {
- return key == KeyTy(getShape(), elementType, operand);
- }
-
- /// Construction.
- static MMAMatrixStorageType *construct(TypeStorageAllocator &allocator,
- const KeyTy &key) {
- ArrayRef<int64_t> shape = allocator.copyInto(std::get<0>(key));
- StringRef operand = allocator.copyInto(std::get<2>(key));
-
- return new (allocator.allocate<MMAMatrixStorageType>())
- MMAMatrixStorageType(shape.size(), shape.data(), std::get<1>(key),
- operand);
- }
-
- ArrayRef<int64_t> getShape() const {
- return ArrayRef<int64_t>(dimShapes, numDims);
- }
-
- StringRef getOperand() const { return operand; }
-
- /// Reference to the shape of the MMA matrix.
- const int64_t *dimShapes;
-
- /// Number of dimensions in the MMA matrix.
- unsigned numDims;
-
- /// Element type of elements held in the MMA matrix.
- Type elementType;
-
- /// MMA operand that this MMAMatrix holds. The general form of operation this
- /// type supports is given by the equation C += A*B. This field specifies
- /// which operand in the given equation is held by this type. The valid values
- /// are "AOp", "BOp" and "COp".
- StringRef operand;
-};
-
-/// MMAMatrix represents a matrix held by a subgroup for matrix-matrix multiply
-/// accumulate operations. MMAMatrices are taken as direct operands by these
-/// operations and are also produced as results. These matrices are meant to
-/// reside in the registers. A limited number of pointwise operations can be
-/// performed on these matrices, i.e., operations which operate uniformly on
-/// all the elements in the matrix and do not change the order of matrix
-/// elements. The above conditions exist because the layout of matrix elements
-/// inside the matrix is opaque i.e., the elements may be present in the
-/// matrix in any order. The general usage of this type is shown as follows:-
-///
-/// %0 = gpu.subgroup_mma_load_matrix %arg0[%c0, %c0] {leadDimension = 16 :
-/// index} : memref<16x16xf16> -> !gpu.mma_matrix<16x16xf16, "AOp">
-///
-/// The MMAMatrixType describes the shape of the matrix being loaded and the
-/// operand being loaded too. The operand needs to be specified to aid the
-/// lowering of this type to dialects such as NVVM where each workitem may
-/// hold different amount of elements depending on the elementType of the
-/// matrix. For e.g., Each workitem holds 4 vector<2xf16>s for f16 data type
-/// and 8 f32s for f32 data type of MMAMatrix. Some other instances of usage
-/// are:-
-///
-/// %3 = gpu.subgroup_mma_compute %0, %1, %2 :
-/// !gpu.mma_matrix<16x16xf16, "AOp">, !gpu.mma_matrix<16x16xf16, "BOp">
-/// -> !gpu.mma_matrix<16x16xf32, "COp">
-///
-///
-/// gpu.subgroup_mma_store_matrix %3, %arg22[%c0, %c0] {leadDimension = 16
-/// : index}: !gpu.mma_matrix<16x16xf32, "COp">, memref<16x16xf32>
-///
-/// MMA matrices may be memref elements on targets that support storing them,
-/// for example SPIRV can represent arrays of MMA matrices.
-// TODO: consider moving this to ODS.
-class MMAMatrixType
- : public Type::TypeBase<MMAMatrixType, Type, MMAMatrixStorageType,
- MemRefElementTypeInterface::Trait> {
-public:
- using Base::Base;
-
- static constexpr StringLiteral name = "gpu.mma_matrix";
-
- /// Get MMAMatrixType and verify construction Invariants.
- static MMAMatrixType get(ArrayRef<int64_t> shape, Type elementType,
- StringRef operand);
-
- /// Get MMAMatrixType at a particular location and verify construction
- /// Invariants.
- static MMAMatrixType getChecked(function_ref<InFlightDiagnostic()> emitError,
- ArrayRef<int64_t> shape, Type elementType,
- StringRef operand);
-
- /// Check if a type is valid a MMAMatrixType elementType.
- static bool isValidElementType(Type elementType);
-
- /// Verify that shape and elementType are actually allowed for the
- /// MMAMatrixType.
- static LogicalResult
- verifyInvariants(function_ref<InFlightDiagnostic()> emitError,
- ArrayRef<int64_t> shape, Type elementType,
- StringRef operand);
-
- /// Get number of dims.
- unsigned getNumDims() const;
-
- /// Get shape of the matrix.
- ArrayRef<int64_t> getShape() const;
-
- /// Get elementType of a single element.
- Type getElementType() const;
-
- /// The general form of operation this type supports is given by the equation
- /// C += A*B. This function returns which operand in the given equation is
- /// held by this type. String returned can be one of"AOp", "BOp" and "COp".
- StringRef getOperand() const;
-};
-
// Adds a `gpu.async.token` to the front of the argument list.
void addAsyncDependency(Operation *op, Value token);
-
-// Handle types for sparse.
-enum class SparseHandleKind { SpMat, DnTensor, SpGEMMOp };
-
-class SparseDnTensorHandleType
- : public Type::TypeBase<SparseDnTensorHandleType, Type, TypeStorage> {
-public:
- using Base =
- Type::TypeBase<SparseDnTensorHandleType, Type, TypeStorage>::Base;
- using Base::Base;
-
- static constexpr StringLiteral name = "gpu.sparse.dntensor_handle";
-};
-
-class SparseSpMatHandleType
- : public Type::TypeBase<SparseSpMatHandleType, Type, TypeStorage> {
-public:
- using Base = Type::TypeBase<SparseSpMatHandleType, Type, TypeStorage>::Base;
- using Base::Base;
-
- static constexpr StringLiteral name = "gpu.sparse.spmat_handle";
-};
-
-class SparseSpGEMMOpHandleType
- : public Type::TypeBase<SparseSpGEMMOpHandleType, Type, TypeStorage> {
-public:
- using Base =
- Type::TypeBase<SparseSpGEMMOpHandleType, Type, TypeStorage>::Base;
- using Base::Base;
-
- static constexpr StringLiteral name = "gpu.sparse.spgemmop_handle";
-};
-
} // namespace gpu
} // namespace mlir
#include "mlir/Dialect/GPU/IR/GPUOpInterfaces.h.inc"
-#include "mlir/Dialect/SCF/IR/DeviceMappingInterface.h"
-
-#define GET_ATTRDEF_CLASSES
-#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.h.inc"
-
#define GET_OP_CLASSES
#include "mlir/Dialect/GPU/IR/GPUOps.h.inc"
diff --git a/mlir/include/mlir/Dialect/OpenACC/Extensions/AllExtensions.h b/mlir/include/mlir/Dialect/OpenACC/Extensions/AllExtensions.h
new file mode 100644
index 00000000000000..a84fa1ab3277f4
--- /dev/null
+++ b/mlir/include/mlir/Dialect/OpenACC/Extensions/AllExtensions.h
@@ -0,0 +1,22 @@
+//===- AllExtensions.h - OpenACC external interfaces ---------*- 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_OPENACC_EXTENSIONS_ALLEXTENSIONS_H
+#define MLIR_DIALECT_OPENACC_EXTENSIONS_ALLEXTENSIONS_H
+
+namespace mlir {
+class DialectRegistry;
+namespace acc {
+/// Register the standard OpenACC models for GPU, LLVM and MemRef IR. Models
+/// for MemRef and LLVM also load the dialects used by their IR builders when
+/// OpenACC is loaded. The GPU model attaches when GPU and OpenACC are loaded.
+void registerAllExtensions(DialectRegistry ®istry);
+} // namespace acc
+} // namespace mlir
+
+#endif // MLIR_DIALECT_OPENACC_EXTENSIONS_ALLEXTENSIONS_H
diff --git a/mlir/include/mlir/Dialect/OpenACC/Extensions/GPUInterfaces.h b/mlir/include/mlir/Dialect/OpenACC/Extensions/GPUInterfaces.h
new file mode 100644
index 00000000000000..0661135673680e
--- /dev/null
+++ b/mlir/include/mlir/Dialect/OpenACC/Extensions/GPUInterfaces.h
@@ -0,0 +1,20 @@
+//===- GPUInterfaces.h - OpenACC external models -----------*- C++ -*-===//
+//
+// Part of the MLIR 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_OPENACC_EXTENSIONS_GPUINTERFACES_H
+#define MLIR_DIALECT_OPENACC_EXTENSIONS_GPUINTERFACES_H
+
+namespace mlir {
+class DialectRegistry;
+namespace acc {
+/// Register OpenACC interfaces for GPU IR. This does not load dialects.
+void registerGPUInterfaces(DialectRegistry ®istry);
+} // namespace acc
+} // namespace mlir
+
+#endif // MLIR_DIALECT_OPENACC_EXTENSIONS_GPUINTERFACES_H
diff --git a/mlir/include/mlir/Dialect/OpenACC/Extensions/LLVMInterfaces.h b/mlir/include/mlir/Dialect/OpenACC/Extensions/LLVMInterfaces.h
new file mode 100644
index 00000000000000..5d156dc77ef515
--- /dev/null
+++ b/mlir/include/mlir/Dialect/OpenACC/Extensions/LLVMInterfaces.h
@@ -0,0 +1,21 @@
+//===- LLVMInterfaces.h - OpenACC external models -----------*- C++ -*-===//
+//
+// Part of the MLIR 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_OPENACC_EXTENSIONS_LLVMINTERFACES_H
+#define MLIR_DIALECT_OPENACC_EXTENSIONS_LLVMINTERFACES_H
+
+namespace mlir {
+class DialectRegistry;
+namespace acc {
+/// Register OpenACC interfaces for LLVM IR. Loading OpenACC loads LLVM and
+/// Arith so that these interfaces can generate IR.
+void registerLLVMInterfaces(DialectRegistry ®istry);
+} // namespace acc
+} // namespace mlir
+
+#endif // MLIR_DIALECT_OPENACC_EXTENSIONS_LLVMINTERFACES_H
diff --git a/mlir/include/mlir/Dialect/OpenACC/Extensions/MemRefInterfaces.h b/mlir/include/mlir/Dialect/OpenACC/Extensions/MemRefInterfaces.h
new file mode 100644
index 00000000000000..44cb28bbe0a3f5
--- /dev/null
+++ b/mlir/include/mlir/Dialect/OpenACC/Extensions/MemRefInterfaces.h
@@ -0,0 +1,21 @@
+//===- MemRefInterfaces.h - OpenACC external models -----------*- C++ -*-===//
+//
+// Part of the MLIR 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_OPENACC_EXTENSIONS_MEMREFINTERFACES_H
+#define MLIR_DIALECT_OPENACC_EXTENSIONS_MEMREFINTERFACES_H
+
+namespace mlir {
+class DialectRegistry;
+namespace acc {
+/// Register OpenACC interfaces for MemRef IR. Loading OpenACC loads MemRef and
+/// Arith so that these interfaces can generate IR.
+void registerMemRefInterfaces(DialectRegistry ®istry);
+} // namespace acc
+} // namespace mlir
+
+#endif // MLIR_DIALECT_OPENACC_EXTENSIONS_MEMREFINTERFACES_H
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACC.h b/mlir/include/mlir/Dialect/OpenACC/OpenACC.h
index 3cfa8457759de1..bba456ceac70c1 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACC.h
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACC.h
@@ -22,7 +22,7 @@
#include "mlir/Bytecode/BytecodeOpInterface.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
-#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUBase.h"
#include "mlir/Dialect/OpenACC/OpenACCOpsDialect.h.inc"
#include "mlir/Dialect/OpenACC/OpenACCOpsEnums.h.inc"
#include "mlir/Dialect/OpenACC/OpenACCOpsInterfaces.h.inc"
@@ -79,7 +79,6 @@
namespace mlir {
namespace acc {
-
/// Enumeration used to encode the execution mapping on a loop construct.
/// They refer directly to the OpenACC 3.3 standard:
/// 2.9.2. gang
@@ -297,7 +296,6 @@ bool isGangWorkerVectorAllOne(ComputeOpT op) {
return isConstantIntValue(stripIndexCasts(numWorkers), 1) &&
isConstantIntValue(stripIndexCasts(vectorLength), 1);
}
-
} // namespace acc
} // namespace mlir
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCBase.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCBase.td
index 5810759c54298d..d18e5d28eb3091 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCBase.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCBase.td
@@ -21,9 +21,6 @@ def OpenACC_Dialect : Dialect {
let useDefaultAttributePrinterParser = 1;
let useDefaultTypePrinterParser = 1;
let cppNamespace = "::mlir::acc";
- let dependentDialects = ["::mlir::memref::MemRefDialect",
- "::mlir::LLVM::LLVMDialect",
- "::mlir::gpu::GPUDialect"];
}
#endif // OPENACC_BASE
diff --git a/mlir/lib/CAPI/Dialect/CMakeLists.txt b/mlir/lib/CAPI/Dialect/CMakeLists.txt
index 551f5a5a3df771..769b2f44ccfbff 100644
--- a/mlir/lib/CAPI/Dialect/CMakeLists.txt
+++ b/mlir/lib/CAPI/Dialect/CMakeLists.txt
@@ -291,6 +291,8 @@ add_mlir_upstream_c_api_library(MLIRCAPIOpenACC
LINK_LIBS PUBLIC
MLIRCAPIIR
+ MLIRGPUDialect
+ MLIROpenACCAllExtensions
MLIROpenACCDialect
MLIROpenACCTransforms
)
diff --git a/mlir/lib/CAPI/Dialect/OpenACC.cpp b/mlir/lib/CAPI/Dialect/OpenACC.cpp
index b304a8a5aacb13..03d814edb1e78a 100644
--- a/mlir/lib/CAPI/Dialect/OpenACC.cpp
+++ b/mlir/lib/CAPI/Dialect/OpenACC.cpp
@@ -8,8 +8,42 @@
#include "mlir-c/Dialect/OpenACC.h"
#include "mlir/CAPI/Registration.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
+using namespace mlir;
using namespace mlir::acc;
-MLIR_DEFINE_CAPI_DIALECT_REGISTRATION(OpenACC, acc, mlir::acc::OpenACCDialect)
+namespace {
+// Keep the C API's complete registration behavior. Frontends that only need
+// OpenACC IR can link MLIROpenACCDialect directly instead.
+void registerOpenACC(DialectRegistry ®istry) {
+ registry.insert<OpenACCDialect, gpu::GPUDialect>();
+ acc::registerAllExtensions(registry);
+ registry.addExtension(+[](MLIRContext *context, OpenACCDialect *) {
+ context->getOrLoadDialect<gpu::GPUDialect>();
+ });
+}
+
+void insertOpenACC(MlirDialectRegistry registry) {
+ registerOpenACC(*unwrap(registry));
+}
+
+MlirDialect loadOpenACC(MlirContext context) {
+ DialectRegistry registry;
+ registerOpenACC(registry);
+ unwrap(context)->appendDialectRegistry(registry);
+ return wrap(unwrap(context)->getOrLoadDialect<OpenACCDialect>());
+}
+
+MlirStringRef getNamespace() {
+ return wrap(OpenACCDialect::getDialectNamespace());
+}
+} // namespace
+
+MlirDialectHandle mlirGetDialectHandle__acc__() {
+ static MlirDialectRegistrationHooks hooks = {insertOpenACC, loadOpenACC,
+ getNamespace};
+ return MlirDialectHandle{&hooks};
+}
diff --git a/mlir/lib/Dialect/GPU/CMakeLists.txt b/mlir/lib/Dialect/GPU/CMakeLists.txt
index ae7ed6f4257b5b..67cb2f1e6dacd0 100644
--- a/mlir/lib/Dialect/GPU/CMakeLists.txt
+++ b/mlir/lib/Dialect/GPU/CMakeLists.txt
@@ -1,3 +1,26 @@
+add_mlir_library(MLIRGPUBase
+ IR/GPUBase.cpp
+
+ ADDITIONAL_HEADER_DIRS
+ ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/GPU
+
+ DEPENDS
+ MLIRGPUOpsIncGen
+ MLIRGPUOpsAttributesIncGen
+ MLIRGPUOpsEnumsGen
+ MLIRGPUOpInterfacesIncGen
+ MLIRGPUCompilationAttrInterfacesIncGen
+ MLIRParallelLoopMapperEnumsGen
+ MLIRGPUDeviceMapperEnumsGen
+ MLIRDeviceMappingInterfacesIncGen
+ intrinsics_gen
+
+ LINK_LIBS PUBLIC
+ MLIRFunctionInterfaces
+ MLIRIR
+ MLIRSupport
+ )
+
add_mlir_dialect_library(MLIRGPUDialect
IR/GPUDialect.cpp
IR/InferIntRangeInterfaceImpls.cpp
@@ -18,6 +41,7 @@ add_mlir_dialect_library(MLIRGPUDialect
intrinsics_gen
LINK_LIBS PUBLIC
+ MLIRGPUBase
MLIRArithDialect
MLIRDLTIDialect
MLIRControlFlowInterfaces
diff --git a/mlir/lib/Dialect/GPU/IR/GPUBase.cpp b/mlir/lib/Dialect/GPU/IR/GPUBase.cpp
new file mode 100644
index 00000000000000..c06792ca53b0a1
--- /dev/null
+++ b/mlir/lib/Dialect/GPU/IR/GPUBase.cpp
@@ -0,0 +1,509 @@
+//===- GPUBase.cpp - GPU types and attributes -----------------------------===//
+//
+// 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 GPU types, attributes, and enums.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/GPU/IR/GPUBase.h"
+
+#include "mlir/IR/Builders.h"
+#include "mlir/IR/DialectImplementation.h"
+#include "llvm/ADT/TypeSwitch.h"
+#include "llvm/Support/CommandLine.h"
+#include "llvm/Support/StringSaver.h"
+
+using namespace mlir;
+using namespace mlir::gpu;
+
+//===----------------------------------------------------------------------===//
+// GPU Device Mapping Attributes
+//===----------------------------------------------------------------------===//
+
+int64_t GPUBlockMappingAttr::getMappingId() const {
+ return static_cast<int64_t>(getBlock());
+}
+
+bool GPUBlockMappingAttr::isLinearMapping() const {
+ return getMappingId() >= static_cast<int64_t>(MappingId::LinearDim0);
+}
+
+int64_t GPUBlockMappingAttr::getRelativeIndex() const {
+ return isLinearMapping()
+ ? getMappingId() - static_cast<int64_t>(MappingId::LinearDim0)
+ : getMappingId();
+}
+
+int64_t GPUWarpgroupMappingAttr::getMappingId() const {
+ return static_cast<int64_t>(getWarpgroup());
+}
+
+bool GPUWarpgroupMappingAttr::isLinearMapping() const {
+ return getMappingId() >= static_cast<int64_t>(MappingId::LinearDim0);
+}
+
+int64_t GPUWarpgroupMappingAttr::getRelativeIndex() const {
+ return isLinearMapping()
+ ? getMappingId() - static_cast<int64_t>(MappingId::LinearDim0)
+ : getMappingId();
+}
+
+int64_t GPUWarpMappingAttr::getMappingId() const {
+ return static_cast<int64_t>(getWarp());
+}
+
+bool GPUWarpMappingAttr::isLinearMapping() const {
+ return getMappingId() >= static_cast<int64_t>(MappingId::LinearDim0);
+}
+
+int64_t GPUWarpMappingAttr::getRelativeIndex() const {
+ return isLinearMapping()
+ ? getMappingId() - static_cast<int64_t>(MappingId::LinearDim0)
+ : getMappingId();
+}
+
+int64_t GPUThreadMappingAttr::getMappingId() const {
+ return static_cast<int64_t>(getThread());
+}
+
+bool GPUThreadMappingAttr::isLinearMapping() const {
+ return getMappingId() >= static_cast<int64_t>(MappingId::LinearDim0);
+}
+
+int64_t GPUThreadMappingAttr::getRelativeIndex() const {
+ return isLinearMapping()
+ ? getMappingId() - static_cast<int64_t>(MappingId::LinearDim0)
+ : getMappingId();
+}
+
+int64_t GPULaneMappingAttr::getMappingId() const {
+ return static_cast<int64_t>(getLane());
+}
+
+bool GPULaneMappingAttr::isLinearMapping() const {
+ return getMappingId() >= static_cast<int64_t>(MappingId::LinearDim0);
+}
+
+int64_t GPULaneMappingAttr::getRelativeIndex() const {
+ return isLinearMapping()
+ ? getMappingId() - static_cast<int64_t>(MappingId::LinearDim0)
+ : getMappingId();
+}
+
+int64_t GPUMemorySpaceMappingAttr::getMappingId() const {
+ return static_cast<int64_t>(getAddressSpace());
+}
+
+bool GPUMemorySpaceMappingAttr::isLinearMapping() const {
+ llvm_unreachable("GPUMemorySpaceMappingAttr does not support linear mapping");
+}
+
+int64_t GPUMemorySpaceMappingAttr::getRelativeIndex() const {
+ llvm_unreachable("GPUMemorySpaceMappingAttr does not support relative index");
+}
+
+//===----------------------------------------------------------------------===//
+// MMAMatrixType
+//===----------------------------------------------------------------------===//
+
+MMAMatrixType MMAMatrixType::get(ArrayRef<int64_t> shape, Type elementType,
+ StringRef operand) {
+ return Base::get(elementType.getContext(), shape, elementType, operand);
+}
+
+MMAMatrixType
+MMAMatrixType::getChecked(function_ref<InFlightDiagnostic()> emitError,
+ ArrayRef<int64_t> shape, Type elementType,
+ StringRef operand) {
+ return Base::getChecked(emitError, elementType.getContext(), shape,
+ elementType, operand);
+}
+
+unsigned MMAMatrixType::getNumDims() const { return getImpl()->numDims; }
+
+ArrayRef<int64_t> MMAMatrixType::getShape() const {
+ return getImpl()->getShape();
+}
+
+Type MMAMatrixType::getElementType() const { return getImpl()->elementType; }
+
+StringRef MMAMatrixType::getOperand() const { return getImpl()->getOperand(); }
+
+bool MMAMatrixType::isValidElementType(Type elementType) {
+ return elementType.isF16() || elementType.isF32() || elementType.isF64() ||
+ elementType.isUnsignedInteger(8) || elementType.isSignedInteger(8) ||
+ elementType.isInteger(32);
+}
+
+LogicalResult
+MMAMatrixType::verifyInvariants(function_ref<InFlightDiagnostic()> emitError,
+ ArrayRef<int64_t> shape, Type elementType,
+ StringRef operand) {
+ if (operand != "AOp" && operand != "BOp" && operand != "COp")
+ return emitError() << "operand expected to be one of AOp, BOp or COp";
+
+ if (shape.size() != 2)
+ return emitError() << "MMAMatrixType must have exactly two dimensions";
+
+ if (!MMAMatrixType::isValidElementType(elementType))
+ return emitError()
+ << "MMAMatrixType elements must be SI8, UI8, I32, F16, F32, or F64";
+
+ return success();
+}
+
+//===----------------------------------------------------------------------===//
+// GPUDialect
+//===----------------------------------------------------------------------===//
+
+bool GPUDialect::isWorkgroupMemoryAddressSpace(Attribute memorySpace) {
+ if (!memorySpace)
+ return false;
+ if (auto gpuAttr = llvm::dyn_cast<gpu::AddressSpaceAttr>(memorySpace))
+ return gpuAttr.getValue() == getWorkgroupAddressSpace();
+ return false;
+}
+
+bool GPUDialect::hasWorkgroupMemoryAddressSpace(MemRefType type) {
+ Attribute memorySpace = type.getMemorySpace();
+ return isWorkgroupMemoryAddressSpace(memorySpace);
+}
+
+bool GPUDialect::isConstantMemoryAddressSpace(Attribute memorySpace) {
+ if (!memorySpace)
+ return false;
+ if (auto gpuAttr = llvm::dyn_cast<gpu::AddressSpaceAttr>(memorySpace))
+ return gpuAttr.getValue() == getConstantAddressSpace();
+ return false;
+}
+
+bool GPUDialect::hasConstantMemoryAddressSpace(MemRefType type) {
+ Attribute memorySpace = type.getMemorySpace();
+ return isConstantMemoryAddressSpace(memorySpace);
+}
+
+//===----------------------------------------------------------------------===//
+// GPU object attribute
+//===----------------------------------------------------------------------===//
+
+LogicalResult ObjectAttr::verify(function_ref<InFlightDiagnostic()> emitError,
+ Attribute target, CompilationTarget format,
+ StringAttr object, DictionaryAttr properties,
+ KernelTableAttr kernels) {
+ if (!target)
+ return emitError() << "the target attribute cannot be null";
+ if (target.hasPromiseOrImplementsInterface<TargetAttrInterface>())
+ return success();
+ return emitError() << "the target attribute must implement or promise the "
+ "`gpu::TargetAttrInterface`";
+}
+
+namespace {
+ParseResult parseObject(AsmParser &odsParser, CompilationTarget &format,
+ StringAttr &object) {
+ std::optional<CompilationTarget> formatResult;
+ StringRef enumKeyword;
+ auto loc = odsParser.getCurrentLocation();
+ if (failed(odsParser.parseOptionalKeyword(&enumKeyword)))
+ formatResult = CompilationTarget::Fatbin;
+ if (!formatResult &&
+ (formatResult =
+ gpu::symbolizeEnum<gpu::CompilationTarget>(enumKeyword)) &&
+ odsParser.parseEqual())
+ return odsParser.emitError(loc, "expected an equal sign");
+ if (!formatResult)
+ return odsParser.emitError(loc, "expected keyword for GPU object format");
+ FailureOr<StringAttr> objectResult =
+ FieldParser<StringAttr>::parse(odsParser);
+ if (failed(objectResult))
+ return odsParser.emitError(odsParser.getCurrentLocation(),
+ "failed to parse GPU_ObjectAttr parameter "
+ "'object' which is to be a `StringAttr`");
+ format = *formatResult;
+ object = *objectResult;
+ return success();
+}
+
+void printObject(AsmPrinter &odsParser, CompilationTarget format,
+ StringAttr object) {
+ if (format != CompilationTarget::Fatbin)
+ odsParser << stringifyEnum(format) << " = ";
+ odsParser << object;
+}
+} // namespace
+
+//===----------------------------------------------------------------------===//
+// GPU select object attribute
+//===----------------------------------------------------------------------===//
+
+LogicalResult
+gpu::SelectObjectAttr::verify(function_ref<InFlightDiagnostic()> emitError,
+ Attribute target) {
+ // Check `target`, it can be null, an integer attr or a GPU Target attribute.
+ if (target) {
+ if (auto intAttr = mlir::dyn_cast<IntegerAttr>(target)) {
+ if (intAttr.getInt() < 0) {
+ return emitError() << "the object index must be positive";
+ }
+ } else if (!target.hasPromiseOrImplementsInterface<TargetAttrInterface>()) {
+ return emitError()
+ << "the target attribute must be a GPU Target attribute";
+ }
+ }
+ return success();
+}
+
+//===----------------------------------------------------------------------===//
+// GPU KernelMetadataAttr
+//===----------------------------------------------------------------------===//
+
+KernelMetadataAttr KernelMetadataAttr::get(FunctionOpInterface kernel,
+ DictionaryAttr metadata) {
+ assert(kernel && "invalid kernel");
+ return get(kernel.getNameAttr(), kernel.getFunctionType(),
+ kernel.getAllArgAttrs(), metadata);
+}
+
+KernelMetadataAttr
+KernelMetadataAttr::getChecked(function_ref<InFlightDiagnostic()> emitError,
+ FunctionOpInterface kernel,
+ DictionaryAttr metadata) {
+ assert(kernel && "invalid kernel");
+ return getChecked(emitError, kernel.getNameAttr(), kernel.getFunctionType(),
+ kernel.getAllArgAttrs(), metadata);
+}
+
+KernelMetadataAttr
+KernelMetadataAttr::appendMetadata(ArrayRef<NamedAttribute> attrs) const {
+ if (attrs.empty())
+ return *this;
+ NamedAttrList attrList;
+ if (DictionaryAttr dict = getMetadata())
+ attrList.append(dict);
+ attrList.append(attrs);
+ return KernelMetadataAttr::get(getName(), getFunctionType(), getArgAttrs(),
+ attrList.getDictionary(getContext()));
+}
+
+LogicalResult
+KernelMetadataAttr::verify(function_ref<InFlightDiagnostic()> emitError,
+ StringAttr name, Type functionType,
+ ArrayAttr argAttrs, DictionaryAttr metadata) {
+ if (name.empty())
+ return emitError() << "the kernel name can't be empty";
+ if (argAttrs) {
+ if (llvm::any_of(argAttrs, [](Attribute attr) {
+ return !llvm::isa<DictionaryAttr>(attr);
+ }))
+ return emitError()
+ << "all attributes in the array must be a dictionary attribute";
+ }
+ return success();
+}
+
+//===----------------------------------------------------------------------===//
+// GPU KernelTableAttr
+//===----------------------------------------------------------------------===//
+
+KernelTableAttr KernelTableAttr::get(MLIRContext *context,
+ ArrayRef<KernelMetadataAttr> kernels,
+ bool isSorted) {
+ // Note that `is_sorted` is always only invoked once even with assertions ON.
+ assert((!isSorted || llvm::is_sorted(kernels)) &&
+ "expected a sorted kernel array");
+ // Immediately return the attribute if the array is sorted.
+ if (isSorted || llvm::is_sorted(kernels))
+ return Base::get(context, kernels);
+ // Sort the array.
+ SmallVector<KernelMetadataAttr> kernelsTmp(kernels);
+ llvm::array_pod_sort(kernelsTmp.begin(), kernelsTmp.end());
+ return Base::get(context, kernelsTmp);
+}
+
+KernelTableAttr KernelTableAttr::getChecked(
+ function_ref<InFlightDiagnostic()> emitError, MLIRContext *context,
+ ArrayRef<KernelMetadataAttr> kernels, bool isSorted) {
+ // Note that `is_sorted` is always only invoked once even with assertions ON.
+ assert((!isSorted || llvm::is_sorted(kernels)) &&
+ "expected a sorted kernel array");
+ // Immediately return the attribute if the array is sorted.
+ if (isSorted || llvm::is_sorted(kernels))
+ return Base::getChecked(emitError, context, kernels);
+ // Sort the array.
+ SmallVector<KernelMetadataAttr> kernelsTmp(kernels);
+ llvm::array_pod_sort(kernelsTmp.begin(), kernelsTmp.end());
+ return Base::getChecked(emitError, context, kernelsTmp);
+}
+
+LogicalResult
+KernelTableAttr::verify(function_ref<InFlightDiagnostic()> emitError,
+ ArrayRef<KernelMetadataAttr> kernels) {
+ if (kernels.size() < 2)
+ return success();
+ // Check that the kernels are uniquely named.
+ if (std::adjacent_find(kernels.begin(), kernels.end(),
+ [](KernelMetadataAttr l, KernelMetadataAttr r) {
+ return l.getName() == r.getName();
+ }) != kernels.end()) {
+ return emitError() << "expected all kernels to be uniquely named";
+ }
+ return success();
+}
+
+KernelMetadataAttr KernelTableAttr::lookup(StringRef key) const {
+ auto [iterator, found] = impl::findAttrSorted(begin(), end(), key);
+ return found ? *iterator : KernelMetadataAttr();
+}
+
+KernelMetadataAttr KernelTableAttr::lookup(StringAttr key) const {
+ auto [iterator, found] = impl::findAttrSorted(begin(), end(), key);
+ return found ? *iterator : KernelMetadataAttr();
+}
+
+//===----------------------------------------------------------------------===//
+// GPU target options
+//===----------------------------------------------------------------------===//
+
+TargetOptions::TargetOptions(
+ StringRef toolkitPath, ArrayRef<Attribute> librariesToLink,
+ StringRef cmdOptions, StringRef elfSection,
+ CompilationTarget compilationTarget,
+ function_ref<SymbolTable *()> getSymbolTableCallback,
+ function_ref<void(llvm::Module &)> initialLlvmIRCallback,
+ function_ref<void(llvm::Module &)> linkedLlvmIRCallback,
+ function_ref<void(llvm::Module &)> optimizedLlvmIRCallback,
+ function_ref<void(StringRef)> isaCallback)
+ : TargetOptions(TypeID::get<TargetOptions>(), toolkitPath, librariesToLink,
+ cmdOptions, elfSection, compilationTarget,
+ getSymbolTableCallback, initialLlvmIRCallback,
+ linkedLlvmIRCallback, optimizedLlvmIRCallback,
+ isaCallback) {}
+
+TargetOptions::TargetOptions(
+ TypeID typeID, StringRef toolkitPath, ArrayRef<Attribute> librariesToLink,
+ StringRef cmdOptions, StringRef elfSection,
+ CompilationTarget compilationTarget,
+ function_ref<SymbolTable *()> getSymbolTableCallback,
+ function_ref<void(llvm::Module &)> initialLlvmIRCallback,
+ function_ref<void(llvm::Module &)> linkedLlvmIRCallback,
+ function_ref<void(llvm::Module &)> optimizedLlvmIRCallback,
+ function_ref<void(StringRef)> isaCallback)
+ : toolkitPath(toolkitPath.str()), librariesToLink(librariesToLink),
+ cmdOptions(cmdOptions.str()), elfSection(elfSection.str()),
+ compilationTarget(compilationTarget),
+ getSymbolTableCallback(getSymbolTableCallback),
+ initialLlvmIRCallback(initialLlvmIRCallback),
+ linkedLlvmIRCallback(linkedLlvmIRCallback),
+ optimizedLlvmIRCallback(optimizedLlvmIRCallback),
+ isaCallback(isaCallback), typeID(typeID) {}
+
+TypeID TargetOptions::getTypeID() const { return typeID; }
+
+StringRef TargetOptions::getToolkitPath() const { return toolkitPath; }
+
+ArrayRef<Attribute> TargetOptions::getLibrariesToLink() const {
+ return librariesToLink;
+}
+
+StringRef TargetOptions::getCmdOptions() const { return cmdOptions; }
+
+StringRef TargetOptions::getELFSection() const { return elfSection; }
+
+SymbolTable *TargetOptions::getSymbolTable() const {
+ return getSymbolTableCallback ? getSymbolTableCallback() : nullptr;
+}
+
+function_ref<void(llvm::Module &)>
+TargetOptions::getInitialLlvmIRCallback() const {
+ return initialLlvmIRCallback;
+}
+
+function_ref<void(llvm::Module &)>
+TargetOptions::getLinkedLlvmIRCallback() const {
+ return linkedLlvmIRCallback;
+}
+
+function_ref<void(llvm::Module &)>
+TargetOptions::getOptimizedLlvmIRCallback() const {
+ return optimizedLlvmIRCallback;
+}
+
+function_ref<void(StringRef)> TargetOptions::getISACallback() const {
+ return isaCallback;
+}
+
+CompilationTarget TargetOptions::getCompilationTarget() const {
+ return compilationTarget;
+}
+
+CompilationTarget TargetOptions::getDefaultCompilationTarget() {
+ return CompilationTarget::Fatbin;
+}
+
+std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
+TargetOptions::tokenizeCmdOptions(const std::string &cmdOptions) {
+ std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>> options;
+ llvm::StringSaver stringSaver(options.first);
+ StringRef opts = cmdOptions;
+ // For a correct tokenization of the command line options `opts` must be
+ // unquoted, otherwise the tokenization function returns a single string: the
+ // unquoted `cmdOptions` -which is not the desired behavior.
+ // Remove any quotes if they are at the beginning and end of the string:
+ if (!opts.empty() && opts.front() == '"' && opts.back() == '"')
+ opts.consume_front("\""), opts.consume_back("\"");
+ if (!opts.empty() && opts.front() == '\'' && opts.back() == '\'')
+ opts.consume_front("'"), opts.consume_back("'");
+#ifdef _WIN32
+ llvm::cl::TokenizeWindowsCommandLine(opts, stringSaver, options.second,
+ /*MarkEOLs=*/false);
+#else
+ llvm::cl::TokenizeGNUCommandLine(opts, stringSaver, options.second,
+ /*MarkEOLs=*/false);
+#endif // _WIN32
+ return options;
+}
+
+std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
+TargetOptions::tokenizeCmdOptions() const {
+ return tokenizeCmdOptions(cmdOptions);
+}
+
+std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
+TargetOptions::tokenizeAndRemoveSuffixCmdOptions(llvm::StringRef startsWith) {
+ size_t startPos = cmdOptions.find(startsWith);
+ if (startPos == std::string::npos)
+ return {llvm::BumpPtrAllocator(), SmallVector<const char *>()};
+
+ auto tokenized =
+ tokenizeCmdOptions(cmdOptions.substr(startPos + startsWith.size()));
+ cmdOptions.resize(startPos);
+ return tokenized;
+}
+
+MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions)
+
+#include "mlir/Dialect/GPU/IR/GPUOpsEnums.cpp.inc"
+
+#define GET_ATTRDEF_CLASSES
+#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
+
+namespace {
+template <typename... Attrs>
+void registerAttributeStorageImpl(MLIRContext *context) {
+ (mlir::detail::AttributeUniquer::registerAttribute<Attrs>(context), ...);
+}
+} // namespace
+
+void mlir::gpu::detail::registerAttributeStorage(MLIRContext *context) {
+ registerAttributeStorageImpl<
+#define GET_ATTRDEF_LIST
+#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
+ >(context);
+}
+
+#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.cpp.inc"
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index ec31892edf1919..a17692832e9c14 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -36,11 +36,9 @@
#include "mlir/Transforms/InliningUtils.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/TypeSwitch.h"
-#include "llvm/Support/CommandLine.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/FormatVariadic.h"
#include "llvm/Support/InterleavedRange.h"
-#include "llvm/Support/StringSaver.h"
#include <cassert>
#include <numeric>
#include <optional>
@@ -51,6 +49,17 @@ using namespace mlir::gpu;
#include "mlir/Dialect/GPU/IR/GPUOpsDialect.cpp.inc"
namespace {
+// Attribute storage definitions live in MLIRGPUBase. Keep interface metadata
+// here: some attribute interfaces build GPU/Arith/Math IR, and their models
+// belong with the operation implementation rather than the representations.
+template <typename... Attrs>
+void registerAttributeMetadata(
+ Dialect &dialect,
+ function_ref<void(TypeID, AbstractAttribute &&)> addAttribute) {
+ (addAttribute(Attrs::getTypeID(), AbstractAttribute::get<Attrs>(dialect)),
+ ...);
+}
+
struct GPUDialectBytecodeInterface : BytecodeDialectInterface {
using BytecodeDialectInterface::BytecodeDialectInterface;
@@ -72,80 +81,6 @@ struct GPUDialectBytecodeInterface : BytecodeDialectInterface {
};
} // namespace
-//===----------------------------------------------------------------------===//
-// GPU Device Mapping Attributes
-//===----------------------------------------------------------------------===//
-
-int64_t GPUBlockMappingAttr::getMappingId() const {
- return static_cast<int64_t>(getBlock());
-}
-
-bool GPUBlockMappingAttr::isLinearMapping() const {
- return getMappingId() >= static_cast<int64_t>(MappingId::LinearDim0);
-}
-
-int64_t GPUBlockMappingAttr::getRelativeIndex() const {
- return isLinearMapping()
- ? getMappingId() - static_cast<int64_t>(MappingId::LinearDim0)
- : getMappingId();
-}
-
-int64_t GPUWarpgroupMappingAttr::getMappingId() const {
- return static_cast<int64_t>(getWarpgroup());
-}
-
-bool GPUWarpgroupMappingAttr::isLinearMapping() const {
- return getMappingId() >= static_cast<int64_t>(MappingId::LinearDim0);
-}
-
-int64_t GPUWarpgroupMappingAttr::getRelativeIndex() const {
- return isLinearMapping()
- ? getMappingId() - static_cast<int64_t>(MappingId::LinearDim0)
- : getMappingId();
-}
-
-int64_t GPUWarpMappingAttr::getMappingId() const {
- return static_cast<int64_t>(getWarp());
-}
-
-bool GPUWarpMappingAttr::isLinearMapping() const {
- return getMappingId() >= static_cast<int64_t>(MappingId::LinearDim0);
-}
-
-int64_t GPUWarpMappingAttr::getRelativeIndex() const {
- return isLinearMapping()
- ? getMappingId() - static_cast<int64_t>(MappingId::LinearDim0)
- : getMappingId();
-}
-
-int64_t GPUThreadMappingAttr::getMappingId() const {
- return static_cast<int64_t>(getThread());
-}
-
-bool GPUThreadMappingAttr::isLinearMapping() const {
- return getMappingId() >= static_cast<int64_t>(MappingId::LinearDim0);
-}
-
-int64_t GPUThreadMappingAttr::getRelativeIndex() const {
- return isLinearMapping()
- ? getMappingId() - static_cast<int64_t>(MappingId::LinearDim0)
- : getMappingId();
-}
-
-int64_t GPULaneMappingAttr::getMappingId() const {
- return static_cast<int64_t>(getLane());
-}
-
-bool GPULaneMappingAttr::isLinearMapping() const {
- return getMappingId() >= static_cast<int64_t>(MappingId::LinearDim0);
-}
-
-int64_t GPULaneMappingAttr::getRelativeIndex() const {
- return isLinearMapping()
- ? getMappingId() - static_cast<int64_t>(MappingId::LinearDim0)
- : getMappingId();
-}
-
int64_t GPUMappingMaskAttr::getMaxNumPhysicalIds() const { return 64; }
/// 8 4 0
@@ -193,98 +128,6 @@ Value GPUMappingMaskAttr::createIsActiveIdPredicate(
zero);
}
-int64_t GPUMemorySpaceMappingAttr::getMappingId() const {
- return static_cast<int64_t>(getAddressSpace());
-}
-
-bool GPUMemorySpaceMappingAttr::isLinearMapping() const {
- llvm_unreachable("GPUMemorySpaceMappingAttr does not support linear mapping");
-}
-
-int64_t GPUMemorySpaceMappingAttr::getRelativeIndex() const {
- llvm_unreachable("GPUMemorySpaceMappingAttr does not support relative index");
-}
-
-//===----------------------------------------------------------------------===//
-// MMAMatrixType
-//===----------------------------------------------------------------------===//
-
-MMAMatrixType MMAMatrixType::get(ArrayRef<int64_t> shape, Type elementType,
- StringRef operand) {
- return Base::get(elementType.getContext(), shape, elementType, operand);
-}
-
-MMAMatrixType
-MMAMatrixType::getChecked(function_ref<InFlightDiagnostic()> emitError,
- ArrayRef<int64_t> shape, Type elementType,
- StringRef operand) {
- return Base::getChecked(emitError, elementType.getContext(), shape,
- elementType, operand);
-}
-
-unsigned MMAMatrixType::getNumDims() const { return getImpl()->numDims; }
-
-ArrayRef<int64_t> MMAMatrixType::getShape() const {
- return getImpl()->getShape();
-}
-
-Type MMAMatrixType::getElementType() const { return getImpl()->elementType; }
-
-StringRef MMAMatrixType::getOperand() const { return getImpl()->getOperand(); }
-
-bool MMAMatrixType::isValidElementType(Type elementType) {
- return elementType.isF16() || elementType.isF32() || elementType.isF64() ||
- elementType.isUnsignedInteger(8) || elementType.isSignedInteger(8) ||
- elementType.isInteger(32);
-}
-
-LogicalResult
-MMAMatrixType::verifyInvariants(function_ref<InFlightDiagnostic()> emitError,
- ArrayRef<int64_t> shape, Type elementType,
- StringRef operand) {
- if (operand != "AOp" && operand != "BOp" && operand != "COp")
- return emitError() << "operand expected to be one of AOp, BOp or COp";
-
- if (shape.size() != 2)
- return emitError() << "MMAMatrixType must have exactly two dimensions";
-
- if (!MMAMatrixType::isValidElementType(elementType))
- return emitError()
- << "MMAMatrixType elements must be SI8, UI8, I32, F16, F32, or F64";
-
- return success();
-}
-
-//===----------------------------------------------------------------------===//
-// GPUDialect
-//===----------------------------------------------------------------------===//
-
-bool GPUDialect::isWorkgroupMemoryAddressSpace(Attribute memorySpace) {
- if (!memorySpace)
- return false;
- if (auto gpuAttr = llvm::dyn_cast<gpu::AddressSpaceAttr>(memorySpace))
- return gpuAttr.getValue() == getWorkgroupAddressSpace();
- return false;
-}
-
-bool GPUDialect::hasWorkgroupMemoryAddressSpace(MemRefType type) {
- Attribute memorySpace = type.getMemorySpace();
- return isWorkgroupMemoryAddressSpace(memorySpace);
-}
-
-bool GPUDialect::isConstantMemoryAddressSpace(Attribute memorySpace) {
- if (!memorySpace)
- return false;
- if (auto gpuAttr = llvm::dyn_cast<gpu::AddressSpaceAttr>(memorySpace))
- return gpuAttr.getValue() == getConstantAddressSpace();
- return false;
-}
-
-bool GPUDialect::hasConstantMemoryAddressSpace(MemRefType type) {
- Attribute memorySpace = type.getMemorySpace();
- return isConstantMemoryAddressSpace(memorySpace);
-}
-
bool GPUDialect::isKernel(Operation *op) {
if (auto gpuFunc = dyn_cast<GPUFuncOp>(op))
return gpuFunc.isKernel();
@@ -317,10 +160,15 @@ void GPUDialect::initialize() {
#define GET_OP_LIST
#include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc"
>();
- addAttributes<
+ // Match addAttributes while keeping the storage and interface instantiations
+ // in their respective implementation libraries.
+ registerAttributeMetadata<
#define GET_ATTRDEF_LIST
#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
- >();
+ >(*this, [this](TypeID id, AbstractAttribute &&attribute) {
+ addAttribute(id, std::move(attribute));
+ });
+ gpu::detail::registerAttributeStorage(getContext());
addInterfaces<GPUInlinerInterface>();
declarePromisedInterface<bufferization::BufferDeallocationOpInterface,
TerminatorOp>();
@@ -2137,7 +1985,6 @@ LogicalResult MemcpyOp::verify() {
}
namespace {
-
/// Erases a common case of copy ops where a destination value is used only by
/// the copy op, alloc and dealloc ops.
struct EraseTrivialCopyOp : public OpRewritePattern<MemcpyOp> {
@@ -2170,7 +2017,6 @@ struct EraseTrivialCopyOp : public OpRewritePattern<MemcpyOp> {
return success();
}
};
-
} // end anonymous namespace
void MemcpyOp::getCanonicalizationPatterns(RewritePatternSet &results,
@@ -2262,7 +2108,6 @@ LogicalResult MemsetOp::fold(FoldAdaptor adaptor,
//===----------------------------------------------------------------------===//
namespace {
-
/// Remove gpu.wait op use of gpu.wait op def without async dependencies.
/// %t = gpu.wait async [] // No async dependencies.
/// ... gpu.wait ... [%t, ...] // %t can be removed.
@@ -2322,7 +2167,6 @@ struct SimplifyGpuWaitOp : public OpRewritePattern<WaitOp> {
return failure();
}
};
-
} // end anonymous namespace
void WaitOp::getCanonicalizationPatterns(RewritePatternSet &results,
@@ -2353,7 +2197,6 @@ LogicalResult AllocOp::verify() {
}
namespace {
-
/// Folding of memref.dim(gpu.alloc(%size), %idx) -> %size similar to
/// `memref::AllocOp`.
struct SimplifyDimOfAllocOp : public OpRewritePattern<memref::DimOp> {
@@ -2381,7 +2224,6 @@ struct SimplifyDimOfAllocOp : public OpRewritePattern<memref::DimOp> {
return success();
}
};
-
} // namespace
void AllocOp::getCanonicalizationPatterns(RewritePatternSet &results,
@@ -2389,77 +2231,6 @@ void AllocOp::getCanonicalizationPatterns(RewritePatternSet &results,
results.add<SimplifyDimOfAllocOp>(context);
}
-//===----------------------------------------------------------------------===//
-// GPU object attribute
-//===----------------------------------------------------------------------===//
-
-LogicalResult ObjectAttr::verify(function_ref<InFlightDiagnostic()> emitError,
- Attribute target, CompilationTarget format,
- StringAttr object, DictionaryAttr properties,
- KernelTableAttr kernels) {
- if (!target)
- return emitError() << "the target attribute cannot be null";
- if (target.hasPromiseOrImplementsInterface<TargetAttrInterface>())
- return success();
- return emitError() << "the target attribute must implement or promise the "
- "`gpu::TargetAttrInterface`";
-}
-
-namespace {
-ParseResult parseObject(AsmParser &odsParser, CompilationTarget &format,
- StringAttr &object) {
- std::optional<CompilationTarget> formatResult;
- StringRef enumKeyword;
- auto loc = odsParser.getCurrentLocation();
- if (failed(odsParser.parseOptionalKeyword(&enumKeyword)))
- formatResult = CompilationTarget::Fatbin;
- if (!formatResult &&
- (formatResult =
- gpu::symbolizeEnum<gpu::CompilationTarget>(enumKeyword)) &&
- odsParser.parseEqual())
- return odsParser.emitError(loc, "expected an equal sign");
- if (!formatResult)
- return odsParser.emitError(loc, "expected keyword for GPU object format");
- FailureOr<StringAttr> objectResult =
- FieldParser<StringAttr>::parse(odsParser);
- if (failed(objectResult))
- return odsParser.emitError(odsParser.getCurrentLocation(),
- "failed to parse GPU_ObjectAttr parameter "
- "'object' which is to be a `StringAttr`");
- format = *formatResult;
- object = *objectResult;
- return success();
-}
-
-void printObject(AsmPrinter &odsParser, CompilationTarget format,
- StringAttr object) {
- if (format != CompilationTarget::Fatbin)
- odsParser << stringifyEnum(format) << " = ";
- odsParser << object;
-}
-} // namespace
-
-//===----------------------------------------------------------------------===//
-// GPU select object attribute
-//===----------------------------------------------------------------------===//
-
-LogicalResult
-gpu::SelectObjectAttr::verify(function_ref<InFlightDiagnostic()> emitError,
- Attribute target) {
- // Check `target`, it can be null, an integer attr or a GPU Target attribute.
- if (target) {
- if (auto intAttr = mlir::dyn_cast<IntegerAttr>(target)) {
- if (intAttr.getInt() < 0) {
- return emitError() << "the object index must be positive";
- }
- } else if (!target.hasPromiseOrImplementsInterface<TargetAttrInterface>()) {
- return emitError()
- << "the target attribute must be a GPU Target attribute";
- }
- }
- return success();
-}
-
//===----------------------------------------------------------------------===//
// DynamicSharedMemoryOp
//===----------------------------------------------------------------------===//
@@ -2722,242 +2493,7 @@ OpFoldResult gpu::SubgroupBroadcastOp::fold(FoldAdaptor /*adaptor*/) {
// No custom implementations needed; ballot uses default behavior from ODS.
-//===----------------------------------------------------------------------===//
-// GPU KernelMetadataAttr
-//===----------------------------------------------------------------------===//
-
-KernelMetadataAttr KernelMetadataAttr::get(FunctionOpInterface kernel,
- DictionaryAttr metadata) {
- assert(kernel && "invalid kernel");
- return get(kernel.getNameAttr(), kernel.getFunctionType(),
- kernel.getAllArgAttrs(), metadata);
-}
-
-KernelMetadataAttr
-KernelMetadataAttr::getChecked(function_ref<InFlightDiagnostic()> emitError,
- FunctionOpInterface kernel,
- DictionaryAttr metadata) {
- assert(kernel && "invalid kernel");
- return getChecked(emitError, kernel.getNameAttr(), kernel.getFunctionType(),
- kernel.getAllArgAttrs(), metadata);
-}
-
-KernelMetadataAttr
-KernelMetadataAttr::appendMetadata(ArrayRef<NamedAttribute> attrs) const {
- if (attrs.empty())
- return *this;
- NamedAttrList attrList;
- if (DictionaryAttr dict = getMetadata())
- attrList.append(dict);
- attrList.append(attrs);
- return KernelMetadataAttr::get(getName(), getFunctionType(), getArgAttrs(),
- attrList.getDictionary(getContext()));
-}
-
-LogicalResult
-KernelMetadataAttr::verify(function_ref<InFlightDiagnostic()> emitError,
- StringAttr name, Type functionType,
- ArrayAttr argAttrs, DictionaryAttr metadata) {
- if (name.empty())
- return emitError() << "the kernel name can't be empty";
- if (argAttrs) {
- if (llvm::any_of(argAttrs, [](Attribute attr) {
- return !llvm::isa<DictionaryAttr>(attr);
- }))
- return emitError()
- << "all attributes in the array must be a dictionary attribute";
- }
- return success();
-}
-
-//===----------------------------------------------------------------------===//
-// GPU KernelTableAttr
-//===----------------------------------------------------------------------===//
-
-KernelTableAttr KernelTableAttr::get(MLIRContext *context,
- ArrayRef<KernelMetadataAttr> kernels,
- bool isSorted) {
- // Note that `is_sorted` is always only invoked once even with assertions ON.
- assert((!isSorted || llvm::is_sorted(kernels)) &&
- "expected a sorted kernel array");
- // Immediately return the attribute if the array is sorted.
- if (isSorted || llvm::is_sorted(kernels))
- return Base::get(context, kernels);
- // Sort the array.
- SmallVector<KernelMetadataAttr> kernelsTmp(kernels);
- llvm::array_pod_sort(kernelsTmp.begin(), kernelsTmp.end());
- return Base::get(context, kernelsTmp);
-}
-
-KernelTableAttr KernelTableAttr::getChecked(
- function_ref<InFlightDiagnostic()> emitError, MLIRContext *context,
- ArrayRef<KernelMetadataAttr> kernels, bool isSorted) {
- // Note that `is_sorted` is always only invoked once even with assertions ON.
- assert((!isSorted || llvm::is_sorted(kernels)) &&
- "expected a sorted kernel array");
- // Immediately return the attribute if the array is sorted.
- if (isSorted || llvm::is_sorted(kernels))
- return Base::getChecked(emitError, context, kernels);
- // Sort the array.
- SmallVector<KernelMetadataAttr> kernelsTmp(kernels);
- llvm::array_pod_sort(kernelsTmp.begin(), kernelsTmp.end());
- return Base::getChecked(emitError, context, kernelsTmp);
-}
-
-LogicalResult
-KernelTableAttr::verify(function_ref<InFlightDiagnostic()> emitError,
- ArrayRef<KernelMetadataAttr> kernels) {
- if (kernels.size() < 2)
- return success();
- // Check that the kernels are uniquely named.
- if (std::adjacent_find(kernels.begin(), kernels.end(),
- [](KernelMetadataAttr l, KernelMetadataAttr r) {
- return l.getName() == r.getName();
- }) != kernels.end()) {
- return emitError() << "expected all kernels to be uniquely named";
- }
- return success();
-}
-
-KernelMetadataAttr KernelTableAttr::lookup(StringRef key) const {
- auto [iterator, found] = impl::findAttrSorted(begin(), end(), key);
- return found ? *iterator : KernelMetadataAttr();
-}
-
-KernelMetadataAttr KernelTableAttr::lookup(StringAttr key) const {
- auto [iterator, found] = impl::findAttrSorted(begin(), end(), key);
- return found ? *iterator : KernelMetadataAttr();
-}
-
-//===----------------------------------------------------------------------===//
-// GPU target options
-//===----------------------------------------------------------------------===//
-
-TargetOptions::TargetOptions(
- StringRef toolkitPath, ArrayRef<Attribute> librariesToLink,
- StringRef cmdOptions, StringRef elfSection,
- CompilationTarget compilationTarget,
- function_ref<SymbolTable *()> getSymbolTableCallback,
- function_ref<void(llvm::Module &)> initialLlvmIRCallback,
- function_ref<void(llvm::Module &)> linkedLlvmIRCallback,
- function_ref<void(llvm::Module &)> optimizedLlvmIRCallback,
- function_ref<void(StringRef)> isaCallback)
- : TargetOptions(TypeID::get<TargetOptions>(), toolkitPath, librariesToLink,
- cmdOptions, elfSection, compilationTarget,
- getSymbolTableCallback, initialLlvmIRCallback,
- linkedLlvmIRCallback, optimizedLlvmIRCallback,
- isaCallback) {}
-
-TargetOptions::TargetOptions(
- TypeID typeID, StringRef toolkitPath, ArrayRef<Attribute> librariesToLink,
- StringRef cmdOptions, StringRef elfSection,
- CompilationTarget compilationTarget,
- function_ref<SymbolTable *()> getSymbolTableCallback,
- function_ref<void(llvm::Module &)> initialLlvmIRCallback,
- function_ref<void(llvm::Module &)> linkedLlvmIRCallback,
- function_ref<void(llvm::Module &)> optimizedLlvmIRCallback,
- function_ref<void(StringRef)> isaCallback)
- : toolkitPath(toolkitPath.str()), librariesToLink(librariesToLink),
- cmdOptions(cmdOptions.str()), elfSection(elfSection.str()),
- compilationTarget(compilationTarget),
- getSymbolTableCallback(getSymbolTableCallback),
- initialLlvmIRCallback(initialLlvmIRCallback),
- linkedLlvmIRCallback(linkedLlvmIRCallback),
- optimizedLlvmIRCallback(optimizedLlvmIRCallback),
- isaCallback(isaCallback), typeID(typeID) {}
-
-TypeID TargetOptions::getTypeID() const { return typeID; }
-
-StringRef TargetOptions::getToolkitPath() const { return toolkitPath; }
-
-ArrayRef<Attribute> TargetOptions::getLibrariesToLink() const {
- return librariesToLink;
-}
-
-StringRef TargetOptions::getCmdOptions() const { return cmdOptions; }
-
-StringRef TargetOptions::getELFSection() const { return elfSection; }
-
-SymbolTable *TargetOptions::getSymbolTable() const {
- return getSymbolTableCallback ? getSymbolTableCallback() : nullptr;
-}
-
-function_ref<void(llvm::Module &)>
-TargetOptions::getInitialLlvmIRCallback() const {
- return initialLlvmIRCallback;
-}
-
-function_ref<void(llvm::Module &)>
-TargetOptions::getLinkedLlvmIRCallback() const {
- return linkedLlvmIRCallback;
-}
-
-function_ref<void(llvm::Module &)>
-TargetOptions::getOptimizedLlvmIRCallback() const {
- return optimizedLlvmIRCallback;
-}
-
-function_ref<void(StringRef)> TargetOptions::getISACallback() const {
- return isaCallback;
-}
-
-CompilationTarget TargetOptions::getCompilationTarget() const {
- return compilationTarget;
-}
-
-CompilationTarget TargetOptions::getDefaultCompilationTarget() {
- return CompilationTarget::Fatbin;
-}
-
-std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
-TargetOptions::tokenizeCmdOptions(const std::string &cmdOptions) {
- std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>> options;
- llvm::StringSaver stringSaver(options.first);
- StringRef opts = cmdOptions;
- // For a correct tokenization of the command line options `opts` must be
- // unquoted, otherwise the tokenization function returns a single string: the
- // unquoted `cmdOptions` -which is not the desired behavior.
- // Remove any quotes if they are at the beginning and end of the string:
- if (!opts.empty() && opts.front() == '"' && opts.back() == '"')
- opts.consume_front("\""), opts.consume_back("\"");
- if (!opts.empty() && opts.front() == '\'' && opts.back() == '\'')
- opts.consume_front("'"), opts.consume_back("'");
-#ifdef _WIN32
- llvm::cl::TokenizeWindowsCommandLine(opts, stringSaver, options.second,
- /*MarkEOLs=*/false);
-#else
- llvm::cl::TokenizeGNUCommandLine(opts, stringSaver, options.second,
- /*MarkEOLs=*/false);
-#endif // _WIN32
- return options;
-}
-
-std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
-TargetOptions::tokenizeCmdOptions() const {
- return tokenizeCmdOptions(cmdOptions);
-}
-
-std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
-TargetOptions::tokenizeAndRemoveSuffixCmdOptions(llvm::StringRef startsWith) {
- size_t startPos = cmdOptions.find(startsWith);
- if (startPos == std::string::npos)
- return {llvm::BumpPtrAllocator(), SmallVector<const char *>()};
-
- auto tokenized =
- tokenizeCmdOptions(cmdOptions.substr(startPos + startsWith.size()));
- cmdOptions.resize(startPos);
- return tokenized;
-}
-
-MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions)
-
#include "mlir/Dialect/GPU/IR/GPUOpInterfaces.cpp.inc"
-#include "mlir/Dialect/GPU/IR/GPUOpsEnums.cpp.inc"
-
-#define GET_ATTRDEF_CLASSES
-#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
#define GET_OP_CLASSES
#include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc"
-
-#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.cpp.inc"
diff --git a/mlir/lib/Dialect/OpenACC/CMakeLists.txt b/mlir/lib/Dialect/OpenACC/CMakeLists.txt
index e8a916e824d713..f8208e5a03ea91 100644
--- a/mlir/lib/Dialect/OpenACC/CMakeLists.txt
+++ b/mlir/lib/Dialect/OpenACC/CMakeLists.txt
@@ -1,4 +1,5 @@
add_subdirectory(Analysis)
add_subdirectory(IR)
+add_subdirectory(Extensions)
add_subdirectory(Utils)
add_subdirectory(Transforms)
diff --git a/mlir/lib/Dialect/OpenACC/Extensions/AllExtensions.cpp b/mlir/lib/Dialect/OpenACC/Extensions/AllExtensions.cpp
new file mode 100644
index 00000000000000..a584a7eef58a1b
--- /dev/null
+++ b/mlir/lib/Dialect/OpenACC/Extensions/AllExtensions.cpp
@@ -0,0 +1,19 @@
+//===- AllExtensions.cpp - Standard OpenACC external interfaces
+//------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
+#include "mlir/Dialect/OpenACC/Extensions/GPUInterfaces.h"
+#include "mlir/Dialect/OpenACC/Extensions/LLVMInterfaces.h"
+#include "mlir/Dialect/OpenACC/Extensions/MemRefInterfaces.h"
+
+void mlir::acc::registerAllExtensions(DialectRegistry ®istry) {
+ registerGPUInterfaces(registry);
+ registerLLVMInterfaces(registry);
+ registerMemRefInterfaces(registry);
+}
diff --git a/mlir/lib/Dialect/OpenACC/Extensions/CMakeLists.txt b/mlir/lib/Dialect/OpenACC/Extensions/CMakeLists.txt
new file mode 100644
index 00000000000000..6e06d5b8d9d05e
--- /dev/null
+++ b/mlir/lib/Dialect/OpenACC/Extensions/CMakeLists.txt
@@ -0,0 +1,41 @@
+set(LLVM_OPTIONAL_SOURCES
+ AllExtensions.cpp
+ GPUInterfaces.cpp
+ LLVMInterfaces.cpp
+ MemRefInterfaces.cpp
+ )
+
+add_mlir_extension_library(MLIROpenACCMemRefInterfaces
+ MemRefInterfaces.cpp
+
+ LINK_LIBS PUBLIC
+ MLIRArithDialect
+ MLIROpenACCDialect
+ MLIRMemRefDialect
+ )
+
+add_mlir_extension_library(MLIROpenACCLLVMInterfaces
+ LLVMInterfaces.cpp
+
+ LINK_LIBS PUBLIC
+ MLIRArithDialect
+ MLIROpenACCDialect
+ MLIRLLVMDialect
+ )
+
+add_mlir_extension_library(MLIROpenACCGPUInterfaces
+ GPUInterfaces.cpp
+
+ LINK_LIBS PUBLIC
+ MLIROpenACCDialect
+ MLIRGPUDialect
+ )
+
+add_mlir_extension_library(MLIROpenACCAllExtensions
+ AllExtensions.cpp
+
+ LINK_LIBS PUBLIC
+ MLIROpenACCGPUInterfaces
+ MLIROpenACCLLVMInterfaces
+ MLIROpenACCMemRefInterfaces
+ )
diff --git a/mlir/lib/Dialect/OpenACC/Extensions/GPUInterfaces.cpp b/mlir/lib/Dialect/OpenACC/Extensions/GPUInterfaces.cpp
new file mode 100644
index 00000000000000..92a711b3a8834d
--- /dev/null
+++ b/mlir/lib/Dialect/OpenACC/Extensions/GPUInterfaces.cpp
@@ -0,0 +1,32 @@
+//===- GPUInterfaces.cpp - OpenACC external models --------------------===//
+//
+// Part of the MLIR 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/OpenACC/Extensions/GPUInterfaces.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/DialectRegistry.h"
+
+using namespace mlir;
+using namespace mlir::acc;
+
+namespace {
+struct GPULaunchOffloadRegionModel
+ : public acc::OffloadRegionOpInterface::ExternalModel<
+ GPULaunchOffloadRegionModel, gpu::LaunchOp> {
+ mlir::Region &getOffloadRegion(mlir::Operation *op) const {
+ return cast<gpu::LaunchOp>(op).getBody();
+ }
+};
+} // namespace
+
+void mlir::acc::registerGPUInterfaces(DialectRegistry ®istry) {
+ registry.addExtension(
+ +[](MLIRContext *ctx, OpenACCDialect *, gpu::GPUDialect *) {
+ gpu::LaunchOp::attachInterface<GPULaunchOffloadRegionModel>(*ctx);
+ });
+}
diff --git a/mlir/lib/Dialect/OpenACC/Extensions/LLVMInterfaces.cpp b/mlir/lib/Dialect/OpenACC/Extensions/LLVMInterfaces.cpp
new file mode 100644
index 00000000000000..4d99f31d6edc7f
--- /dev/null
+++ b/mlir/lib/Dialect/OpenACC/Extensions/LLVMInterfaces.cpp
@@ -0,0 +1,98 @@
+//===- LLVMInterfaces.cpp - OpenACC external models --------------------===//
+//
+// Part of the MLIR 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/OpenACC/Extensions/LLVMInterfaces.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/DialectRegistry.h"
+
+using namespace mlir;
+using namespace mlir::acc;
+
+namespace {
+struct LLVMPointerPointerLikeModel
+ : public PointerLikeType::ExternalModel<LLVMPointerPointerLikeModel,
+ LLVM::LLVMPointerType> {
+ Type getElementType(Type pointer) const { return Type(); }
+
+ mlir::Value genLoad(Type pointer, OpBuilder &builder, Location loc,
+ TypedValue<PointerLikeType> srcPtr,
+ Type valueType) const {
+ // For LLVM pointers, we need the valueType to determine what to load
+ if (!valueType)
+ return {};
+
+ return LLVM::LoadOp::create(builder, loc, valueType, srcPtr);
+ }
+
+ bool genStore(Type pointer, OpBuilder &builder, Location loc,
+ Value valueToStore, TypedValue<PointerLikeType> destPtr) const {
+ LLVM::StoreOp::create(builder, loc, valueToStore, destPtr);
+ return true;
+ }
+
+ Value genCast(Type, OpBuilder &builder, Location loc, Value value,
+ Type resultType) const {
+ if (value.getType() == resultType)
+ return value;
+
+ auto srcPtrTy = dyn_cast<LLVM::LLVMPointerType>(value.getType());
+ auto dstPtrTy = dyn_cast<LLVM::LLVMPointerType>(resultType);
+ if (srcPtrTy && dstPtrTy) {
+ if (srcPtrTy.getAddressSpace() != dstPtrTy.getAddressSpace())
+ return LLVM::AddrSpaceCastOp::create(builder, loc, resultType, value);
+ return value;
+ }
+
+ if (srcPtrTy && isa<IntegerType>(resultType))
+ return LLVM::PtrToIntOp::create(builder, loc, resultType, value);
+
+ if (dstPtrTy) {
+ Value intVal = value;
+ if (isa<IndexType>(value.getType()))
+ intVal = arith::IndexCastUIOp::create(builder, loc,
+ builder.getI64Type(), value);
+ if (isa<IntegerType>(intVal.getType()))
+ return LLVM::IntToPtrOp::create(builder, loc, resultType, intVal);
+ }
+
+ if (auto resPtrLike = dyn_cast<PointerLikeType>(resultType))
+ if (!isa<LLVM::LLVMPointerType>(resPtrLike))
+ if (Value v = resPtrLike.genCast(builder, loc, value, resultType))
+ return v;
+ if (auto valPtrLike = dyn_cast<PointerLikeType>(value.getType()))
+ if (!isa<LLVM::LLVMPointerType>(valPtrLike))
+ if (Value v = valPtrLike.genCast(builder, loc, value, resultType))
+ return v;
+
+ return UnrealizedConversionCastOp::create(builder, loc,
+ TypeRange(resultType), value)
+ .getResult(0);
+ }
+};
+
+struct LLVMAddressOfGlobalModel
+ : public AddressOfGlobalOpInterface::ExternalModel<LLVMAddressOfGlobalModel,
+ LLVM::AddressOfOp> {
+ SymbolRefAttr getSymbol(Operation *op) const {
+ auto addressOfOp = cast<LLVM::AddressOfOp>(op);
+ return addressOfOp.getGlobalNameAttr();
+ }
+};
+} // namespace
+
+void mlir::acc::registerLLVMInterfaces(DialectRegistry ®istry) {
+ registry.addExtension(+[](MLIRContext *ctx, OpenACCDialect *) {
+ ctx->getOrLoadDialect<LLVM::LLVMDialect>();
+ ctx->getOrLoadDialect<arith::ArithDialect>();
+ LLVM::LLVMPointerType::attachInterface<LLVMPointerPointerLikeModel>(*ctx);
+ LLVM::AddressOfOp::attachInterface<LLVMAddressOfGlobalModel>(*ctx);
+ });
+}
diff --git a/mlir/lib/Dialect/OpenACC/Extensions/MemRefInterfaces.cpp b/mlir/lib/Dialect/OpenACC/Extensions/MemRefInterfaces.cpp
new file mode 100644
index 00000000000000..49bb69e1c43c73
--- /dev/null
+++ b/mlir/lib/Dialect/OpenACC/Extensions/MemRefInterfaces.cpp
@@ -0,0 +1,328 @@
+//===- MemRefInterfaces.cpp - OpenACC external models --------------------===//
+//
+// Part of the MLIR 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/OpenACC/Extensions/MemRefInterfaces.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/GPU/IR/GPUBase.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/DialectRegistry.h"
+
+using namespace mlir;
+using namespace mlir::acc;
+
+namespace {
+static bool isScalarLikeType(Type type) {
+ return type.isIntOrIndexOrFloat() || isa<ComplexType>(type);
+}
+
+/// Helper function to attach the `VarName` attribute to an operation
+/// if a variable name is provided.
+static void attachVarNameAttr(Operation *op, OpBuilder &builder,
+ StringRef varName) {
+ if (!varName.empty()) {
+ auto varNameAttr = acc::VarNameAttr::get(builder.getContext(), varName);
+ op->setDiscardableAttr(acc::getVarNameAttrName(), varNameAttr);
+ }
+}
+
+template <typename T>
+struct MemRefPointerLikeModel
+ : public PointerLikeType::ExternalModel<MemRefPointerLikeModel<T>, T> {
+ Type getElementType(Type pointer) const {
+ return cast<T>(pointer).getElementType();
+ }
+
+ mlir::acc::VariableTypeCategory
+ getPointeeTypeCategory(Type pointer, TypedValue<PointerLikeType> varPtr,
+ Type varType) const {
+ if (auto mappableTy = dyn_cast<MappableType>(varType)) {
+ return mappableTy.getTypeCategory(varPtr);
+ }
+ auto memrefTy = cast<T>(pointer);
+ if (!memrefTy.hasRank()) {
+ // This memref is unranked - aka it could have any rank, including a
+ // rank of 0 which could mean scalar. For now, return uncategorized.
+ return mlir::acc::VariableTypeCategory::uncategorized;
+ }
+
+ if (memrefTy.getRank() == 0) {
+ if (isScalarLikeType(memrefTy.getElementType())) {
+ return mlir::acc::VariableTypeCategory::scalar;
+ }
+ // Zero-rank non-scalar - need further analysis to determine the type
+ // category. For now, return uncategorized.
+ return mlir::acc::VariableTypeCategory::uncategorized;
+ }
+
+ // It has a rank - must be an array.
+ assert(memrefTy.getRank() > 0 && "rank expected to be positive");
+ return mlir::acc::VariableTypeCategory::array;
+ }
+
+ mlir::Value genAllocate(Type pointer, OpBuilder &builder, Location loc,
+ StringRef varName, Type varType, Value originalVar,
+ bool &needsFree) const {
+ auto memrefTy = cast<MemRefType>(pointer);
+
+ // Check if this is a static memref (all dimensions are known) - if yes
+ // then we can generate an alloca operation.
+ if (memrefTy.hasStaticShape()) {
+ needsFree = false; // alloca doesn't need deallocation
+ auto allocaOp = memref::AllocaOp::create(builder, loc, memrefTy);
+ attachVarNameAttr(allocaOp, builder, varName);
+ return allocaOp.getResult();
+ }
+
+ // For dynamic memrefs, extract sizes from the original variable if
+ // provided. Otherwise they cannot be handled.
+ if (originalVar && originalVar.getType() == memrefTy &&
+ memrefTy.hasRank()) {
+ SmallVector<Value> dynamicSizes;
+ for (int64_t i = 0; i < memrefTy.getRank(); ++i) {
+ if (memrefTy.isDynamicDim(i)) {
+ // Extract the size of dimension i from the original variable
+ auto indexValue = arith::ConstantIndexOp::create(builder, loc, i);
+ auto dimSize =
+ memref::DimOp::create(builder, loc, originalVar, indexValue);
+ dynamicSizes.push_back(dimSize);
+ }
+ // Note: We only add dynamic sizes to the dynamicSizes array
+ // Static dimensions are handled automatically by AllocOp
+ }
+ needsFree = true; // alloc needs deallocation
+ auto allocOp =
+ memref::AllocOp::create(builder, loc, memrefTy, dynamicSizes);
+ attachVarNameAttr(allocOp, builder, varName);
+ return allocOp.getResult();
+ }
+
+ // TODO: Unranked not yet supported.
+ return {};
+ }
+
+ bool genFree(Type pointer, OpBuilder &builder, Location loc,
+ TypedValue<PointerLikeType> varToFree, Value allocRes,
+ Type varType) const {
+ if (auto memrefValue = dyn_cast<TypedValue<MemRefType>>(varToFree)) {
+ // Use allocRes if provided to determine the allocation type
+ Value valueToInspect = allocRes ? allocRes : memrefValue;
+
+ // Walk through casts to find the original allocation
+ Value currentValue = valueToInspect;
+ Operation *originalAlloc = nullptr;
+
+ // Follow the chain of operations to find the original allocation
+ // even if a casted result is provided.
+ while (currentValue) {
+ if (auto *definingOp = currentValue.getDefiningOp()) {
+ // Check if this is an allocation operation
+ if (isa<memref::AllocOp, memref::AllocaOp>(definingOp)) {
+ originalAlloc = definingOp;
+ break;
+ }
+
+ // Check if this is a cast operation we can look through
+ if (auto castOp = dyn_cast<memref::CastOp>(definingOp)) {
+ currentValue = castOp.getSource();
+ continue;
+ }
+
+ // Check for other cast-like operations
+ if (auto reinterpretCastOp =
+ dyn_cast<memref::ReinterpretCastOp>(definingOp)) {
+ currentValue = reinterpretCastOp.getSource();
+ continue;
+ }
+
+ // If we can't look through this operation, stop
+ break;
+ }
+ // This is a block argument or similar - can't trace further.
+ break;
+ }
+
+ if (originalAlloc) {
+ if (isa<memref::AllocaOp>(originalAlloc)) {
+ // This is an alloca - no dealloc needed, but return true (success)
+ return true;
+ }
+ if (isa<memref::AllocOp>(originalAlloc)) {
+ // This is an alloc - generate dealloc on varToFree
+ memref::DeallocOp::create(builder, loc, memrefValue);
+ return true;
+ }
+ }
+ }
+
+ return false;
+ }
+
+ bool genCopy(Type pointer, OpBuilder &builder, Location loc,
+ TypedValue<PointerLikeType> destination,
+ TypedValue<PointerLikeType> source, Type varType) const {
+ // Generate a copy operation between two memrefs
+ auto destMemref = dyn_cast_if_present<TypedValue<MemRefType>>(destination);
+ auto srcMemref = dyn_cast_if_present<TypedValue<MemRefType>>(source);
+
+ // As per memref documentation, source and destination must have same
+ // element type and shape in order to be compatible. We do not want to fail
+ // with an IR verification error - thus check that before generating the
+ // copy operation.
+ if (destMemref && srcMemref &&
+ destMemref.getType().getElementType() ==
+ srcMemref.getType().getElementType() &&
+ destMemref.getType().getShape() == srcMemref.getType().getShape()) {
+ memref::CopyOp::create(builder, loc, srcMemref, destMemref);
+ return true;
+ }
+
+ return false;
+ }
+
+ mlir::Value genLoad(Type pointer, OpBuilder &builder, Location loc,
+ TypedValue<PointerLikeType> srcPtr,
+ Type valueType) const {
+ // Load from a memref - only valid for scalar memrefs (rank 0).
+ // This is because the address computation for memrefs is part of the load
+ // (and not computed separately), but the API does not have arguments for
+ // indexing.
+ auto memrefValue = dyn_cast_if_present<TypedValue<MemRefType>>(srcPtr);
+ if (!memrefValue)
+ return {};
+
+ auto memrefTy = memrefValue.getType();
+
+ // Only load from scalar memrefs (rank 0)
+ if (memrefTy.getRank() != 0)
+ return {};
+
+ return memref::LoadOp::create(builder, loc, memrefValue, ValueRange{});
+ }
+
+ bool genStore(Type pointer, OpBuilder &builder, Location loc,
+ Value valueToStore, TypedValue<PointerLikeType> destPtr) const {
+ // Store to a memref - only valid for scalar memrefs (rank 0)
+ // This is because the address computation for memrefs is part of the store
+ // (and not computed separately), but the API does not have arguments for
+ // indexing.
+ auto memrefValue = dyn_cast_if_present<TypedValue<MemRefType>>(destPtr);
+ if (!memrefValue)
+ return false;
+
+ auto memrefTy = memrefValue.getType();
+
+ // Only store to scalar memrefs (rank 0)
+ if (memrefTy.getRank() != 0)
+ return false;
+
+ memref::StoreOp::create(builder, loc, valueToStore, memrefValue);
+ return true;
+ }
+
+ Value genCast(Type, OpBuilder &builder, Location loc, Value value,
+ Type resultType) const {
+ if (value.getType() == resultType)
+ return value;
+
+ if (isa<BaseMemRefType>(value.getType()) &&
+ isa<BaseMemRefType>(resultType)) {
+ if (memref::CastOp::areCastCompatible(TypeRange(value.getType()),
+ TypeRange(resultType)))
+ return memref::CastOp::create(builder, loc, resultType, value);
+ if (memref::MemorySpaceCastOp::areCastCompatible(
+ TypeRange(value.getType()), TypeRange(resultType)))
+ return memref::MemorySpaceCastOp::create(builder, loc, resultType,
+ value);
+ }
+
+ // If one side is not a memref, try the other type's `PointerLikeType`
+ // implementation (since it may be an out-of-tree reference type that
+ // we cannot generate here).
+ if (auto resPtrLike = dyn_cast<PointerLikeType>(resultType))
+ if (!isa<BaseMemRefType>(resPtrLike))
+ if (Value v = resPtrLike.genCast(builder, loc, value, resultType))
+ return v;
+ if (auto valPtrLike = dyn_cast<PointerLikeType>(value.getType()))
+ if (!isa<BaseMemRefType>(valPtrLike))
+ if (Value v = valPtrLike.genCast(builder, loc, value, resultType))
+ return v;
+
+ return {};
+ }
+
+ bool isDeviceAccessible(Type pointer, Value var) const {
+ auto memrefTy = cast<T>(pointer);
+ Attribute memSpace = memrefTy.getMemorySpace();
+ return isa_and_nonnull<gpu::AddressSpaceAttr>(memSpace);
+ }
+
+ MemRefType getAsMemRefType(Type pointer, ModuleOp module) const {
+ (void)module;
+ return dyn_cast<MemRefType>(pointer);
+ }
+};
+
+struct MemrefAddressOfGlobalModel
+ : public AddressOfGlobalOpInterface::ExternalModel<
+ MemrefAddressOfGlobalModel, memref::GetGlobalOp> {
+ SymbolRefAttr getSymbol(Operation *op) const {
+ auto getGlobalOp = cast<memref::GetGlobalOp>(op);
+ return getGlobalOp.getNameAttr();
+ }
+};
+
+struct MemrefGlobalVariableModel
+ : public GlobalVariableOpInterface::ExternalModel<MemrefGlobalVariableModel,
+ memref::GlobalOp> {
+ bool isConstant(Operation *op) const {
+ auto globalOp = cast<memref::GlobalOp>(op);
+ return globalOp.getConstant();
+ }
+
+ bool hasInitializer(Operation *op) const {
+ auto globalOp = cast<memref::GlobalOp>(op);
+ return globalOp.getInitialValue().has_value();
+ }
+
+ Region *getInitRegion(Operation *op) const {
+ // GlobalOp uses attributes for initialization, not regions
+ return nullptr;
+ }
+
+ bool isDeviceAccessible(Operation *op) const {
+ auto globalOp = cast<memref::GlobalOp>(op);
+ Attribute memSpace = globalOp.getType().getMemorySpace();
+ return isa_and_nonnull<gpu::AddressSpaceAttr>(memSpace);
+ }
+
+ bool isInDeviceMemory(Operation *op) const {
+ // A memref address space models storage that is physically resident on the
+ // device, so a device-accessible global is also in device memory. (There
+ // is no host-shared/migratable address space to exclude here.)
+ return isDeviceAccessible(op);
+ }
+
+ bool isCompilerGenerated(Operation *op) const { return false; }
+};
+} // namespace
+
+void mlir::acc::registerMemRefInterfaces(DialectRegistry ®istry) {
+ // MemRef types belong to BuiltinDialect, so IR can use them without any
+ // memref operations. Install their models as soon as OpenACC is loaded.
+ registry.addExtension(+[](MLIRContext *ctx, OpenACCDialect *) {
+ ctx->getOrLoadDialect<memref::MemRefDialect>();
+ ctx->getOrLoadDialect<arith::ArithDialect>();
+ MemRefType::attachInterface<MemRefPointerLikeModel<MemRefType>>(*ctx);
+ UnrankedMemRefType::attachInterface<
+ MemRefPointerLikeModel<UnrankedMemRefType>>(*ctx);
+ memref::GetGlobalOp::attachInterface<MemrefAddressOfGlobalModel>(*ctx);
+ memref::GlobalOp::attachInterface<MemrefGlobalVariableModel>(*ctx);
+ });
+}
diff --git a/mlir/lib/Dialect/OpenACC/IR/CMakeLists.txt b/mlir/lib/Dialect/OpenACC/IR/CMakeLists.txt
index 409155ec18fe6e..881628c5ab98ce 100644
--- a/mlir/lib/Dialect/OpenACC/IR/CMakeLists.txt
+++ b/mlir/lib/Dialect/OpenACC/IR/CMakeLists.txt
@@ -16,9 +16,7 @@ add_mlir_dialect_library(MLIROpenACCDialect
LINK_LIBS PUBLIC
MLIRIR
MLIRDialectUtils
- MLIRGPUDialect
- MLIRLLVMDialect
- MLIRMemRefDialect
+ MLIRGPUBase
MLIRTransformUtils
MLIROpenACCMPCommon
)
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 4f0aec22a00b23..a9c4b738af0b0f 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -8,10 +8,6 @@
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
-#include "mlir/Dialect/GPU/IR/GPUDialect.h"
-#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/LLVMIR/LLVMTypes.h"
-#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/Utils/StaticValueUtils.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinAttributes.h"
@@ -39,319 +35,6 @@ using namespace acc;
#include "mlir/Dialect/OpenACCMPCommon/Interfaces/OpenACCMPOpsInterfaces.cpp.inc"
namespace {
-
-static bool isScalarLikeType(Type type) {
- return type.isIntOrIndexOrFloat() || isa<ComplexType>(type);
-}
-
-/// Helper function to attach the `VarName` attribute to an operation
-/// if a variable name is provided.
-static void attachVarNameAttr(Operation *op, OpBuilder &builder,
- StringRef varName) {
- if (!varName.empty()) {
- auto varNameAttr = acc::VarNameAttr::get(builder.getContext(), varName);
- op->setDiscardableAttr(acc::getVarNameAttrName(), varNameAttr);
- }
-}
-
-template <typename T>
-struct MemRefPointerLikeModel
- : public PointerLikeType::ExternalModel<MemRefPointerLikeModel<T>, T> {
- Type getElementType(Type pointer) const {
- return cast<T>(pointer).getElementType();
- }
-
- mlir::acc::VariableTypeCategory
- getPointeeTypeCategory(Type pointer, TypedValue<PointerLikeType> varPtr,
- Type varType) const {
- if (auto mappableTy = dyn_cast<MappableType>(varType)) {
- return mappableTy.getTypeCategory(varPtr);
- }
- auto memrefTy = cast<T>(pointer);
- if (!memrefTy.hasRank()) {
- // This memref is unranked - aka it could have any rank, including a
- // rank of 0 which could mean scalar. For now, return uncategorized.
- return mlir::acc::VariableTypeCategory::uncategorized;
- }
-
- if (memrefTy.getRank() == 0) {
- if (isScalarLikeType(memrefTy.getElementType())) {
- return mlir::acc::VariableTypeCategory::scalar;
- }
- // Zero-rank non-scalar - need further analysis to determine the type
- // category. For now, return uncategorized.
- return mlir::acc::VariableTypeCategory::uncategorized;
- }
-
- // It has a rank - must be an array.
- assert(memrefTy.getRank() > 0 && "rank expected to be positive");
- return mlir::acc::VariableTypeCategory::array;
- }
-
- mlir::Value genAllocate(Type pointer, OpBuilder &builder, Location loc,
- StringRef varName, Type varType, Value originalVar,
- bool &needsFree) const {
- auto memrefTy = cast<MemRefType>(pointer);
-
- // Check if this is a static memref (all dimensions are known) - if yes
- // then we can generate an alloca operation.
- if (memrefTy.hasStaticShape()) {
- needsFree = false; // alloca doesn't need deallocation
- auto allocaOp = memref::AllocaOp::create(builder, loc, memrefTy);
- attachVarNameAttr(allocaOp, builder, varName);
- return allocaOp.getResult();
- }
-
- // For dynamic memrefs, extract sizes from the original variable if
- // provided. Otherwise they cannot be handled.
- if (originalVar && originalVar.getType() == memrefTy &&
- memrefTy.hasRank()) {
- SmallVector<Value> dynamicSizes;
- for (int64_t i = 0; i < memrefTy.getRank(); ++i) {
- if (memrefTy.isDynamicDim(i)) {
- // Extract the size of dimension i from the original variable
- auto indexValue = arith::ConstantIndexOp::create(builder, loc, i);
- auto dimSize =
- memref::DimOp::create(builder, loc, originalVar, indexValue);
- dynamicSizes.push_back(dimSize);
- }
- // Note: We only add dynamic sizes to the dynamicSizes array
- // Static dimensions are handled automatically by AllocOp
- }
- needsFree = true; // alloc needs deallocation
- auto allocOp =
- memref::AllocOp::create(builder, loc, memrefTy, dynamicSizes);
- attachVarNameAttr(allocOp, builder, varName);
- return allocOp.getResult();
- }
-
- // TODO: Unranked not yet supported.
- return {};
- }
-
- bool genFree(Type pointer, OpBuilder &builder, Location loc,
- TypedValue<PointerLikeType> varToFree, Value allocRes,
- Type varType) const {
- if (auto memrefValue = dyn_cast<TypedValue<MemRefType>>(varToFree)) {
- // Use allocRes if provided to determine the allocation type
- Value valueToInspect = allocRes ? allocRes : memrefValue;
-
- // Walk through casts to find the original allocation
- Value currentValue = valueToInspect;
- Operation *originalAlloc = nullptr;
-
- // Follow the chain of operations to find the original allocation
- // even if a casted result is provided.
- while (currentValue) {
- if (auto *definingOp = currentValue.getDefiningOp()) {
- // Check if this is an allocation operation
- if (isa<memref::AllocOp, memref::AllocaOp>(definingOp)) {
- originalAlloc = definingOp;
- break;
- }
-
- // Check if this is a cast operation we can look through
- if (auto castOp = dyn_cast<memref::CastOp>(definingOp)) {
- currentValue = castOp.getSource();
- continue;
- }
-
- // Check for other cast-like operations
- if (auto reinterpretCastOp =
- dyn_cast<memref::ReinterpretCastOp>(definingOp)) {
- currentValue = reinterpretCastOp.getSource();
- continue;
- }
-
- // If we can't look through this operation, stop
- break;
- }
- // This is a block argument or similar - can't trace further.
- break;
- }
-
- if (originalAlloc) {
- if (isa<memref::AllocaOp>(originalAlloc)) {
- // This is an alloca - no dealloc needed, but return true (success)
- return true;
- }
- if (isa<memref::AllocOp>(originalAlloc)) {
- // This is an alloc - generate dealloc on varToFree
- memref::DeallocOp::create(builder, loc, memrefValue);
- return true;
- }
- }
- }
-
- return false;
- }
-
- bool genCopy(Type pointer, OpBuilder &builder, Location loc,
- TypedValue<PointerLikeType> destination,
- TypedValue<PointerLikeType> source, Type varType) const {
- // Generate a copy operation between two memrefs
- auto destMemref = dyn_cast_if_present<TypedValue<MemRefType>>(destination);
- auto srcMemref = dyn_cast_if_present<TypedValue<MemRefType>>(source);
-
- // As per memref documentation, source and destination must have same
- // element type and shape in order to be compatible. We do not want to fail
- // with an IR verification error - thus check that before generating the
- // copy operation.
- if (destMemref && srcMemref &&
- destMemref.getType().getElementType() ==
- srcMemref.getType().getElementType() &&
- destMemref.getType().getShape() == srcMemref.getType().getShape()) {
- memref::CopyOp::create(builder, loc, srcMemref, destMemref);
- return true;
- }
-
- return false;
- }
-
- mlir::Value genLoad(Type pointer, OpBuilder &builder, Location loc,
- TypedValue<PointerLikeType> srcPtr,
- Type valueType) const {
- // Load from a memref - only valid for scalar memrefs (rank 0).
- // This is because the address computation for memrefs is part of the load
- // (and not computed separately), but the API does not have arguments for
- // indexing.
- auto memrefValue = dyn_cast_if_present<TypedValue<MemRefType>>(srcPtr);
- if (!memrefValue)
- return {};
-
- auto memrefTy = memrefValue.getType();
-
- // Only load from scalar memrefs (rank 0)
- if (memrefTy.getRank() != 0)
- return {};
-
- return memref::LoadOp::create(builder, loc, memrefValue, ValueRange{});
- }
-
- bool genStore(Type pointer, OpBuilder &builder, Location loc,
- Value valueToStore, TypedValue<PointerLikeType> destPtr) const {
- // Store to a memref - only valid for scalar memrefs (rank 0)
- // This is because the address computation for memrefs is part of the store
- // (and not computed separately), but the API does not have arguments for
- // indexing.
- auto memrefValue = dyn_cast_if_present<TypedValue<MemRefType>>(destPtr);
- if (!memrefValue)
- return false;
-
- auto memrefTy = memrefValue.getType();
-
- // Only store to scalar memrefs (rank 0)
- if (memrefTy.getRank() != 0)
- return false;
-
- memref::StoreOp::create(builder, loc, valueToStore, memrefValue);
- return true;
- }
-
- Value genCast(Type, OpBuilder &builder, Location loc, Value value,
- Type resultType) const {
- if (value.getType() == resultType)
- return value;
-
- if (isa<BaseMemRefType>(value.getType()) &&
- isa<BaseMemRefType>(resultType)) {
- if (memref::CastOp::areCastCompatible(TypeRange(value.getType()),
- TypeRange(resultType)))
- return memref::CastOp::create(builder, loc, resultType, value);
- if (memref::MemorySpaceCastOp::areCastCompatible(
- TypeRange(value.getType()), TypeRange(resultType)))
- return memref::MemorySpaceCastOp::create(builder, loc, resultType,
- value);
- }
-
- // If one side is not a memref, try the other type's `PointerLikeType`
- // implementation (since it may be an out-of-tree reference type that
- // we cannot generate here).
- if (auto resPtrLike = dyn_cast<PointerLikeType>(resultType))
- if (!isa<BaseMemRefType>(resPtrLike))
- if (Value v = resPtrLike.genCast(builder, loc, value, resultType))
- return v;
- if (auto valPtrLike = dyn_cast<PointerLikeType>(value.getType()))
- if (!isa<BaseMemRefType>(valPtrLike))
- if (Value v = valPtrLike.genCast(builder, loc, value, resultType))
- return v;
-
- return {};
- }
-
- bool isDeviceAccessible(Type pointer, Value var) const {
- auto memrefTy = cast<T>(pointer);
- Attribute memSpace = memrefTy.getMemorySpace();
- return isa_and_nonnull<gpu::AddressSpaceAttr>(memSpace);
- }
-
- MemRefType getAsMemRefType(Type pointer, ModuleOp module) const {
- (void)module;
- return dyn_cast<MemRefType>(pointer);
- }
-};
-
-struct LLVMPointerPointerLikeModel
- : public PointerLikeType::ExternalModel<LLVMPointerPointerLikeModel,
- LLVM::LLVMPointerType> {
- Type getElementType(Type pointer) const { return Type(); }
-
- mlir::Value genLoad(Type pointer, OpBuilder &builder, Location loc,
- TypedValue<PointerLikeType> srcPtr,
- Type valueType) const {
- // For LLVM pointers, we need the valueType to determine what to load
- if (!valueType)
- return {};
-
- return LLVM::LoadOp::create(builder, loc, valueType, srcPtr);
- }
-
- bool genStore(Type pointer, OpBuilder &builder, Location loc,
- Value valueToStore, TypedValue<PointerLikeType> destPtr) const {
- LLVM::StoreOp::create(builder, loc, valueToStore, destPtr);
- return true;
- }
-
- Value genCast(Type, OpBuilder &builder, Location loc, Value value,
- Type resultType) const {
- if (value.getType() == resultType)
- return value;
-
- auto srcPtrTy = dyn_cast<LLVM::LLVMPointerType>(value.getType());
- auto dstPtrTy = dyn_cast<LLVM::LLVMPointerType>(resultType);
- if (srcPtrTy && dstPtrTy) {
- if (srcPtrTy.getAddressSpace() != dstPtrTy.getAddressSpace())
- return LLVM::AddrSpaceCastOp::create(builder, loc, resultType, value);
- return value;
- }
-
- if (srcPtrTy && isa<IntegerType>(resultType))
- return LLVM::PtrToIntOp::create(builder, loc, resultType, value);
-
- if (dstPtrTy) {
- Value intVal = value;
- if (isa<IndexType>(value.getType()))
- intVal = arith::IndexCastUIOp::create(builder, loc,
- builder.getI64Type(), value);
- if (isa<IntegerType>(intVal.getType()))
- return LLVM::IntToPtrOp::create(builder, loc, resultType, intVal);
- }
-
- if (auto resPtrLike = dyn_cast<PointerLikeType>(resultType))
- if (!isa<LLVM::LLVMPointerType>(resPtrLike))
- if (Value v = resPtrLike.genCast(builder, loc, value, resultType))
- return v;
- if (auto valPtrLike = dyn_cast<PointerLikeType>(value.getType()))
- if (!isa<LLVM::LLVMPointerType>(valPtrLike))
- if (Value v = valPtrLike.genCast(builder, loc, value, resultType))
- return v;
-
- return UnrealizedConversionCastOp::create(builder, loc,
- TypeRange(resultType), value)
- .getResult(0);
- }
-};
-
struct PrivateTypePointerLikeModel
: public PointerLikeType::ExternalModel<PrivateTypePointerLikeModel,
PrivateType> {
@@ -378,66 +61,6 @@ struct PrivateTypePointerLikeModel
}
};
-struct MemrefAddressOfGlobalModel
- : public AddressOfGlobalOpInterface::ExternalModel<
- MemrefAddressOfGlobalModel, memref::GetGlobalOp> {
- SymbolRefAttr getSymbol(Operation *op) const {
- auto getGlobalOp = cast<memref::GetGlobalOp>(op);
- return getGlobalOp.getNameAttr();
- }
-};
-
-struct LLVMAddressOfGlobalModel
- : public AddressOfGlobalOpInterface::ExternalModel<LLVMAddressOfGlobalModel,
- LLVM::AddressOfOp> {
- SymbolRefAttr getSymbol(Operation *op) const {
- auto addressOfOp = cast<LLVM::AddressOfOp>(op);
- return addressOfOp.getGlobalNameAttr();
- }
-};
-
-struct MemrefGlobalVariableModel
- : public GlobalVariableOpInterface::ExternalModel<MemrefGlobalVariableModel,
- memref::GlobalOp> {
- bool isConstant(Operation *op) const {
- auto globalOp = cast<memref::GlobalOp>(op);
- return globalOp.getConstant();
- }
-
- bool hasInitializer(Operation *op) const {
- auto globalOp = cast<memref::GlobalOp>(op);
- return globalOp.getInitialValue().has_value();
- }
-
- Region *getInitRegion(Operation *op) const {
- // GlobalOp uses attributes for initialization, not regions
- return nullptr;
- }
-
- bool isDeviceAccessible(Operation *op) const {
- auto globalOp = cast<memref::GlobalOp>(op);
- Attribute memSpace = globalOp.getType().getMemorySpace();
- return isa_and_nonnull<gpu::AddressSpaceAttr>(memSpace);
- }
-
- bool isInDeviceMemory(Operation *op) const {
- // A memref address space models storage that is physically resident on the
- // device, so a device-accessible global is also in device memory. (There
- // is no host-shared/migratable address space to exclude here.)
- return isDeviceAccessible(op);
- }
-
- bool isCompilerGenerated(Operation *op) const { return false; }
-};
-
-struct GPULaunchOffloadRegionModel
- : public acc::OffloadRegionOpInterface::ExternalModel<
- GPULaunchOffloadRegionModel, gpu::LaunchOp> {
- mlir::Region &getOffloadRegion(mlir::Operation *op) const {
- return cast<gpu::LaunchOp>(op).getBody();
- }
-};
-
/// Helper function for any of the times we need to modify an ArrayAttr based on
/// a device type list. Returns a new ArrayAttr with all of the
/// existingDeviceTypes, plus the effective new ones(or an added none if hte new
@@ -520,23 +143,7 @@ void OpenACCDialect::initialize() {
#include "mlir/Dialect/OpenACC/OpenACCOpsTypes.cpp.inc"
>();
- // By attaching interfaces here, we make the OpenACC dialect dependent on
- // the other dialects. This is probably better than having dialects like LLVM
- // and memref be dependent on OpenACC.
- MemRefType::attachInterface<MemRefPointerLikeModel<MemRefType>>(
- *getContext());
- UnrankedMemRefType::attachInterface<
- MemRefPointerLikeModel<UnrankedMemRefType>>(*getContext());
- LLVM::LLVMPointerType::attachInterface<LLVMPointerPointerLikeModel>(
- *getContext());
PrivateType::attachInterface<PrivateTypePointerLikeModel>(*getContext());
-
- // Attach operation interfaces
- memref::GetGlobalOp::attachInterface<MemrefAddressOfGlobalModel>(
- *getContext());
- LLVM::AddressOfOp::attachInterface<LLVMAddressOfGlobalModel>(*getContext());
- memref::GlobalOp::attachInterface<MemrefGlobalVariableModel>(*getContext());
- gpu::LaunchOp::attachInterface<GPULaunchOffloadRegionModel>(*getContext());
}
//===----------------------------------------------------------------------===//
@@ -1926,7 +1533,6 @@ createDestroyRegion(OpBuilder &builder, Location loc, Region &destroyRegion,
acc::TerminatorOp::create(builder, loc);
return success();
}
-
} // namespace
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACCCG.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACCCG.cpp
index c88d7b0e67b1c0..dd2aae7e121ebe 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACCCG.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACCCG.cpp
@@ -15,7 +15,7 @@
//===----------------------------------------------------------------------===//
#include "mlir/Dialect/Arith/IR/Arith.h"
-#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUBase.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/Utils/StaticValueUtils.h"
#include "mlir/IR/BuiltinAttributes.h"
@@ -33,7 +33,6 @@ using namespace mlir;
using namespace acc;
namespace {
-
/// Generic helper for single-region OpenACC ops that execute their body once
/// and then continue after the operation with their results (if any).
static void
@@ -371,7 +370,6 @@ static void printGPUParallelDimList(AsmPrinter &printer,
});
printer << "]";
}
-
} // namespace
//===----------------------------------------------------------------------===//
@@ -976,8 +974,7 @@ LogicalResult GPUSharedMemoryOp::verify() {
auto resultTy = cast<MemRefType>(getResult().getType());
auto addrSpace =
dyn_cast_if_present<gpu::AddressSpaceAttr>(resultTy.getMemorySpace());
- if (!addrSpace ||
- addrSpace.getValue() != gpu::GPUDialect::getWorkgroupAddressSpace())
+ if (!addrSpace || addrSpace.getValue() != gpu::AddressSpace::Workgroup)
return emitOpError("result memref must use #gpu.address_space<workgroup>");
return success();
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCBindRoutine.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCBindRoutine.cpp
index 7aace49e4e7ac1..4c6ed1099743c3 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCBindRoutine.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCBindRoutine.cpp
@@ -28,6 +28,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
@@ -52,7 +53,6 @@ using namespace mlir;
using namespace mlir::acc;
namespace {
-
static RoutineOp getFirstAccRoutineOp(FunctionOpInterface funcOp,
const SymbolTable &fallbackSymTab) {
SymbolRefAttr routineRef;
@@ -176,6 +176,12 @@ class ACCBindRoutine : public acc::impl::ACCBindRoutineBase<ACCBindRoutine> {
public:
using acc::impl::ACCBindRoutineBase<ACCBindRoutine>::ACCBindRoutineBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ acc::impl::ACCBindRoutineBase<ACCBindRoutine>::getDependentDialects(
+ registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
FunctionOpInterface func = getOperation();
ModuleOp module = func->getParentOfType<ModuleOp>();
@@ -254,5 +260,4 @@ class ACCBindRoutine : public acc::impl::ACCBindRoutineBase<ACCBindRoutine> {
signalPassFailure();
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index f8560857bb47b7..fd99bd71795d6c 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -87,6 +87,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
@@ -4380,6 +4381,11 @@ class ACCCGToGPU : public acc::impl::ACCCGToGPUBase<ACCCGToGPU> {
public:
using acc::impl::ACCCGToGPUBase<ACCCGToGPU>::ACCCGToGPUBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ acc::impl::ACCCGToGPUBase<ACCCGToGPU>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
FunctionOpInterface funcOp = getOperation();
MLIRContext *context = funcOp->getContext();
@@ -4427,5 +4433,4 @@ class ACCCGToGPU : public acc::impl::ACCCGToGPUBase<ACCCGToGPU> {
}
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCComputeLowering.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCComputeLowering.cpp
index f22d6f7b64afa9..9be9d2b97c98f9 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCComputeLowering.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCComputeLowering.cpp
@@ -43,6 +43,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
@@ -75,7 +76,6 @@ using namespace mlir;
using namespace mlir::acc;
namespace {
-
//===----------------------------------------------------------------------===//
// Helper functions
//===----------------------------------------------------------------------===//
@@ -485,6 +485,11 @@ class ACCComputeLowering
public:
using ACCComputeLoweringBase::ACCComputeLoweringBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCComputeLoweringBase::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
auto op = getOperation();
auto *context = op.getContext();
@@ -515,5 +520,4 @@ class ACCComputeLowering
return signalPassFailure();
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
index 1d4441a1476d0b..a7aa45697cd0f9 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCDeclareGPUModuleInsertion.cpp
@@ -49,6 +49,7 @@
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/IR/BuiltinOps.h"
@@ -68,7 +69,6 @@ namespace acc {
using namespace mlir;
namespace {
-
static bool hasAccDeclareGlobals(ModuleOp mod) {
for (Operation &op : mod.getBody()->getOperations())
if (op.getDiscardableAttr(acc::getDeclareAttrName()))
@@ -94,6 +94,12 @@ class ACCDeclareGPUModuleInsertion
using acc::impl::ACCDeclareGPUModuleInsertionBase<
ACCDeclareGPUModuleInsertion>::ACCDeclareGPUModuleInsertionBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ acc::impl::ACCDeclareGPUModuleInsertionBase<
+ ACCDeclareGPUModuleInsertion>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
LogicalResult copyGlobalsToGPUModule(gpu::GPUModuleOp gpuMod, ModuleOp mod,
acc::OpenACCSupport &accSupport) const {
SymbolTable gpuSymTable(gpuMod);
@@ -188,5 +194,4 @@ class ACCDeclareGPUModuleInsertion
return;
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksData.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksData.cpp
index ab251c75513552..1abb606353d501 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksData.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksData.cpp
@@ -13,6 +13,7 @@
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "llvm/ADT/STLExtras.h"
@@ -30,7 +31,6 @@ namespace acc {
using namespace mlir;
namespace {
-
static bool isStructuredDeclareEnter(acc::DeclareEnterOp op) {
return !op.getToken().getUsers().empty();
}
@@ -229,6 +229,11 @@ class ACCEmitRemarksData
public:
using ACCEmitRemarksDataBase<ACCEmitRemarksData>::ACCEmitRemarksDataBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCEmitRemarksDataBase<ACCEmitRemarksData>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
func::FuncOp func = getOperation();
@@ -268,5 +273,4 @@ class ACCEmitRemarksData
});
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksLoop.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksLoop.cpp
index 069bb00b06445c..07f8162df969a5 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksLoop.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksLoop.cpp
@@ -13,6 +13,7 @@
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenACC/OpenACCParMapping.h"
#include "mlir/Dialect/OpenACC/OpenACCUtilsLoop.h"
@@ -35,7 +36,6 @@ namespace acc {
using namespace mlir;
namespace {
-
static bool
computeRegionInSpecializedAccRoutine(acc::ComputeRegionOp computeRegion) {
if (auto func = computeRegion->getParentOfType<FunctionOpInterface>())
@@ -158,6 +158,11 @@ class ACCEmitRemarksLoop
public:
using ACCEmitRemarksLoopBase<ACCEmitRemarksLoop>::ACCEmitRemarksLoopBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCEmitRemarksLoopBase<ACCEmitRemarksLoop>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
func::FuncOp func = getOperation();
acc::OpenACCSupport &accSupport = getAnalysis<acc::OpenACCSupport>();
@@ -179,5 +184,4 @@ class ACCEmitRemarksLoop
});
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksPrivate.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksPrivate.cpp
index 645080aea0b79f..bd3d4229e1889a 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksPrivate.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCEmitRemarksPrivate.cpp
@@ -13,6 +13,7 @@
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "llvm/ADT/StringExtras.h"
@@ -31,7 +32,6 @@ namespace acc {
using namespace mlir;
namespace {
-
template <typename OpTy>
static void reportPrivatization(Operation *accOp, ValueRange operands,
acc::OpenACCSupport &accSupport,
@@ -80,6 +80,12 @@ class ACCEmitRemarksPrivate
using ACCEmitRemarksPrivateBase<
ACCEmitRemarksPrivate>::ACCEmitRemarksPrivateBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCEmitRemarksPrivateBase<ACCEmitRemarksPrivate>::getDependentDialects(
+ registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
func::FuncOp func = getOperation();
@@ -96,5 +102,4 @@ class ACCEmitRemarksPrivate
});
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCIfClauseLowering.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCIfClauseLowering.cpp
index 8adf441f7bd443..30ad8171e4139d 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCIfClauseLowering.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCIfClauseLowering.cpp
@@ -54,6 +54,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
@@ -82,11 +83,16 @@ using namespace mlir;
using namespace mlir::acc;
namespace {
-
class ACCIfClauseLowering
: public acc::impl::ACCIfClauseLoweringBase<ACCIfClauseLowering> {
using ACCIfClauseLoweringBase<ACCIfClauseLowering>::ACCIfClauseLoweringBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCIfClauseLoweringBase<ACCIfClauseLowering>::getDependentDialects(
+ registry);
+ acc::registerAllExtensions(registry);
+ }
+
private:
OpenACCSupport *accSupport = nullptr;
@@ -325,5 +331,4 @@ void ACCIfClauseLowering::runOnOperation() {
}
} while (erased);
}
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitData.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitData.cpp
index 99970b49ebca64..bc5a36ff4f3355 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitData.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitData.cpp
@@ -197,6 +197,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "llvm/ADT/SmallVectorExtras.h"
@@ -232,11 +233,16 @@ namespace acc {
using namespace mlir;
namespace {
-
class ACCImplicitData : public acc::impl::ACCImplicitDataBase<ACCImplicitData> {
public:
using acc::impl::ACCImplicitDataBase<ACCImplicitData>::ACCImplicitDataBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ acc::impl::ACCImplicitDataBase<ACCImplicitData>::getDependentDialects(
+ registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override;
private:
@@ -802,5 +808,4 @@ void ACCImplicitData::runOnOperation() {
}
});
}
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitDeclare.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitDeclare.cpp
index 88ff5da9903774..748988b504d014 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitDeclare.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitDeclare.cpp
@@ -175,6 +175,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
@@ -201,7 +202,6 @@ namespace acc {
using namespace mlir;
namespace {
-
using GlobalOpSetT = llvm::SmallSetVector<Operation *, 16>;
/// Checks whether a use of the requested `globalOp` should be considered
@@ -383,6 +383,11 @@ class ACCImplicitDeclare
public:
using ACCImplicitDeclareBase<ACCImplicitDeclare>::ACCImplicitDeclareBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCImplicitDeclareBase<ACCImplicitDeclare>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
ModuleOp mod = getOperation();
MLIRContext *context = &getContext();
@@ -482,5 +487,4 @@ class ACCImplicitDeclare
}
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitRoutine.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitRoutine.cpp
index c5a4e35616f86d..547a661e77850a 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitRoutine.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCImplicitRoutine.cpp
@@ -49,6 +49,7 @@
// should pre-register the `acc::OpenACCSupport` analysis.
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
@@ -73,7 +74,6 @@ namespace acc {
} // namespace mlir
namespace {
-
using namespace mlir;
class ACCImplicitRoutine
@@ -246,6 +246,11 @@ class ACCImplicitRoutine
public:
using ACCImplicitRoutineBase<ACCImplicitRoutine>::ACCImplicitRoutineBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCImplicitRoutineBase<ACCImplicitRoutine>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
auto module = getOperation();
mlir::OpBuilder builder(module.getContext());
@@ -278,5 +283,4 @@ class ACCImplicitRoutine
return signalPassFailure();
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCLegalizeSerial.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCLegalizeSerial.cpp
index f41ce276f994f0..1005b63affd4ce 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCLegalizeSerial.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCLegalizeSerial.cpp
@@ -17,6 +17,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
@@ -105,6 +106,11 @@ class ACCLegalizeSerial
: public mlir::acc::impl::ACCLegalizeSerialBase<ACCLegalizeSerial> {
public:
using ACCLegalizeSerialBase<ACCLegalizeSerial>::ACCLegalizeSerialBase;
+
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCLegalizeSerialBase<ACCLegalizeSerial>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
void runOnOperation() override {
func::FuncOp funcOp = getOperation();
MLIRContext *context = funcOp.getContext();
@@ -113,5 +119,4 @@ class ACCLegalizeSerial
(void)applyPatternsGreedily(funcOp, std::move(patterns));
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCLoopTiling.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCLoopTiling.cpp
index 766d50fe543482..9f71b5d77cd558 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCLoopTiling.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCLoopTiling.cpp
@@ -60,6 +60,7 @@
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenACC/OpenACCUtilsTiling.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
@@ -212,6 +213,11 @@ class ACCLoopTiling : public acc::impl::ACCLoopTilingBase<ACCLoopTiling> {
public:
using ACCLoopTilingBase<ACCLoopTiling>::ACCLoopTilingBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCLoopTilingBase<ACCLoopTiling>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
func::FuncOp funcOp = getOperation();
MLIRContext *context = funcOp.getContext();
@@ -225,5 +231,4 @@ class ACCLoopTiling : public acc::impl::ACCLoopTilingBase<ACCLoopTiling> {
(void)applyPatternsGreedily(funcOp, std::move(patterns), grc);
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCRecipeMaterialization.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCRecipeMaterialization.cpp
index 14422f079f17c7..32d2f672b5d461 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCRecipeMaterialization.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCRecipeMaterialization.cpp
@@ -48,6 +48,7 @@
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Arith/Utils/Utils.h"
#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenACC/OpenACCParMapping.h"
#include "mlir/Dialect/OpenACC/OpenACCUtils.h"
@@ -79,7 +80,6 @@ namespace acc {
#define DEBUG_TYPE "acc-recipe-materialization"
namespace {
-
using namespace mlir;
static void setLocation(Region ®ion, Location loc) {
@@ -173,6 +173,12 @@ class ACCRecipeMaterialization
public:
using acc::impl::ACCRecipeMaterializationBase<
ACCRecipeMaterialization>::ACCRecipeMaterializationBase;
+
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ acc::impl::ACCRecipeMaterializationBase<
+ ACCRecipeMaterialization>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
void runOnOperation() override;
private:
@@ -643,5 +649,4 @@ void ACCRecipeMaterialization::runOnOperation() {
removeRecipe(recipe, moduleOp, usedSymbols);
});
}
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCRoutineLowering.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCRoutineLowering.cpp
index 4e3d7e7e77cbe6..8fb9003ac62b0b 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCRoutineLowering.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCRoutineLowering.cpp
@@ -33,6 +33,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
@@ -58,7 +59,6 @@ using namespace mlir;
using namespace mlir::acc;
namespace {
-
/// Compute the ParLevel from an acc.routine op for specialization.
static ParLevel computeParLevel(RoutineOp routineOp, DeviceType deviceType) {
auto gangDim = routineOp.getGangDimValue(deviceType);
@@ -184,6 +184,11 @@ class ACCRoutineLowering
public:
using ACCRoutineLoweringBase::ACCRoutineLoweringBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCRoutineLoweringBase::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
ModuleOp mod = getOperation();
if (mod.getOps<RoutineOp>().empty()) {
@@ -232,5 +237,4 @@ class ACCRoutineLowering
finalizeRoutines(accRoutineInfo, ctx);
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCRoutineToGPUFunc.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCRoutineToGPUFunc.cpp
index 6d9528ce57549f..d771eff93ac8e3 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCRoutineToGPUFunc.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCRoutineToGPUFunc.cpp
@@ -61,6 +61,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
@@ -89,7 +90,6 @@ using namespace mlir;
using namespace mlir::acc;
namespace {
-
/// Create a gpu.func from a func.func by cloning the body.
static gpu::GPUFuncOp createGPUFuncFromFunc(OpBuilder &builder,
func::FuncOp sourceFunc) {
@@ -328,6 +328,12 @@ class ACCRoutineToGPUFunc
using acc::impl::ACCRoutineToGPUFuncBase<
ACCRoutineToGPUFunc>::ACCRoutineToGPUFuncBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ acc::impl::ACCRoutineToGPUFuncBase<
+ ACCRoutineToGPUFunc>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
ModuleOp mod = getOperation();
if (mod.getOps<RoutineOp>().empty()) {
@@ -368,5 +374,4 @@ class ACCRoutineToGPUFunc
cleanupHostModule(funcsToClone);
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCSpecializeForDevice.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCSpecializeForDevice.cpp
index 9f114a2941fb75..526415189607e4 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCSpecializeForDevice.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCSpecializeForDevice.cpp
@@ -58,6 +58,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
@@ -79,7 +80,6 @@ using namespace mlir;
using namespace mlir::acc;
namespace {
-
/// Fold acc.on_device with a constant device type in device code.
class FoldAccOnDeviceOpConversion : public OpRewritePattern<acc::OnDeviceOp> {
using OpRewritePattern<acc::OnDeviceOp>::OpRewritePattern;
@@ -111,6 +111,12 @@ class ACCSpecializeForDevice
using ACCSpecializeForDeviceBase<
ACCSpecializeForDevice>::ACCSpecializeForDeviceBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCSpecializeForDeviceBase<ACCSpecializeForDevice>::getDependentDialects(
+ registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
func::FuncOp func = getOperation();
@@ -151,7 +157,6 @@ class ACCSpecializeForDevice
}
}
};
-
} // namespace
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCSpecializeForHost.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCSpecializeForHost.cpp
index 9019827ac3964f..96322e82cf6499 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCSpecializeForHost.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCSpecializeForHost.cpp
@@ -76,6 +76,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
@@ -118,7 +119,6 @@ static bool isAtomicConversionDeferred(Operation *op) {
}
namespace {
-
// Lower orphan acc.atomic.update by: load from addr, clone region expr with
// the loaded value, then store the computed result back to addr.
// Only matches outside compute regions and atomic captures.
@@ -336,6 +336,12 @@ class ACCSpecializeForHost
using ACCSpecializeForHostBase<
ACCSpecializeForHost>::ACCSpecializeForHostBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ ACCSpecializeForHostBase<ACCSpecializeForHost>::getDependentDialects(
+ registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
LLVM_DEBUG(llvm::dbgs() << "Enter ACCSpecializeForHost()\n");
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/CMakeLists.txt b/mlir/lib/Dialect/OpenACC/Transforms/CMakeLists.txt
index 3fd2672e82c502..cb68dbbc73dc62 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/CMakeLists.txt
+++ b/mlir/lib/Dialect/OpenACC/Transforms/CMakeLists.txt
@@ -41,6 +41,7 @@ add_mlir_dialect_library(MLIROpenACCTransforms
MLIRComplexDialect
MLIROpenACCAnalysis
MLIROpenACCDialect
+ MLIROpenACCAllExtensions
MLIROpenACCUtils
MLIRFuncDialect
MLIRGPUDialect
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/LegalizeDataValues.cpp b/mlir/lib/Dialect/OpenACC/Transforms/LegalizeDataValues.cpp
index 8b15d2c1cc7f2e..7d012e4eb84272 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/LegalizeDataValues.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/LegalizeDataValues.cpp
@@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
@@ -25,7 +26,6 @@ namespace acc {
using namespace mlir;
namespace {
-
static bool insideAccComputeRegion(mlir::Operation *op) {
return op->getParentOfType<ACC_COMPUTE_CONSTRUCT_OPS>();
}
@@ -152,6 +152,12 @@ class LegalizeDataValuesInRegion
using LegalizeDataValuesInRegionBase<
LegalizeDataValuesInRegion>::LegalizeDataValuesInRegionBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ LegalizeDataValuesInRegionBase<
+ LegalizeDataValuesInRegion>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
void runOnOperation() override {
func::FuncOp funcOp = getOperation();
bool replaceHostVsDevice = this->hostToDevice.getValue();
@@ -196,5 +202,4 @@ class LegalizeDataValuesInRegion
});
}
};
-
} // end anonymous namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/OffloadLiveInValueCanonicalization.cpp b/mlir/lib/Dialect/OpenACC/Transforms/OffloadLiveInValueCanonicalization.cpp
index 4b57354b8b5d8a..3a08706c835087 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/OffloadLiveInValueCanonicalization.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/OffloadLiveInValueCanonicalization.cpp
@@ -66,6 +66,7 @@
//
//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/Analysis/TopologicalSortUtils.h"
@@ -96,7 +97,6 @@ namespace acc {
using namespace mlir;
namespace {
-
/// Returns true if all users of the given value are inside the region.
static bool allUsersAreInsideRegion(Value val, Region ®ion) {
for (Operation *user : val.getUsers())
@@ -206,6 +206,12 @@ class OffloadLiveInValueCanonicalization
OffloadLiveInValueCanonicalization>::
OffloadLiveInValueCanonicalizationBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ acc::impl::OffloadLiveInValueCanonicalizationBase<
+ OffloadLiveInValueCanonicalization>::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
/// Canonicalizes live-in values for a region by sinking or rematerializing
/// operations. Returns true if any changes were made.
bool canonicalizeLiveInValues(Region ®ion,
@@ -326,5 +332,4 @@ class OffloadLiveInValueCanonicalization
LLVM_DEBUG(llvm::dbgs() << "Exit OffloadLiveInValueCanonicalization\n");
}
};
-
} // namespace
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/OffloadTargetVerifier.cpp b/mlir/lib/Dialect/OpenACC/Transforms/OffloadTargetVerifier.cpp
index 62941686890ada..1465ea20ef9b0e 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/OffloadTargetVerifier.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/OffloadTargetVerifier.cpp
@@ -54,6 +54,7 @@
#include "mlir/Analysis/Liveness.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
#include "mlir/IR/SymbolTable.h"
@@ -71,12 +72,16 @@ namespace acc {
using namespace mlir;
namespace {
-
class OffloadTargetVerifier
: public acc::impl::OffloadTargetVerifierBase<OffloadTargetVerifier> {
public:
using OffloadTargetVerifierBase::OffloadTargetVerifierBase;
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ OffloadTargetVerifierBase::getDependentDialects(registry);
+ acc::registerAllExtensions(registry);
+ }
+
/// Returns true if the target device type corresponds to host execution.
bool isHostTarget() const {
return deviceType == acc::DeviceType::Host ||
@@ -255,5 +260,4 @@ class OffloadTargetVerifier
LLVM_DEBUG(llvm::dbgs() << "Exit OffloadTargetVerifier()\n");
}
};
-
} // namespace
diff --git a/mlir/lib/RegisterAllExtensions.cpp b/mlir/lib/RegisterAllExtensions.cpp
index 67df66c0136096..ff7810c8c6430a 100644
--- a/mlir/lib/RegisterAllExtensions.cpp
+++ b/mlir/lib/RegisterAllExtensions.cpp
@@ -46,6 +46,7 @@
#include "mlir/Dialect/Math/TransformOps/MathTransformOps.h"
#include "mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.h"
#include "mlir/Dialect/NVGPU/TransformOps/NVGPUTransformOps.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/SCF/TransformOps/SCFTransformOps.h"
#include "mlir/Dialect/SparseTensor/TransformOps/SparseTensorTransformOps.h"
#include "mlir/Dialect/Tensor/Extensions/AllExtensions.h"
@@ -72,6 +73,7 @@
/// individually register the specific extensions that are useful for the
/// pipelines and transformations you are using.
void mlir::registerAllExtensions(DialectRegistry ®istry) {
+ acc::registerAllExtensions(registry);
// Register all conversions to LLVM extensions.
registerConvertArithToEmitCInterface(registry);
arith::registerConvertArithToLLVMInterface(registry);
diff --git a/mlir/test/CAPI/CMakeLists.txt b/mlir/test/CAPI/CMakeLists.txt
index bd8b58eb442bf4..66849266d57558 100644
--- a/mlir/test/CAPI/CMakeLists.txt
+++ b/mlir/test/CAPI/CMakeLists.txt
@@ -64,6 +64,14 @@ _add_capi_test_executable(mlir-capi-llvm-test
MLIRCAPIRegisterEverything
)
+_add_capi_test_executable(mlir-capi-openacc-test
+ openacc.c
+ LINK_LIBS PRIVATE
+ MLIRCAPIIR
+ MLIRCAPIFunc
+ MLIRCAPIOpenACC
+)
+
_add_capi_test_executable(mlir-capi-remarks-test
remarks.c
LINK_LIBS PRIVATE
diff --git a/mlir/test/CAPI/openacc.c b/mlir/test/CAPI/openacc.c
new file mode 100644
index 00000000000000..b24a3bf66c7ef3
--- /dev/null
+++ b/mlir/test/CAPI/openacc.c
@@ -0,0 +1,55 @@
+//===- openacc.c - OpenACC dialect registration tests
+//----------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// RUN: mlir-capi-openacc-test | FileCheck %s
+
+#include "mlir-c/Dialect/OpenACC.h"
+#include "mlir-c/Dialect/Func.h"
+#include "mlir-c/IR.h"
+#include <stdio.h>
+
+int main(void) {
+ for (int useRegistry = 0; useRegistry != 2; ++useRegistry) {
+ MlirContext context = mlirContextCreate();
+ MlirDialectHandle acc = mlirGetDialectHandle__acc__();
+ mlirDialectHandleRegisterDialect(mlirGetDialectHandle__func__(), context);
+ if (useRegistry)
+ mlirDialectHandleRegisterDialect(acc, context);
+ else
+ mlirDialectHandleLoadDialect(acc, context);
+
+ // Verifying the copyin requires the standard MemRef pointer model even
+ // though the input has no MemRef operations.
+ MlirModule module = mlirModuleCreateParse(
+ context,
+ mlirStringRefCreateFromCString(
+ "module { func.func @f(%arg : memref<f32>) { "
+ "%0 = acc.copyin varPtr(%arg : memref<f32>) -> memref<f32> "
+ "return } }"));
+ if (mlirModuleIsNull(module) ||
+ !mlirOperationVerify(mlirModuleGetOperation(module)))
+ return 1;
+
+ MlirAttribute space = mlirAttributeParseGet(
+ context,
+ mlirStringRefCreateFromCString("#gpu.address_space<workgroup>"));
+ MlirAttribute dim = mlirAttributeParseGet(
+ context, mlirStringRefCreateFromCString("#acc.par_dim<thread_y>"));
+ MlirType token = mlirTypeParseGet(
+ context, mlirStringRefCreateFromCString("!gpu.async.token"));
+ if (mlirAttributeIsNull(space) || mlirAttributeIsNull(dim) ||
+ mlirTypeIsNull(token))
+ return 2;
+ mlirModuleDestroy(module);
+ mlirContextDestroy(context);
+ }
+ // CHECK: OpenACC C API registration passed
+ puts("OpenACC C API registration passed");
+ return 0;
+}
diff --git a/mlir/test/CMakeLists.txt b/mlir/test/CMakeLists.txt
index 747d0cf68116b4..0329b46b52a476 100644
--- a/mlir/test/CMakeLists.txt
+++ b/mlir/test/CMakeLists.txt
@@ -124,6 +124,7 @@ set(MLIR_TEST_DEPENDS
mlir-capi-ir-test
mlir-capi-irdl-test
mlir-capi-llvm-test
+ mlir-capi-openacc-test
mlir-capi-pass-test
mlir-capi-quant-test
mlir-capi-remarks-test
diff --git a/mlir/unittests/Dialect/OpenACC/CMakeLists.txt b/mlir/unittests/Dialect/OpenACC/CMakeLists.txt
index 0fe641c2da5566..8dd25bb56d4196 100644
--- a/mlir/unittests/Dialect/OpenACC/CMakeLists.txt
+++ b/mlir/unittests/Dialect/OpenACC/CMakeLists.txt
@@ -1,5 +1,8 @@
+add_subdirectory(Standalone)
+
add_mlir_unittest(MLIROpenACCTests
OpenACCCGOpsTest.cpp
+ OpenACCExtensionsTest.cpp
OpenACCOpsTest.cpp
OpenACCOpsInterfacesTest.cpp
OpenACCRuntimeUtilsTest.cpp
@@ -26,7 +29,9 @@ mlir_target_link_libraries(MLIROpenACCTests
MLIRArithValueBoundsOpInterfaceImpl
MLIROpenACCAnalysis
MLIROpenACCDialect
+ MLIROpenACCAllExtensions
MLIROpenACCUtils
+ MLIROpenACCTransforms
MLIRSCFDialect
MLIRControlFlowDialect
)
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCCGOpsTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCCGOpsTest.cpp
index 44283c9091b4d8..cec30db0c389b9 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCCGOpsTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCCGOpsTest.cpp
@@ -9,6 +9,7 @@
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenACC/OpenACCUtilsCG.h"
#include "mlir/IR/BuiltinOps.h"
@@ -30,6 +31,9 @@ using namespace mlir::acc;
class OpenACCCGOpsTest : public ::testing::Test {
protected:
OpenACCCGOpsTest() : b(&context), loc(UnknownLoc::get(&context)) {
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context.loadDialect<OpenACCDialect, arith::ArithDialect, func::FuncDialect,
gpu::GPUDialect>();
}
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCExtensionsTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCExtensionsTest.cpp
new file mode 100644
index 00000000000000..caa98eb8051bb2
--- /dev/null
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCExtensionsTest.cpp
@@ -0,0 +1,96 @@
+//===- OpenACCExtensionsTest.cpp - Extension tests ----------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
+#include "mlir/Dialect/OpenACC/Extensions/MemRefInterfaces.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/Dialect/OpenACC/Transforms/Passes.h"
+#include "mlir/IR/DialectRegistry.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/Pass/Pass.h"
+#include "gtest/gtest.h"
+
+using namespace mlir;
+
+TEST(OpenACCExtensions, PassRegistersStandardModels) {
+ DialectRegistry registry;
+ acc::createOffloadTargetVerifier()->getDependentDialects(registry);
+ MLIRContext context(registry);
+ context.loadDialect<acc::OpenACCDialect, gpu::GPUDialect>();
+ EXPECT_TRUE(OperationName("gpu.launch", &context)
+ .hasInterface<acc::OffloadRegionOpInterface>());
+ EXPECT_TRUE(isa<acc::PointerLikeType>(
+ MemRefType::get({4}, IntegerType::get(&context, 32))));
+}
+
+TEST(OpenACCExtensions, MemRefTypesWithoutMemRefOperations) {
+ DialectRegistry registry;
+ acc::registerMemRefInterfaces(registry);
+ MLIRContext context(registry);
+ context.loadDialect<acc::OpenACCDialect>();
+ auto type = MemRefType::get({4}, IntegerType::get(&context, 32));
+ EXPECT_TRUE(isa<acc::PointerLikeType>(type));
+ EXPECT_EQ(context.getLoadedDialect("gpu"), nullptr);
+}
+
+TEST(OpenACCExtensions, EitherLoadingOrder) {
+ for (bool accFirst : {false, true}) {
+ DialectRegistry registry;
+ acc::registerAllExtensions(registry);
+ MLIRContext context(registry);
+ EXPECT_EQ(context.getLoadedDialect("gpu"), nullptr);
+ if (accFirst) {
+ context.loadDialect<acc::OpenACCDialect>();
+ EXPECT_EQ(context.getLoadedDialect("gpu"), nullptr);
+ context.loadDialect<gpu::GPUDialect, LLVM::LLVMDialect,
+ memref::MemRefDialect>();
+ } else {
+ context.loadDialect<gpu::GPUDialect, LLVM::LLVMDialect,
+ memref::MemRefDialect>();
+ context.loadDialect<acc::OpenACCDialect>();
+ }
+
+ auto space =
+ gpu::AddressSpaceAttr::get(&context, gpu::AddressSpace::Workgroup);
+ EXPECT_EQ(space.getDialect().getNamespace(), "gpu");
+ EXPECT_EQ(gpu::AsyncTokenType::get(&context).getDialect().getNamespace(),
+ "gpu");
+ auto type = MemRefType::get({4}, IntegerType::get(&context, 32),
+ MemRefLayoutAttrInterface{}, space);
+ auto pointer = dyn_cast<acc::PointerLikeType>(type);
+ ASSERT_TRUE(pointer);
+ EXPECT_TRUE(pointer.isDeviceAccessible({}));
+ EXPECT_TRUE(
+ isa<acc::PointerLikeType>(LLVM::LLVMPointerType::get(&context)));
+
+ OperationName launch("gpu.launch", &context);
+ EXPECT_TRUE(launch.hasInterface<acc::OffloadRegionOpInterface>());
+ EXPECT_TRUE(OperationName("memref.global", &context)
+ .hasInterface<acc::GlobalVariableOpInterface>());
+ EXPECT_TRUE(OperationName("llvm.mlir.addressof", &context)
+ .hasInterface<acc::AddressOfGlobalOpInterface>());
+ }
+}
+
+TEST(OpenACCExtensions, RegistrationAfterLoading) {
+ MLIRContext context;
+ context.loadDialect<acc::OpenACCDialect, gpu::GPUDialect,
+ memref::MemRefDialect, LLVM::LLVMDialect>();
+ EXPECT_FALSE(OperationName("gpu.launch", &context)
+ .hasInterface<acc::OffloadRegionOpInterface>());
+ DialectRegistry registry;
+ acc::registerAllExtensions(registry);
+ context.appendDialectRegistry(registry);
+ EXPECT_TRUE(OperationName("gpu.launch", &context)
+ .hasInterface<acc::OffloadRegionOpInterface>());
+ EXPECT_TRUE(isa<acc::PointerLikeType>(
+ MemRefType::get({4}, IntegerType::get(&context, 32))));
+}
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCOpsInterfacesTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCOpsInterfacesTest.cpp
index 26ca426dd1158e..dc1361d2a53212 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCOpsInterfacesTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCOpsInterfacesTest.cpp
@@ -8,6 +8,7 @@
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinTypes.h"
@@ -26,6 +27,9 @@ class OpenACCOpsInterfacesTest : public ::testing::Test {
protected:
OpenACCOpsInterfacesTest()
: context(), builder(&context), loc(UnknownLoc::get(&context)) {
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context.loadDialect<acc::OpenACCDialect, memref::MemRefDialect,
LLVM::LLVMDialect>();
}
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCOpsTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCOpsTest.cpp
index df1109d7f6dcfb..17c20e91c7e1a1 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCOpsTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCOpsTest.cpp
@@ -8,6 +8,7 @@
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/Diagnostics.h"
@@ -36,6 +37,9 @@ static OpTy createEmptyOp(OpBuilder &builder, Location loc) {
class OpenACCOpsTest : public ::testing::Test {
protected:
OpenACCOpsTest() : b(&context), loc(UnknownLoc::get(&context)) {
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context.loadDialect<acc::OpenACCDialect, arith::ArithDialect,
memref::MemRefDialect>();
}
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCTypeInterfacesTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCTypeInterfacesTest.cpp
index 29407122049308..0b5309b9cdfaf5 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCTypeInterfacesTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCTypeInterfacesTest.cpp
@@ -10,6 +10,7 @@
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/BuiltinDialect.h"
#include "mlir/IR/BuiltinOps.h"
@@ -23,7 +24,6 @@ using namespace mlir::acc;
using namespace mlir::LLVM;
namespace {
-
/// Test model that attaches ReducibleType to IntegerType for testing purposes.
/// This only implements a small subset of reduction operators to exercise the
/// interface - a real implementation would handle all valid operators.
@@ -40,7 +40,6 @@ struct TestReducibleIntegerModel
}
}
};
-
} // namespace
//===----------------------------------------------------------------------===//
@@ -56,6 +55,9 @@ class OpenACCTypeInterfacesTest : public ::testing::Test {
IntegerType::attachInterface<TestReducibleIntegerModel>(*ctx);
});
context.appendDialectRegistry(registry);
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context
.loadDialect<acc::OpenACCDialect, arith::ArithDialect,
memref::MemRefDialect, func::FuncDialect, LLVMDialect>();
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsCGTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsCGTest.cpp
index 9de3158accee1d..540cb88f42786d 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsCGTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsCGTest.cpp
@@ -14,6 +14,7 @@
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRefDialect.h"
#include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenACC/OpenACCParMapping.h"
#include "mlir/Dialect/SCF/IR/SCF.h"
@@ -38,6 +39,9 @@ class OpenACCUtilsCGTest : public ::testing::Test {
DialectRegistry registry;
arith::registerValueBoundsOpInterfaceExternalModels(registry);
context.appendDialectRegistry(registry);
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context.loadDialect<acc::OpenACCDialect, arith::ArithDialect,
func::FuncDialect, scf::SCFDialect, gpu::GPUDialect,
memref::MemRefDialect, DLTIDialect>();
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsGPUTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsGPUTest.cpp
index e4243ed04f8f10..a21b580e2cee57 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsGPUTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsGPUTest.cpp
@@ -9,6 +9,7 @@
#include "mlir/Dialect/OpenACC/OpenACCUtilsGPU.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/MLIRContext.h"
#include "mlir/IR/OwningOpRef.h"
@@ -24,6 +25,9 @@ using namespace mlir::acc;
class OpenACCUtilsGPUTest : public ::testing::Test {
protected:
OpenACCUtilsGPUTest() : b(&context), loc(UnknownLoc::get(&context)) {
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context.loadDialect<arith::ArithDialect, gpu::GPUDialect>();
}
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsLoopTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsLoopTest.cpp
index e4d94855a06ec3..fc48be27fa6872 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsLoopTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsLoopTest.cpp
@@ -12,6 +12,7 @@
#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/MemRef/IR/MemRefDialect.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/SCF/IR/SCF.h"
#include "mlir/IR/BuiltinOps.h"
@@ -43,6 +44,9 @@ static OpTy createEmptyOp(OpBuilder &builder, Location loc) {
class OpenACCUtilsLoopTest : public ::testing::Test {
protected:
OpenACCUtilsLoopTest() : b(&context), loc(UnknownLoc::get(&context)) {
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context.loadDialect<acc::OpenACCDialect, affine::AffineDialect,
arith::ArithDialect, memref::MemRefDialect,
func::FuncDialect, scf::SCFDialect,
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsReductionTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsReductionTest.cpp
index 690b174b5ba039..7f975f9bed908a 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsReductionTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsReductionTest.cpp
@@ -10,6 +10,7 @@
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Complex/IR/Complex.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/Dialect/OpenACC/OpenACCUtilsCG.h"
#include "mlir/IR/BuiltinOps.h"
@@ -28,6 +29,9 @@ using namespace mlir::acc;
class OpenACCUtilsReductionTest : public ::testing::Test {
protected:
OpenACCUtilsReductionTest() : b(&context), loc(UnknownLoc::get(&context)) {
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context.loadDialect<acc::OpenACCDialect, arith::ArithDialect,
complex::ComplexDialect, memref::MemRefDialect>();
}
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTest.cpp
index 6d6d62d74f6124..bccde63bcc92c6 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTest.cpp
@@ -11,6 +11,7 @@
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/BuiltinTypes.h"
@@ -40,6 +41,9 @@ static OpTy createEmptyOp(OpBuilder &builder, Location loc) {
class OpenACCUtilsTest : public ::testing::Test {
protected:
OpenACCUtilsTest() : b(&context), loc(UnknownLoc::get(&context)) {
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context
.loadDialect<acc::OpenACCDialect, arith::ArithDialect, gpu::GPUDialect,
memref::MemRefDialect, func::FuncDialect>();
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp
index d4eeaec63e57a2..c1a431edd9d6e1 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp
@@ -10,6 +10,7 @@
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/MemRef/IR/MemRefDialect.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/BuiltinTypes.h"
@@ -27,6 +28,9 @@ using namespace mlir::acc;
class OpenACCUtilsTilingTest : public ::testing::Test {
protected:
OpenACCUtilsTilingTest() : b(&context), loc(UnknownLoc::get(&context)) {
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context.loadDialect<acc::OpenACCDialect, arith::ArithDialect,
memref::MemRefDialect, func::FuncDialect>();
}
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTypeTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTypeTest.cpp
index 586bda6416f90b..5abbcea1feed6f 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTypeTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTypeTest.cpp
@@ -11,6 +11,7 @@
#include "mlir/Dialect/Func/IR/FuncDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/OpenACC/Extensions/AllExtensions.h"
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/BuiltinTypes.h"
@@ -26,6 +27,9 @@ using namespace mlir::acc;
class OpenACCUtilsTypeTest : public ::testing::Test {
protected:
OpenACCUtilsTypeTest() : b(&context), loc(UnknownLoc::get(&context)) {
+ DialectRegistry extensions;
+ acc::registerAllExtensions(extensions);
+ context.appendDialectRegistry(extensions);
context.loadDialect<acc::OpenACCDialect, DLTIDialect, func::FuncDialect,
memref::MemRefDialect, LLVM::LLVMDialect>();
}
diff --git a/mlir/unittests/Dialect/OpenACC/Standalone/CMakeLists.txt b/mlir/unittests/Dialect/OpenACC/Standalone/CMakeLists.txt
new file mode 100644
index 00000000000000..e14776d5929f77
--- /dev/null
+++ b/mlir/unittests/Dialect/OpenACC/Standalone/CMakeLists.txt
@@ -0,0 +1,7 @@
+add_mlir_unittest(MLIROpenACCStandaloneTests
+ OpenACCStandaloneTest.cpp
+)
+mlir_target_link_libraries(MLIROpenACCStandaloneTests
+ PRIVATE
+ MLIROpenACCDialect
+)
diff --git a/mlir/unittests/Dialect/OpenACC/Standalone/OpenACCStandaloneTest.cpp b/mlir/unittests/Dialect/OpenACC/Standalone/OpenACCStandaloneTest.cpp
new file mode 100644
index 00000000000000..94ec8e8210bd99
--- /dev/null
+++ b/mlir/unittests/Dialect/OpenACC/Standalone/OpenACCStandaloneTest.cpp
@@ -0,0 +1,36 @@
+//===- OpenACCStandaloneTest.cpp - Standalone tests ---------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/MLIRContext.h"
+#include "llvm/Support/raw_ostream.h"
+#include "gtest/gtest.h"
+
+using namespace mlir;
+
+// This executable deliberately links MLIROpenACCDialect without the GPU
+// operation implementation or the standard OpenACC external models.
+TEST(OpenACCStandalone, LoadingDoesNotLoadOtherDialects) {
+ MLIRContext context;
+ context.loadDialect<acc::OpenACCDialect>();
+ EXPECT_EQ(context.getLoadedDialect("gpu"), nullptr);
+ EXPECT_EQ(context.getLoadedDialect("llvm"), nullptr);
+ EXPECT_EQ(context.getLoadedDialect("memref"), nullptr);
+ EXPECT_TRUE(OperationName("acc.parallel", &context).isRegistered());
+ EXPECT_TRUE(OperationName("acc.compute_region", &context).isRegistered());
+ EXPECT_TRUE(OperationName("acc.gpu_shared_memory", &context).isRegistered());
+ EXPECT_FALSE(OperationName("gpu.launch", &context).isRegistered());
+
+ auto dim = acc::GPUParallelDimAttr::get(&context, gpu::Processor::ThreadY);
+ EXPECT_EQ(dim.getProcessor(), gpu::Processor::ThreadY);
+ std::string text;
+ llvm::raw_string_ostream stream(text);
+ stream << dim;
+ EXPECT_EQ(text, "#acc.par_dim<thread_y>");
+ EXPECT_EQ(context.getLoadedDialect("gpu"), nullptr);
+}
More information about the Mlir-commits
mailing list