[Mlir-commits] [mlir] [MLIR][LLVM] Remove typed pointers from the LLVM dialect (PR #71285)
Christian Ulmann
llvmlistbot at llvm.org
Mon Nov 6 01:01:40 PST 2023
https://github.com/Dinistro updated https://github.com/llvm/llvm-project/pull/71285
>From de69c84f48a68fd17ba062f595869e63b86fb58c Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christianulmann at gmail.com>
Date: Sat, 4 Nov 2023 15:20:30 +0100
Subject: [PATCH 1/5] [MLIR][LLVM] Remove typed pointers from the LLVM dialect
This commit removes the support for typed pointers from the LLVM
dialect. Typed pointers have been deprecated for a while and thus this
removal was announced in a PSA:
https://discourse.llvm.org/t/psa-removal-of-typed-pointers-from-the-llvm-dialect/74502
This change includes:
- Changing the ` LLVMPointerType`
- Removing remaining usages of the builders and the now removed element
type
- Fixing assembly formats that require fully qualified pointer types
- Updating ODS pointer constraints
---
.../include/mlir/Dialect/LLVMIR/LLVMDialect.h | 3 +-
.../mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td | 28 +--
.../include/mlir/Dialect/LLVMIR/LLVMOpBase.td | 28 +--
mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td | 56 ++---
mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td | 21 +-
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 54 ++--
.../Conversion/LLVMCommon/TypeConverter.cpp | 8 +-
mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp | 233 +-----------------
mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp | 18 +-
mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp | 84 +------
.../LLVMIR/Transforms/TypeConsistency.cpp | 8 -
mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp | 41 ++-
mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp | 27 +-
mlir/test/Dialect/LLVMIR/global.mlir | 5 -
mlir/test/Dialect/LLVMIR/invalid.mlir | 64 +----
.../unittests/Dialect/LLVMIR/LLVMTypeTest.cpp | 34 ---
16 files changed, 118 insertions(+), 594 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h b/mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h
index bbed1ea5cf62204..06df4a601b7a3f7 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h
@@ -209,8 +209,7 @@ class GEPIndicesAdaptor {
/// global and use it to compute the address of the first character in the
/// string (operations inserted at the builder insertion point).
Value createGlobalString(Location loc, OpBuilder &builder, StringRef name,
- StringRef value, Linkage linkage,
- bool useOpaquePointers = true);
+ StringRef value, Linkage linkage);
/// LLVM requires some operations to be inside of a Module operation. This
/// function confirms that the Operation has the desired properties.
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
index 72c932ac07a2e1a..1123466b7a75e3f 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
@@ -469,16 +469,16 @@ def LLVM_ThreadlocalAddressOp : LLVM_OneResultIntrOp<"threadlocal.address", [],
def LLVM_CoroIdOp : LLVM_IntrOp<"coro.id", [], [], [], 1> {
let arguments = (ins I32:$align,
- LLVM_i8Ptr:$promise,
- LLVM_i8Ptr:$coroaddr,
- LLVM_i8Ptr:$fnaddrs);
+ LLVM_AnyPointer:$promise,
+ LLVM_AnyPointer:$coroaddr,
+ LLVM_AnyPointer:$fnaddrs);
let assemblyFormat = "$align `,` $promise `,` $coroaddr `,` $fnaddrs"
" attr-dict `:` functional-type(operands, results)";
}
def LLVM_CoroBeginOp : LLVM_IntrOp<"coro.begin", [], [], [], 1> {
let arguments = (ins LLVM_TokenType:$token,
- LLVM_i8Ptr:$mem);
+ LLVM_AnyPointer:$mem);
let assemblyFormat = "$token `,` $mem attr-dict `:` functional-type(operands, results)";
}
@@ -491,7 +491,7 @@ def LLVM_CoroAlignOp : LLVM_IntrOp<"coro.align", [0], [], [], 1> {
}
def LLVM_CoroSaveOp : LLVM_IntrOp<"coro.save", [], [], [], 1> {
- let arguments = (ins LLVM_i8Ptr:$handle);
+ let arguments = (ins LLVM_AnyPointer:$handle);
let assemblyFormat = "$handle attr-dict `:` functional-type(operands, results)";
}
@@ -502,7 +502,7 @@ def LLVM_CoroSuspendOp : LLVM_IntrOp<"coro.suspend", [], [], [], 1> {
}
def LLVM_CoroEndOp : LLVM_IntrOp<"coro.end", [], [], [], 1> {
- let arguments = (ins LLVM_i8Ptr:$handle,
+ let arguments = (ins LLVM_AnyPointer:$handle,
I1:$unwind,
LLVM_TokenType:$retvals);
let assemblyFormat = "$handle `,` $unwind `,` $retvals attr-dict `:` functional-type(operands, results)";
@@ -510,12 +510,12 @@ def LLVM_CoroEndOp : LLVM_IntrOp<"coro.end", [], [], [], 1> {
def LLVM_CoroFreeOp : LLVM_IntrOp<"coro.free", [], [], [], 1> {
let arguments = (ins LLVM_TokenType:$id,
- LLVM_i8Ptr:$handle);
+ LLVM_AnyPointer:$handle);
let assemblyFormat = "$id `,` $handle attr-dict `:` functional-type(operands, results)";
}
def LLVM_CoroResumeOp : LLVM_IntrOp<"coro.resume", [], [], [], 0> {
- let arguments = (ins LLVM_i8Ptr:$handle);
+ let arguments = (ins LLVM_AnyPointer:$handle);
let assemblyFormat = "$handle attr-dict `:` qualified(type($handle))";
}
@@ -591,19 +591,19 @@ def LLVM_DbgLabelOp : LLVM_IntrOp<"dbg.label", [], [], [], 0> {
//
def LLVM_VaStartOp : LLVM_ZeroResultIntrOp<"vastart">,
- Arguments<(ins LLVM_i8Ptr:$arg_list)> {
+ Arguments<(ins LLVM_AnyPointer:$arg_list)> {
let assemblyFormat = "$arg_list attr-dict `:` qualified(type($arg_list))";
let summary = "Initializes `arg_list` for subsequent variadic argument extractions.";
}
def LLVM_VaCopyOp : LLVM_ZeroResultIntrOp<"vacopy">,
- Arguments<(ins LLVM_i8Ptr:$dest_list, LLVM_i8Ptr:$src_list)> {
+ Arguments<(ins LLVM_AnyPointer:$dest_list, LLVM_AnyPointer:$src_list)> {
let assemblyFormat = "$src_list `to` $dest_list attr-dict `:` type(operands)";
let summary = "Copies the current argument position from `src_list` to `dest_list`.";
}
def LLVM_VaEndOp : LLVM_ZeroResultIntrOp<"vaend">,
- Arguments<(ins LLVM_i8Ptr:$arg_list)> {
+ Arguments<(ins LLVM_AnyPointer:$arg_list)> {
let assemblyFormat = "$arg_list attr-dict `:` qualified(type($arg_list))";
let summary = "Destroys `arg_list`, which has been initialized by `intr.vastart` or `intr.vacopy`.";
}
@@ -613,7 +613,7 @@ def LLVM_VaEndOp : LLVM_ZeroResultIntrOp<"vaend">,
//
def LLVM_EhTypeidForOp : LLVM_OneResultIntrOp<"eh.typeid.for"> {
- let arguments = (ins LLVM_i8Ptr:$type_info);
+ let arguments = (ins LLVM_AnyPointer:$type_info);
let assemblyFormat = "$type_info attr-dict `:` functional-type(operands, results)";
}
@@ -927,12 +927,12 @@ def LLVM_PtrAnnotation
: LLVM_OneResultIntrOp<"ptr.annotation", [0], [2],
[AllTypesMatch<["res", "ptr"]>,
AllTypesMatch<["annotation", "fileName", "attr"]>]> {
- let arguments = (ins LLVM_PointerTo<AnySignlessInteger>:$ptr,
+ let arguments = (ins LLVM_AnyPointer:$ptr,
LLVM_AnyPointer:$annotation,
LLVM_AnyPointer:$fileName,
I32:$line,
LLVM_AnyPointer:$attr);
- let results = (outs LLVM_PointerTo<AnySignlessInteger>:$res);
+ let results = (outs LLVM_AnyPointer:$res);
}
def LLVM_Annotation
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
index 503988abfc090a2..0166fbb647b41b5 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
@@ -55,43 +55,17 @@ def LLVM_AnyFloat : Type<
def LLVM_AnyPointer : Type<CPred<"::llvm::isa<::mlir::LLVM::LLVMPointerType>($_self)">,
"LLVM pointer type", "::mlir::LLVM::LLVMPointerType">;
-def LLVM_OpaquePointer : Type<
- And<[LLVM_AnyPointer.predicate,
- CPred<"::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).isOpaque()">]>,
- "LLVM opaque pointer", "::mlir::LLVM::LLVMPointerType">;
-
-// Type constraint accepting LLVM pointer type with an additional constraint
-// on the element type.
-class LLVM_PointerTo<Type pointee> : Type<
- And<[LLVM_AnyPointer.predicate,
- Or<[LLVM_OpaquePointer.predicate,
- SubstLeaves<
- "$_self",
- "::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getElementType()",
- pointee.predicate>]>]>,
- "LLVM pointer to " # pointee.summary, "::mlir::LLVM::LLVMPointerType">;
-
// Opaque pointer in a given address space.
class LLVM_OpaquePointerInAddressSpace<int addressSpace> : Type<
- And<[LLVM_OpaquePointer.predicate,
CPred<
"::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace() == "
- # addressSpace>]>,
+ # addressSpace>,
"Opaque LLVM pointer in address space " # addressSpace,
"::mlir::LLVM::LLVMPointerType"> {
let builderCall = "$_builder.getType<::mlir::LLVM::LLVMPointerType>("
# addressSpace # ")";
}
-// Type constraints accepting LLVM pointer type to integer of a specific width.
-class LLVM_IntPtrBase<int width, int addressSpace = 0> : Type<
- And<[LLVM_PointerTo<I<width>>.predicate,
- CPred<"::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace()"
- " == " # addressSpace>]>,
- "LLVM pointer to " # I<width>.summary>;
-
-def LLVM_i8Ptr : LLVM_IntPtrBase<8>;
-
// Type constraint accepting any LLVM structure type.
def LLVM_AnyStruct : Type<CPred<"::llvm::isa<::mlir::LLVM::LLVMStructType>($_self)">,
"LLVM structure type">;
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
index c8549f146d0297a..ec196e1bd13673d 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
@@ -184,16 +184,14 @@ def LLVM_AllocaOp : LLVM_Op<"alloca",
LLVM_MemOpPatterns {
let arguments = (ins AnyInteger:$arraySize,
OptionalAttr<I64Attr>:$alignment,
- OptionalAttr<TypeAttr>:$elem_type,
+ TypeAttr:$elem_type,
UnitAttr:$inalloca);
let results = (outs Res<LLVM_AnyPointer, "",
[MemAlloc<AutomaticAllocationScopeResource>]>:$res);
string llvmInstName = "Alloca";
string llvmBuilder = [{
auto addrSpace = $_resultType->getPointerAddressSpace();
- llvm::Type *elementType = moduleTranslation.convertType(
- $elem_type ? *$elem_type
- : ::llvm::cast<LLVMPointerType>(op.getType()).getElementType());
+ llvm::Type *elementType = moduleTranslation.convertType($elem_type);
auto *inst = builder.CreateAlloca(elementType, addrSpace, $arraySize);
}] # setAlignmentCode # [{
inst->setUsedWithInAlloca($inalloca);
@@ -207,31 +205,16 @@ def LLVM_AllocaOp : LLVM_Op<"alloca",
$res = $_builder.create<LLVM::AllocaOp>(
$_location, $_resultType, $arraySize,
alignment == 0 ? IntegerAttr() : $_builder.getI64IntegerAttr(alignment),
- TypeAttr::get(allocatedType), allocaInst->isUsedWithInAlloca());
+ allocatedType, allocaInst->isUsedWithInAlloca());
}];
let builders = [
- DeprecatedOpBuilder<"the usage of typed pointers is deprecated",
- (ins "Type":$resultType, "Value":$arraySize,
- "unsigned":$alignment),
- [{
- assert(!::llvm::cast<LLVMPointerType>(resultType).isOpaque() &&
- "pass the allocated type explicitly if opaque pointers are used");
- if (alignment == 0)
- return build($_builder, $_state, resultType, arraySize, IntegerAttr(),
- TypeAttr(), false);
- build($_builder, $_state, resultType, arraySize,
- $_builder.getI64IntegerAttr(alignment), TypeAttr(), false);
- }]>,
OpBuilder<(ins "Type":$resultType, "Type":$elementType, "Value":$arraySize,
CArg<"unsigned", "0">:$alignment),
[{
- TypeAttr elemTypeAttr =
- ::llvm::cast<LLVMPointerType>(resultType).isOpaque() ?
- TypeAttr::get(elementType) : TypeAttr();
build($_builder, $_state, resultType, arraySize,
alignment == 0 ? IntegerAttr()
: $_builder.getI64IntegerAttr(alignment),
- elemTypeAttr, false);
+ elementType, false);
}]>
];
@@ -247,7 +230,7 @@ def LLVM_GEPOp : LLVM_Op<"getelementptr", [Pure,
let arguments = (ins LLVM_ScalarOrVectorOf<LLVM_AnyPointer>:$base,
Variadic<LLVM_ScalarOrVectorOf<AnyInteger>>:$dynamicIndices,
DenseI32ArrayAttr:$rawConstantIndices,
- OptionalAttr<TypeAttr>:$elem_type,
+ TypeAttr:$elem_type,
UnitAttr:$inbounds);
let results = (outs LLVM_ScalarOrVectorOf<LLVM_AnyPointer>:$res);
let skipDefaultBuilders = 1;
@@ -282,14 +265,6 @@ def LLVM_GEPOp : LLVM_Op<"getelementptr", [Pure,
OpBuilder<(ins "Type":$resultType, "Type":$basePtrType, "Value":$basePtr,
"ValueRange":$indices, CArg<"bool", "false">:$inbounds,
CArg<"ArrayRef<NamedAttribute>", "{}">:$attributes)>,
- DeprecatedOpBuilder<"the usage of typed pointers is deprecated",
- (ins "Type":$resultType, "Value":$basePtr,
- "ValueRange":$indices, CArg<"bool", "false">:$inbounds,
- CArg<"ArrayRef<NamedAttribute>", "{}">:$attributes)>,
- DeprecatedOpBuilder<"the usage of typed pointers is deprecated",
- (ins "Type":$resultType, "Value":$basePtr,
- "ArrayRef<GEPArg>":$indices, CArg<"bool", "false">:$inbounds,
- CArg<"ArrayRef<NamedAttribute>", "{}">:$attributes)>,
OpBuilder<(ins "Type":$resultType, "Type":$basePtrType, "Value":$basePtr,
"ArrayRef<GEPArg>":$indices, CArg<"bool", "false">:$inbounds,
CArg<"ArrayRef<NamedAttribute>", "{}">:$attributes)>,
@@ -313,7 +288,7 @@ def LLVM_GEPOp : LLVM_Op<"getelementptr", [Pure,
let assemblyFormat = [{
(`inbounds` $inbounds^)?
$base `[` custom<GEPIndices>($dynamicIndices, $rawConstantIndices) `]` attr-dict
- `:` functional-type(operands, results) (`,` $elem_type^)?
+ `:` functional-type(operands, results) `,` $elem_type
}];
let extraClassDeclaration = [{
@@ -332,7 +307,7 @@ def LLVM_LoadOp : LLVM_MemAccessOpBase<"load",
[DeclareOpInterfaceMethods<MemoryEffectsOpInterface>,
DeclareOpInterfaceMethods<PromotableMemOpInterface>,
DeclareOpInterfaceMethods<SafeMemorySlotAccessOpInterface>]> {
- dag args = (ins LLVM_PointerTo<LLVM_LoadableType>:$addr,
+ dag args = (ins LLVM_AnyPointer:$addr,
OptionalAttr<I64Attr>:$alignment,
UnitAttr:$volatile_,
UnitAttr:$nontemporal,
@@ -370,7 +345,8 @@ def LLVM_LoadOp : LLVM_MemAccessOpBase<"load",
let assemblyFormat = [{
(`volatile` $volatile_^)? $addr
(`atomic` (`syncscope` `(` $syncscope^ `)`)? $ordering^)?
- attr-dict `:` custom<LoadType>(type($addr), type($res))
+ attr-dict `:` qualified(type($addr)) `->` type($res)
+
}];
string llvmBuilder = [{
auto *inst = builder.CreateLoad($_resultType, $addr, $volatile_);
@@ -391,9 +367,6 @@ def LLVM_LoadOp : LLVM_MemAccessOpBase<"load",
getLLVMSyncScope(loadInst));
}];
let builders = [
- DeprecatedOpBuilder<"the usage of typed pointers is deprecated",
- (ins "Value":$addr, CArg<"unsigned", "0">:$alignment,
- CArg<"bool", "false">:$isVolatile, CArg<"bool", "false">:$isNonTemporal)>,
OpBuilder<(ins "Type":$type, "Value":$addr,
CArg<"unsigned", "0">:$alignment, CArg<"bool", "false">:$isVolatile,
CArg<"bool", "false">:$isNonTemporal,
@@ -408,7 +381,7 @@ def LLVM_StoreOp : LLVM_MemAccessOpBase<"store",
DeclareOpInterfaceMethods<PromotableMemOpInterface>,
DeclareOpInterfaceMethods<SafeMemorySlotAccessOpInterface>]> {
dag args = (ins LLVM_LoadableType:$value,
- LLVM_PointerTo<LLVM_LoadableType>:$addr,
+ LLVM_AnyPointer:$addr,
OptionalAttr<I64Attr>:$alignment,
UnitAttr:$volatile_,
UnitAttr:$nontemporal,
@@ -445,7 +418,7 @@ def LLVM_StoreOp : LLVM_MemAccessOpBase<"store",
let assemblyFormat = [{
(`volatile` $volatile_^)? $value `,` $addr
(`atomic` (`syncscope` `(` $syncscope^ `)`)? $ordering^)?
- attr-dict `:` custom<StoreType>(type($value), type($addr))
+ attr-dict `:` type($value) `,` qualified(type($addr))
}];
string llvmBuilder = [{
auto *inst = builder.CreateStore($value, $addr, $volatile_);
@@ -651,8 +624,7 @@ def LLVM_CallOp : LLVM_MemAccessOpBase<"call",
OpBuilder<(ins "LLVMFunctionType":$calleeType, "FlatSymbolRefAttr":$callee,
CArg<"ValueRange", "{}">:$args)>,
OpBuilder<(ins "LLVMFunctionType":$calleeType, "StringRef":$callee,
- CArg<"ValueRange", "{}">:$args)>,
- OpBuilder<(ins "Value":$callee, "ValueRange":$args)>
+ CArg<"ValueRange", "{}">:$args)>
];
let hasCustomAssemblyFormat = 1;
let extraClassDeclaration = [{
@@ -1636,7 +1608,7 @@ def LLVM_AtomicRMWOp : LLVM_MemAccessOpBase<"atomicrmw", [
TypesMatchWith<"result #0 and operand #1 have the same type",
"val", "res", "$_self">]> {
dag args = (ins AtomicBinOp:$bin_op,
- LLVM_PointerTo<LLVM_AtomicRMWType>:$ptr,
+ LLVM_AnyPointer:$ptr,
LLVM_AtomicRMWType:$val, AtomicOrdering:$ordering,
OptionalAttr<StrAttr>:$syncscope,
OptionalAttr<I64Attr>:$alignment,
@@ -1687,7 +1659,7 @@ def LLVM_AtomicCmpXchgOp : LLVM_MemAccessOpBase<"cmpxchg", [
TypesMatchWith<"result #0 has an LLVM struct type consisting of "
"the type of operand #2 and a bool", "val", "res",
"getValAndBoolStructType($_self)">]> {
- dag args = (ins LLVM_PointerTo<LLVM_AtomicCmpXchgType>:$ptr,
+ dag args = (ins LLVM_AnyPointer:$ptr,
LLVM_AtomicCmpXchgType:$cmp, LLVM_AtomicCmpXchgType:$val,
AtomicOrdering:$success_ordering,
AtomicOrdering:$failure_ordering,
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
index e31029bfed5a54a..0bd068c1be7c90a 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
@@ -137,30 +137,17 @@ def LLVMPointerType : LLVMType<"LLVMPointer", "ptr", [
```
}];
- let parameters = (ins DefaultValuedParameter<"Type", "Type()">:$elementType,
- DefaultValuedParameter<"unsigned", "0">:$addressSpace);
+ let parameters = (ins DefaultValuedParameter<"unsigned", "0">:$addressSpace);
let assemblyFormat = [{
- (`<` custom<Pointer>($elementType, $addressSpace)^ `>`)?
+ (`<` $addressSpace^ `>`)?
}];
- let genVerifyDecl = 1;
-
+ let skipDefaultBuilders = 1;
let builders = [
- TypeBuilderWithInferredContext<(ins "Type":$elementType,
- CArg<"unsigned", "0">:$addressSpace)>,
TypeBuilder<(ins CArg<"unsigned", "0">:$addressSpace), [{
- return $_get($_ctxt, Type(), addressSpace);
+ return $_get($_ctxt, addressSpace);
}]>
];
-
- let extraClassDeclaration = [{
- /// Returns `true` if this type is the opaque pointer type, i.e., it has no
- /// pointed-to type.
- bool isOpaque() const { return !getElementType(); }
-
- /// Checks if the given type can have a pointer type pointing to it.
- static bool isValidElementType(Type type);
- }];
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index c49decde1638b1c..cc5c1bd7e199304 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -19,10 +19,8 @@ include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
include "mlir/Interfaces/SideEffectInterfaces.td"
include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td"
-def LLVM_i8Ptr_global : LLVM_IntPtrBase<8, 1>;
-def LLVM_i8Ptr_shared : LLVM_IntPtrBase<8, 3>;
-def LLVM_i64ptr_any : LLVM_IntPtrBase<64>;
-def LLVM_i64ptr_shared : LLVM_IntPtrBase<64, 3>;
+def LLVM_ptr_global : LLVM_OpaquePointerInAddressSpace<1>;
+def LLVM_ptr_shared : LLVM_OpaquePointerInAddressSpace<3>;
//===----------------------------------------------------------------------===//
// NVVM dialect definitions
@@ -213,7 +211,7 @@ def NVVM_ReduxOp :
/// mbarrier.init instruction with generic pointer type
def NVVM_MBarrierInitOp : NVVM_PTXBuilder_Op<"mbarrier.init">,
- Arguments<(ins LLVM_i64ptr_any:$addr, I32:$count, PtxPredicate:$predicate)> {
+ Arguments<(ins LLVM_AnyPointer:$addr, I32:$count, PtxPredicate:$predicate)> {
string llvmBuilder = [{
createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_init, {$addr, $count});
}];
@@ -228,7 +226,7 @@ def NVVM_MBarrierInitOp : NVVM_PTXBuilder_Op<"mbarrier.init">,
/// mbarrier.init instruction with shared pointer type
def NVVM_MBarrierInitSharedOp : NVVM_PTXBuilder_Op<"mbarrier.init.shared">,
- Arguments<(ins LLVM_i64ptr_shared:$addr, I32:$count, PtxPredicate:$predicate)> {
+ Arguments<(ins LLVM_ptr_shared:$addr, I32:$count, PtxPredicate:$predicate)> {
string llvmBuilder = [{
createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_init_shared, {$addr, $count});
}];
@@ -240,7 +238,7 @@ def NVVM_MBarrierInitSharedOp : NVVM_PTXBuilder_Op<"mbarrier.init.shared">,
}
def NVVM_MBarrierInvalOp : NVVM_Op<"mbarrier.inval">,
- Arguments<(ins LLVM_i64ptr_any:$addr)> {
+ Arguments<(ins LLVM_AnyPointer:$addr)> {
string llvmBuilder = [{
createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_inval, {$addr});
}];
@@ -248,7 +246,7 @@ def NVVM_MBarrierInvalOp : NVVM_Op<"mbarrier.inval">,
}
def NVVM_MBarrierInvalSharedOp : NVVM_Op<"mbarrier.inval.shared">,
- Arguments<(ins LLVM_i64ptr_shared:$addr)> {
+ Arguments<(ins LLVM_ptr_shared:$addr)> {
string llvmBuilder = [{
createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_inval_shared, {$addr});
}];
@@ -257,7 +255,7 @@ def NVVM_MBarrierInvalSharedOp : NVVM_Op<"mbarrier.inval.shared">,
def NVVM_MBarrierArriveOp : NVVM_Op<"mbarrier.arrive">,
Results<(outs LLVM_Type:$res)>,
- Arguments<(ins LLVM_i64ptr_any:$addr)> {
+ Arguments<(ins LLVM_AnyPointer:$addr)> {
string llvmBuilder = [{
$res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive, {$addr});
}];
@@ -266,16 +264,16 @@ def NVVM_MBarrierArriveOp : NVVM_Op<"mbarrier.arrive">,
def NVVM_MBarrierArriveSharedOp : NVVM_Op<"mbarrier.arrive.shared">,
Results<(outs LLVM_Type:$res)>,
- Arguments<(ins LLVM_i64ptr_shared:$addr)> {
+ Arguments<(ins LLVM_ptr_shared:$addr)> {
string llvmBuilder = [{
$res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_shared, {$addr});
}];
- let assemblyFormat = "$addr attr-dict `:` type($addr) `->` type($res)";
+ let assemblyFormat = "$addr attr-dict `:` qualified(type($addr)) `->` type($res)";
}
def NVVM_MBarrierArriveNocompleteOp : NVVM_Op<"mbarrier.arrive.nocomplete">,
Results<(outs LLVM_Type:$res)>,
- Arguments<(ins LLVM_i64ptr_any:$addr, I32:$count)> {
+ Arguments<(ins LLVM_AnyPointer:$addr, I32:$count)> {
string llvmBuilder = [{
$res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete, {$addr, $count});
}];
@@ -284,7 +282,7 @@ def NVVM_MBarrierArriveNocompleteOp : NVVM_Op<"mbarrier.arrive.nocomplete">,
def NVVM_MBarrierArriveNocompleteSharedOp : NVVM_Op<"mbarrier.arrive.nocomplete.shared">,
Results<(outs LLVM_Type:$res)>,
- Arguments<(ins LLVM_i64ptr_shared:$addr, I32:$count)> {
+ Arguments<(ins LLVM_ptr_shared:$addr, I32:$count)> {
string llvmBuilder = [{
$res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete_shared, {$addr, $count});
}];
@@ -292,7 +290,7 @@ def NVVM_MBarrierArriveNocompleteSharedOp : NVVM_Op<"mbarrier.arrive.nocomplete.
}
def NVVM_MBarrierArriveExpectTxOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_tx">,
- Arguments<(ins LLVM_i64ptr_any:$addr, I32:$txcount, PtxPredicate:$predicate)> {
+ Arguments<(ins LLVM_AnyPointer:$addr, I32:$txcount, PtxPredicate:$predicate)> {
let assemblyFormat = "$addr `,` $txcount (`,` `predicate` `=` $predicate^)? attr-dict `:` type(operands)";
let extraClassDefinition = [{
std::string $cppClass::getPtx() { return std::string("mbarrier.arrive.expect_tx.b64 _, [%0], %1;"); }
@@ -300,7 +298,7 @@ def NVVM_MBarrierArriveExpectTxOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_t
}
def NVVM_MBarrierArriveExpectTxSharedOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_tx.shared">,
- Arguments<(ins LLVM_i64ptr_shared:$addr, I32:$txcount, PtxPredicate:$predicate)> {
+ Arguments<(ins LLVM_ptr_shared:$addr, I32:$txcount, PtxPredicate:$predicate)> {
let assemblyFormat = "$addr `,` $txcount (`,` `predicate` `=` $predicate^)? attr-dict `:` type(operands)";
let extraClassDefinition = [{
std::string $cppClass::getPtx() { return std::string("mbarrier.arrive.expect_tx.shared.b64 _, [%0], %1;"); }
@@ -308,7 +306,7 @@ def NVVM_MBarrierArriveExpectTxSharedOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.ex
}
def NVVM_MBarrierTryWaitParityOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity">,
- Arguments<(ins LLVM_i64ptr_any:$addr, I32:$phase, I32:$ticks)> {
+ Arguments<(ins LLVM_AnyPointer:$addr, I32:$phase, I32:$ticks)> {
let assemblyFormat = "$addr `,` $phase `,` $ticks attr-dict `:` type(operands)";
let extraClassDefinition = [{
std::string $cppClass::getPtx() {
@@ -327,7 +325,7 @@ def NVVM_MBarrierTryWaitParityOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity"
}
def NVVM_MBarrierTryWaitParitySharedOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity.shared">,
- Arguments<(ins LLVM_i64ptr_shared:$addr, I32:$phase, I32:$ticks)> {
+ Arguments<(ins LLVM_ptr_shared:$addr, I32:$phase, I32:$ticks)> {
let assemblyFormat = "$addr `,` $phase `,` $ticks attr-dict `:` type(operands)";
let extraClassDefinition = [{
std::string $cppClass::getPtx() {
@@ -347,7 +345,7 @@ def NVVM_MBarrierTryWaitParitySharedOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.p
def NVVM_MBarrierTestWaitOp : NVVM_Op<"mbarrier.test.wait">,
Results<(outs LLVM_Type:$res)>,
- Arguments<(ins LLVM_i64ptr_any:$addr, LLVM_Type:$state)> {
+ Arguments<(ins LLVM_AnyPointer:$addr, LLVM_Type:$state)> {
string llvmBuilder = [{
$res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_test_wait, {$addr, $state});
}];
@@ -356,7 +354,7 @@ def NVVM_MBarrierTestWaitOp : NVVM_Op<"mbarrier.test.wait">,
def NVVM_MBarrierTestWaitSharedOp : NVVM_Op<"mbarrier.test.wait.shared">,
Results<(outs LLVM_Type:$res)>,
- Arguments<(ins LLVM_i64ptr_shared:$addr, LLVM_Type:$state)> {
+ Arguments<(ins LLVM_ptr_shared:$addr, LLVM_Type:$state)> {
string llvmBuilder = [{
$res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_test_wait_shared, {$addr, $state});
}];
@@ -501,8 +499,8 @@ def LoadCacheModifierKind : I32EnumAttr<"LoadCacheModifierKind",
def LoadCacheModifierAttr : EnumAttr<NVVM_Dialect, LoadCacheModifierKind, "load_cache_modifier">;
def NVVM_CpAsyncOp : NVVM_PTXBuilder_Op<"cp.async.shared.global">,
- Arguments<(ins LLVM_i8Ptr_shared:$dst,
- LLVM_i8Ptr_global:$src,
+ Arguments<(ins LLVM_ptr_shared:$dst,
+ LLVM_ptr_global:$src,
I32Attr:$size,
LoadCacheModifierAttr:$modifier,
Optional<LLVM_Type>:$cpSize)> {
@@ -1187,7 +1185,7 @@ def NVVM_WMMAMmaOp : NVVM_Op<"wmma.mma">,
}
def NVVM_StMatrixOp: NVVM_PTXBuilder_Op<"stmatrix">,
- Arguments<(ins LLVM_i8Ptr_shared:$ptr,
+ Arguments<(ins LLVM_ptr_shared:$ptr,
Variadic<I32>:$sources,
MMALayoutAttr:$layout)> {
let summary = "cooperative matrix store";
@@ -1404,9 +1402,9 @@ def NVVM_CpAsyncBulkTensorGlobalToSharedClusterOp :
NVVM_Op<"cp.async.bulk.tensor.shared.cluster.global",
[DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>,
AttrSizedOperandSegments]>,
- Arguments<(ins LLVM_i64ptr_shared:$dstMem,
- LLVM_i64ptr_any:$tmaDescriptor,
- LLVM_i64ptr_shared:$mbar,
+ Arguments<(ins LLVM_ptr_shared:$dstMem,
+ LLVM_AnyPointer:$tmaDescriptor,
+ LLVM_ptr_shared:$mbar,
Variadic<I32>:$coordinates,
PtxPredicate:$predicate)> {
let assemblyFormat = [{
@@ -1439,8 +1437,8 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
NVVM_Op<"cp.async.bulk.tensor.global.shared.cta",
[DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>,
AttrSizedOperandSegments]>,
- Arguments<(ins LLVM_i64ptr_any:$tmaDescriptor,
- LLVM_i64ptr_shared:$srcMem,
+ Arguments<(ins LLVM_AnyPointer:$tmaDescriptor,
+ LLVM_ptr_shared:$srcMem,
Variadic<I32>:$coordinates,
PtxPredicate:$predicate)> {
let assemblyFormat = [{
@@ -1469,7 +1467,7 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
def NVVM_PrefetchTensorMapOp : NVVM_Op<"prefetch.tensormap",
[DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]>,
- Arguments<(ins LLVM_i64ptr_any:$tmaDescriptor, PtxPredicate:$predicate)> {
+ Arguments<(ins LLVM_AnyPointer:$tmaDescriptor, PtxPredicate:$predicate)> {
let assemblyFormat = "$tmaDescriptor (`,` `predicate` `=` $predicate^)? attr-dict `:` type(operands)";
let extraClassDefinition = [{
std::string $cppClass::getPtx() {
diff --git a/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp b/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp
index 35b95d7a5ebe925..fdd1d993fdb8f62 100644
--- a/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp
+++ b/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp
@@ -78,13 +78,7 @@ LLVMTypeConverter::LLVMTypeConverter(MLIRContext *ctx,
// LLVM container types may (recursively) contain other types that must be
// converted even when the outer type is compatible.
- addConversion([&](LLVM::LLVMPointerType type) -> std::optional<Type> {
- if (type.isOpaque())
- return type;
- if (auto pointee = convertType(type.getElementType()))
- return LLVM::LLVMPointerType::get(pointee, type.getAddressSpace());
- return std::nullopt;
- });
+ addConversion([&](LLVM::LLVMPointerType type) { return type; });
addConversion([&](LLVM::LLVMStructType type, SmallVectorImpl<Type> &results)
-> std::optional<LogicalResult> {
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
index 7f5681e7bdc0592..3ce7e9abccaa387 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
@@ -216,17 +216,13 @@ OpFoldResult ICmpOp::fold(FoldAdaptor adaptor) {
//===----------------------------------------------------------------------===//
void AllocaOp::print(OpAsmPrinter &p) {
- Type elemTy = llvm::cast<LLVM::LLVMPointerType>(getType()).getElementType();
- if (!elemTy)
- elemTy = *getElemType();
-
auto funcTy =
FunctionType::get(getContext(), {getArraySize().getType()}, {getType()});
if (getInalloca())
p << " inalloca";
- p << ' ' << getArraySize() << " x " << elemTy;
+ p << ' ' << getArraySize() << " x " << getElemType();
if (getAlignment() && *getAlignment() != 0)
p.printOptionalAttrDict((*this)->getAttrs(),
{kElemTypeAttrName, getInallocaAttrName()});
@@ -277,40 +273,16 @@ ParseResult AllocaOp::parse(OpAsmParser &parser, OperationState &result) {
return failure();
Type resultType = funcType.getResult(0);
- if (auto ptrResultType = llvm::dyn_cast<LLVMPointerType>(resultType)) {
- if (ptrResultType.isOpaque())
- result.addAttribute(kElemTypeAttrName, TypeAttr::get(elemType));
- }
+ if (auto ptrResultType = llvm::dyn_cast<LLVMPointerType>(resultType))
+ result.addAttribute(kElemTypeAttrName, TypeAttr::get(elemType));
result.addTypes({funcType.getResult(0)});
return success();
}
-/// Checks that the elemental type is present in either the pointer type or
-/// the attribute, but not both.
-static LogicalResult verifyOpaquePtr(Operation *op, LLVMPointerType ptrType,
- std::optional<Type> ptrElementType) {
- if (ptrType.isOpaque() && !ptrElementType.has_value()) {
- return op->emitOpError() << "expected '" << kElemTypeAttrName
- << "' attribute if opaque pointer type is used";
- }
- if (!ptrType.isOpaque() && ptrElementType.has_value()) {
- return op->emitOpError()
- << "unexpected '" << kElemTypeAttrName
- << "' attribute when non-opaque pointer type is used";
- }
- return success();
-}
-
LogicalResult AllocaOp::verify() {
- LLVMPointerType ptrType = llvm::cast<LLVMPointerType>(getType());
- if (failed(verifyOpaquePtr(getOperation(), ptrType, getElemType())))
- return failure();
-
- Type elemTy =
- (ptrType.isOpaque()) ? *getElemType() : ptrType.getElementType();
// Only certain target extension types can be used in 'alloca'.
- if (auto targetExtType = dyn_cast<LLVMTargetExtType>(elemTy);
+ if (auto targetExtType = dyn_cast<LLVMTargetExtType>(getElemType());
targetExtType && !targetExtType.supportsMemOps())
return emitOpError()
<< "this target extension type cannot be used in alloca";
@@ -318,11 +290,7 @@ LogicalResult AllocaOp::verify() {
return success();
}
-Type AllocaOp::getResultPtrElementType() {
- // This will become trivial once non-opaque pointers are gone.
- return getElemType().has_value() ? *getElemType()
- : getResult().getType().getElementType();
-}
+Type AllocaOp::getResultPtrElementType() { return getElemType(); }
//===----------------------------------------------------------------------===//
// LLVM::BrOp
@@ -525,18 +493,6 @@ static Type extractVectorElementType(Type type) {
return type;
}
-void GEPOp::build(OpBuilder &builder, OperationState &result, Type resultType,
- Value basePtr, ArrayRef<GEPArg> indices, bool inbounds,
- ArrayRef<NamedAttribute> attributes) {
- auto ptrType =
- llvm::cast<LLVMPointerType>(extractVectorElementType(basePtr.getType()));
- assert(!ptrType.isOpaque() &&
- "expected non-opaque pointer, provide elementType explicitly when "
- "opaque pointers are used");
- build(builder, result, resultType, ptrType.getElementType(), basePtr, indices,
- inbounds, attributes);
-}
-
/// Destructures the 'indices' parameter into 'rawConstantIndices' and
/// 'dynamicIndices', encoding the former in the process. In the process,
/// dynamic indices which are used to index into a structure type are converted
@@ -602,20 +558,11 @@ void GEPOp::build(OpBuilder &builder, OperationState &result, Type resultType,
result.addAttribute(getInboundsAttrName(result.name),
builder.getUnitAttr());
}
- if (llvm::cast<LLVMPointerType>(extractVectorElementType(basePtr.getType()))
- .isOpaque())
- result.addAttribute(kElemTypeAttrName, TypeAttr::get(elementType));
+ result.addAttribute(kElemTypeAttrName, TypeAttr::get(elementType));
result.addOperands(basePtr);
result.addOperands(dynamicIndices);
}
-void GEPOp::build(OpBuilder &builder, OperationState &result, Type resultType,
- Value basePtr, ValueRange indices, bool inbounds,
- ArrayRef<NamedAttribute> attributes) {
- build(builder, result, resultType, basePtr, SmallVector<GEPArg>(indices),
- inbounds, attributes);
-}
-
void GEPOp::build(OpBuilder &builder, OperationState &result, Type resultType,
Type elementType, Value basePtr, ValueRange indices,
bool inbounds, ArrayRef<NamedAttribute> attributes) {
@@ -712,12 +659,6 @@ verifyStructIndices(Type baseGEPType, GEPIndicesAdaptor<ValueRange> indices,
}
LogicalResult LLVM::GEPOp::verify() {
- if (failed(verifyOpaquePtr(
- getOperation(),
- llvm::cast<LLVMPointerType>(extractVectorElementType(getType())),
- getElemType())))
- return failure();
-
if (static_cast<size_t>(
llvm::count(getRawConstantIndices(), kDynamicIndex)) !=
getDynamicIndices().size())
@@ -728,14 +669,7 @@ LogicalResult LLVM::GEPOp::verify() {
[&] { return emitOpError(); });
}
-Type LLVM::GEPOp::getSourceElementType() {
- if (std::optional<Type> elemType = getElemType())
- return *elemType;
-
- return llvm::cast<LLVMPointerType>(
- extractVectorElementType(getBase().getType()))
- .getElementType();
-}
+Type LLVM::GEPOp::getSourceElementType() { return getElemType(); }
Type GEPOp::getResultPtrElementType() {
// Set the initial type currently being used for indexing. This will be
@@ -836,14 +770,6 @@ LogicalResult LoadOp::verify() {
{AtomicOrdering::release, AtomicOrdering::acq_rel});
}
-void LoadOp::build(OpBuilder &builder, OperationState &state, Value addr,
- unsigned alignment, bool isVolatile, bool isNonTemporal) {
- auto type = llvm::cast<LLVMPointerType>(addr.getType()).getElementType();
- assert(type && "must provide explicit element type to the constructor "
- "when the pointer type is opaque");
- build(builder, state, type, addr, alignment, isVolatile, isNonTemporal);
-}
-
void LoadOp::build(OpBuilder &builder, OperationState &state, Type type,
Value addr, unsigned alignment, bool isVolatile,
bool isNonTemporal, AtomicOrdering ordering,
@@ -857,51 +783,6 @@ void LoadOp::build(OpBuilder &builder, OperationState &state, Type type,
/*tbaa=*/nullptr);
}
-// Extract the pointee type from the LLVM pointer type wrapped in MLIR. Return
-// the resulting type if any, null type if opaque pointers are used, and
-// std::nullopt if the given type is not the pointer type.
-static std::optional<Type>
-getLoadStoreElementType(OpAsmParser &parser, Type type, SMLoc trailingTypeLoc) {
- auto llvmTy = llvm::dyn_cast<LLVM::LLVMPointerType>(type);
- if (!llvmTy) {
- parser.emitError(trailingTypeLoc, "expected LLVM pointer type");
- return std::nullopt;
- }
- return llvmTy.getElementType();
-}
-
-/// Parses the LoadOp type either using the typed or opaque pointer format.
-// TODO: Drop once the typed pointer assembly format is not needed anymore.
-static ParseResult parseLoadType(OpAsmParser &parser, Type &type,
- Type &elementType) {
- SMLoc trailingTypeLoc;
- if (parser.getCurrentLocation(&trailingTypeLoc) || parser.parseType(type))
- return failure();
-
- std::optional<Type> pointerElementType =
- getLoadStoreElementType(parser, type, trailingTypeLoc);
- if (!pointerElementType)
- return failure();
- if (*pointerElementType) {
- elementType = *pointerElementType;
- return success();
- }
-
- if (parser.parseArrow() || parser.parseType(elementType))
- return failure();
- return success();
-}
-
-/// Prints the LoadOp type either using the typed or opaque pointer format.
-// TODO: Drop once the typed pointer assembly format is not needed anymore.
-static void printLoadType(OpAsmPrinter &printer, Operation *op, Type type,
- Type elementType) {
- printer << type;
- auto pointerType = cast<LLVMPointerType>(type);
- if (pointerType.isOpaque())
- printer << " -> " << elementType;
-}
-
//===----------------------------------------------------------------------===//
// StoreOp
//===----------------------------------------------------------------------===//
@@ -940,38 +821,6 @@ void StoreOp::build(OpBuilder &builder, OperationState &state, Value value,
/*alias_scopes=*/nullptr, /*noalias_scopes=*/nullptr, /*tbaa=*/nullptr);
}
-/// Parses the StoreOp type either using the typed or opaque pointer format.
-// TODO: Drop once the typed pointer assembly format is not needed anymore.
-static ParseResult parseStoreType(OpAsmParser &parser, Type &elementType,
- Type &type) {
- SMLoc trailingTypeLoc;
- if (parser.getCurrentLocation(&trailingTypeLoc) ||
- parser.parseType(elementType))
- return failure();
-
- if (succeeded(parser.parseOptionalComma()))
- return parser.parseType(type);
-
- // Extract the element type from the pointer type.
- type = elementType;
- std::optional<Type> pointerElementType =
- getLoadStoreElementType(parser, type, trailingTypeLoc);
- if (!pointerElementType)
- return failure();
- elementType = *pointerElementType;
- return success();
-}
-
-/// Prints the StoreOp type either using the typed or opaque pointer format.
-// TODO: Drop once the typed pointer assembly format is not needed anymore.
-static void printStoreType(OpAsmPrinter &printer, Operation *op,
- Type elementType, Type type) {
- auto pointerType = cast<LLVMPointerType>(type);
- if (pointerType.isOpaque())
- printer << elementType << ", ";
- printer << type;
-}
-
//===----------------------------------------------------------------------===//
// CallOp
//===----------------------------------------------------------------------===//
@@ -1055,22 +904,6 @@ void CallOp::build(OpBuilder &builder, OperationState &state, LLVMFuncOp func,
/*access_groups=*/nullptr, /*alias_scopes=*/nullptr,
/*noalias_scopes=*/nullptr, /*tbaa=*/nullptr);
}
-
-void CallOp::build(OpBuilder &builder, OperationState &state, Value callee,
- ValueRange args) {
- auto calleeType = cast<LLVMFunctionType>(
- cast<LLVMPointerType>(callee.getType()).getElementType());
- SmallVector<Value> operands;
- operands.reserve(1 + args.size());
- operands.push_back(callee);
- llvm::append_range(operands, args);
- return build(builder, state, getCallOpResultTypes(calleeType),
- TypeAttr::get(calleeType), FlatSymbolRefAttr(), operands,
- /*fastmathFlags=*/nullptr, /*branch_weights=*/nullptr,
- /*access_groups=*/nullptr, /*alias_scopes=*/nullptr,
- /*noalias_scopes=*/nullptr, /*tbaa=*/nullptr);
-}
-
CallInterfaceCallable CallOp::getCallableForCallee() {
// Direct call.
if (FlatSymbolRefAttr calleeAttr = getCalleeAttr())
@@ -1145,10 +978,7 @@ LogicalResult CallOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
return emitOpError("indirect call expects a pointer as callee: ")
<< getOperand(0).getType();
- if (ptrType.isOpaque())
- return success();
-
- fnType = ptrType.getElementType();
+ return success();
} else {
Operation *callee =
symbolTable.lookupNearestSymbolFrom(*this, calleeName.getAttr());
@@ -1848,17 +1678,6 @@ AddressOfOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
return emitOpError("pointer address space must match address space of the "
"referenced global");
- if (type.isOpaque())
- return success();
-
- if (global && type.getElementType() != global.getType())
- return emitOpError(
- "the type must be a pointer to the type of the referenced global");
-
- if (function && type.getElementType() != function.getFunctionType())
- return emitOpError(
- "the type must be a pointer to the type of the referenced function");
-
return success();
}
@@ -2135,9 +1954,6 @@ static bool isZeroAttribute(Attribute value) {
}
LogicalResult GlobalOp::verify() {
- if (!LLVMPointerType::isValidElementType(getType()))
- return emitOpError(
- "expects type to be a valid element type for an LLVM pointer");
if ((*this)->getParentOp() && !satisfiesLLVMModule((*this)->getParentOp()))
return emitOpError("must appear at the module level");
@@ -2733,11 +2549,7 @@ void AtomicRMWOp::build(OpBuilder &builder, OperationState &state,
}
LogicalResult AtomicRMWOp::verify() {
- auto ptrType = llvm::cast<LLVM::LLVMPointerType>(getPtr().getType());
auto valType = getVal().getType();
- if (!ptrType.isOpaque() && valType != ptrType.getElementType())
- return emitOpError("expected LLVM IR element type for operand #0 to "
- "match type for operand #1");
if (getBinOp() == AtomicBinOp::fadd || getBinOp() == AtomicBinOp::fsub ||
getBinOp() == AtomicBinOp::fmin || getBinOp() == AtomicBinOp::fmax) {
if (!mlir::LLVM::isCompatibleFloatingPointType(valType))
@@ -2790,9 +2602,6 @@ LogicalResult AtomicCmpXchgOp::verify() {
if (!ptrType)
return emitOpError("expected LLVM IR pointer type for operand #0");
auto valType = getVal().getType();
- if (!ptrType.isOpaque() && valType != ptrType.getElementType())
- return emitOpError("expected LLVM IR element type for operand #0 to "
- "match type for all other operands");
if (!isTypeCompatibleWithAtomicOp(valType,
/*isPointerTypeAllowed=*/true))
return emitOpError("unexpected LLVM IR type");
@@ -3185,14 +2994,7 @@ LogicalResult LLVMDialect::verifyParameterAttribute(Operation *op,
auto checkPointerTypeMatches = [&]() -> LogicalResult {
if (failed(checkPointerType()))
return failure();
- auto ptrType = llvm::cast<LLVMPointerType>(paramType);
- auto typeAttr = llvm::cast<TypeAttr>(paramAttr.getValue());
- if (!ptrType.isOpaque() && ptrType.getElementType() != typeAttr.getValue())
- return op->emitError()
- << name
- << " attribute attached to LLVM pointer argument of "
- "different type";
return success();
};
@@ -3317,8 +3119,7 @@ Operation *LLVMDialect::materializeConstant(OpBuilder &builder, Attribute value,
Value mlir::LLVM::createGlobalString(Location loc, OpBuilder &builder,
StringRef name, StringRef value,
- LLVM::Linkage linkage,
- bool useOpaquePointers) {
+ LLVM::Linkage linkage) {
assert(builder.getInsertionBlock() &&
builder.getInsertionBlock()->getParentOp() &&
"expected builder to point to a block constrained in an op");
@@ -3334,19 +3135,11 @@ Value mlir::LLVM::createGlobalString(Location loc, OpBuilder &builder,
loc, type, /*isConstant=*/true, linkage, name,
builder.getStringAttr(value), /*alignment=*/0);
- LLVMPointerType resultType;
- LLVMPointerType charPtr;
- if (!useOpaquePointers) {
- resultType = LLVMPointerType::get(type);
- charPtr = LLVMPointerType::get(IntegerType::get(ctx, 8));
- } else {
- resultType = charPtr = LLVMPointerType::get(ctx);
- }
-
+ LLVMPointerType ptrType = LLVMPointerType::get(ctx);
// Get the pointer to the first character in the global string.
- Value globalPtr = builder.create<LLVM::AddressOfOp>(loc, resultType,
- global.getSymNameAttr());
- return builder.create<LLVM::GEPOp>(loc, charPtr, type, globalPtr,
+ Value globalPtr =
+ builder.create<LLVM::AddressOfOp>(loc, ptrType, global.getSymNameAttr());
+ return builder.create<LLVM::GEPOp>(loc, ptrType, type, globalPtr,
ArrayRef<GEPArg>{0, 0});
}
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp
index 2cdc5d9fff84e78..cc2a917e58f6394 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp
@@ -32,7 +32,7 @@ llvm::SmallVector<MemorySlot> LLVM::AllocaOp::getPromotableSlots() {
if (!getOperation()->getBlock()->isEntryBlock())
return {};
- return {MemorySlot{getResult(), getResultPtrElementType()}};
+ return {MemorySlot{getResult(), getElemType()}};
}
Value LLVM::AllocaOp::getDefaultValue(const MemorySlot &slot,
@@ -61,8 +61,7 @@ SmallVector<DestructurableMemorySlot> LLVM::AllocaOp::getDestructurableSlots() {
if (!mlir::matchPattern(getArraySize(), m_One()))
return {};
- Type elemType = getResultPtrElementType();
- auto destructurable = dyn_cast<DestructurableTypeInterface>(elemType);
+ auto destructurable = dyn_cast<DestructurableTypeInterface>(getElemType());
if (!destructurable)
return {};
@@ -75,7 +74,8 @@ SmallVector<DestructurableMemorySlot> LLVM::AllocaOp::getDestructurableSlots() {
for (Attribute index : llvm::make_first_range(destructuredType.value()))
allocaTypeMap.insert({index, LLVM::LLVMPointerType::get(getContext())});
- return {DestructurableMemorySlot{{getResult(), elemType}, {allocaTypeMap}}};
+ return {
+ DestructurableMemorySlot{{getResult(), getElemType()}, {allocaTypeMap}}};
}
DenseMap<Attribute, MemorySlot>
@@ -83,12 +83,9 @@ LLVM::AllocaOp::destructure(const DestructurableMemorySlot &slot,
const SmallPtrSetImpl<Attribute> &usedIndices,
RewriterBase &rewriter) {
assert(slot.ptr == getResult());
- Type elemType =
- getElemType() ? *getElemType() : getResult().getType().getElementType();
-
rewriter.setInsertionPointAfter(*this);
- auto destructurableType = cast<DestructurableTypeInterface>(elemType);
+ auto destructurableType = cast<DestructurableTypeInterface>(getElemType());
DenseMap<Attribute, MemorySlot> slotMap;
for (Attribute index : usedIndices) {
Type elemType = destructurableType.getTypeAtIndex(index);
@@ -337,11 +334,6 @@ bool LLVM::GEPOp::canRewire(const DestructurableMemorySlot &slot,
if (!basePtrType)
return false;
- // Typed pointers are not supported. This should be removed once typed
- // pointers are removed from the LLVM dialect.
- if (!basePtrType.isOpaque())
- return false;
-
if (getBase() != slot.ptr || slot.elemType != getElemType())
return false;
if (!isFirstIndexZero(*this))
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp
index bc8300a8b7329ea..8841aa8362569a0 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp
@@ -75,40 +75,6 @@ static void printFunctionTypes(AsmPrinter &p, ArrayRef<Type> params,
p << ')';
}
-//===----------------------------------------------------------------------===//
-// custom<Pointer>
-//===----------------------------------------------------------------------===//
-
-static ParseResult parsePointer(AsmParser &p, Type &elementType,
- unsigned &addressSpace) {
- // `<` addressSpace `>`
- OptionalParseResult result = p.parseOptionalInteger(addressSpace);
- if (result.has_value()) {
- if (failed(result.value()))
- return failure();
- elementType = Type();
- return success();
- }
-
- if (parsePrettyLLVMType(p, elementType))
- return failure();
- if (succeeded(p.parseOptionalComma()))
- return p.parseInteger(addressSpace);
-
- return success();
-}
-
-static void printPointer(AsmPrinter &p, Type elementType,
- unsigned addressSpace) {
- if (elementType)
- printPrettyLLVMType(p, elementType);
- if (addressSpace != 0) {
- if (elementType)
- p << ", ";
- p << addressSpace;
- }
-}
-
//===----------------------------------------------------------------------===//
// custom<ExtTypeParams>
//===----------------------------------------------------------------------===//
@@ -285,33 +251,6 @@ LLVMFunctionType::verify(function_ref<InFlightDiagnostic()> emitError,
return success();
}
-//===----------------------------------------------------------------------===//
-// LLVMPointerType
-//===----------------------------------------------------------------------===//
-
-bool LLVMPointerType::isValidElementType(Type type) {
- if (!type)
- return true;
- return isCompatibleOuterType(type)
- ? !llvm::isa<LLVMVoidType, LLVMTokenType, LLVMMetadataType,
- LLVMLabelType>(type)
- : llvm::isa<PointerElementTypeInterface>(type);
-}
-
-LLVMPointerType LLVMPointerType::get(Type pointee, unsigned addressSpace) {
- assert(pointee && "expected non-null subtype, pass the context instead if "
- "the opaque pointer type is desired");
- return Base::get(pointee.getContext(), pointee, addressSpace);
-}
-
-LogicalResult
-LLVMPointerType::verify(function_ref<InFlightDiagnostic()> emitError,
- Type pointee, unsigned) {
- if (!isValidElementType(pointee))
- return emitError() << "invalid pointer element type: " << pointee;
- return success();
-}
-
//===----------------------------------------------------------------------===//
// DataLayoutTypeInterface
@@ -369,9 +308,7 @@ LLVMPointerType::getTypeSizeInBits(const DataLayout &dataLayout,
// For other memory spaces, use the size of the pointer to the default memory
// space.
- if (isOpaque())
- return dataLayout.getTypeSizeInBits(get(getContext()));
- return dataLayout.getTypeSizeInBits(get(getElementType()));
+ return dataLayout.getTypeSizeInBits(get(getContext()));
}
unsigned LLVMPointerType::getABIAlignment(const DataLayout &dataLayout,
@@ -380,9 +317,7 @@ unsigned LLVMPointerType::getABIAlignment(const DataLayout &dataLayout,
getPointerDataLayoutEntry(params, *this, PtrDLEntryPos::Abi))
return *alignment;
- if (isOpaque())
- return dataLayout.getTypeABIAlignment(get(getContext()));
- return dataLayout.getTypeABIAlignment(get(getElementType()));
+ return dataLayout.getTypeABIAlignment(get(getContext()));
}
unsigned
@@ -392,9 +327,7 @@ LLVMPointerType::getPreferredAlignment(const DataLayout &dataLayout,
getPointerDataLayoutEntry(params, *this, PtrDLEntryPos::Preferred))
return *alignment;
- if (isOpaque())
- return dataLayout.getTypePreferredAlignment(get(getContext()));
- return dataLayout.getTypePreferredAlignment(get(getElementType()));
+ return dataLayout.getTypePreferredAlignment(get(getContext()));
}
bool LLVMPointerType::areCompatible(DataLayoutEntryListRef oldLayout,
@@ -440,7 +373,6 @@ LogicalResult LLVMPointerType::verifyEntries(DataLayoutEntryListRef entries,
for (DataLayoutEntryInterface entry : entries) {
if (!entry.isTypeEntry())
continue;
- auto key = llvm::cast<LLVMPointerType>(entry.getKey().get<Type>());
auto values = llvm::dyn_cast<DenseIntElementsAttr>(entry.getValue());
if (!values || (values.size() != 3 && values.size() != 4)) {
return emitError(loc)
@@ -448,10 +380,6 @@ LogicalResult LLVMPointerType::verifyEntries(DataLayoutEntryListRef entries,
<< " to be a dense integer elements attribute with 3 or 4 "
"elements";
}
- if (key.getElementType() && !key.getElementType().isInteger(8)) {
- return emitError(loc) << "unexpected layout attribute for pointer to "
- << key.getElementType();
- }
if (extractPointerSpecValue(values, PtrDLEntryPos::Abi) >
extractPointerSpecValue(values, PtrDLEntryPos::Preferred)) {
return emitError(loc) << "preferred alignment is expected to be at least "
@@ -869,11 +797,7 @@ static bool isCompatibleImpl(Type type, DenseSet<Type> &compatibleTypes) {
return vecType.getRank() == 1 &&
isCompatible(vecType.getElementType());
})
- .Case<LLVMPointerType>([&](auto pointerType) {
- if (pointerType.isOpaque())
- return true;
- return isCompatible(pointerType.getElementType());
- })
+ .Case<LLVMPointerType>([&](auto pointerType) { return true; })
.Case<LLVMTargetExtType>([&](auto extType) {
return llvm::all_of(extType.getTypeParams(), isCompatible);
})
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
index 9731689e551762b..ee491a6c558b888 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
@@ -92,10 +92,6 @@ LogicalResult AddFieldGetterToStructDirectUse<LoadOp>::matchAndRewrite(
LoadOp load, PatternRewriter &rewriter) const {
PatternRewriter::InsertionGuard guard(rewriter);
- // Load from typed pointers are not supported.
- if (!load.getAddr().getType().isOpaque())
- return failure();
-
Type inconsistentElementType =
isElementTypeInconsistent(load.getAddr(), load.getType());
if (!inconsistentElementType)
@@ -129,10 +125,6 @@ LogicalResult AddFieldGetterToStructDirectUse<StoreOp>::matchAndRewrite(
StoreOp store, PatternRewriter &rewriter) const {
PatternRewriter::InsertionGuard guard(rewriter);
- // Store to typed pointers are not supported.
- if (!store.getAddr().getType().isOpaque())
- return failure();
-
Type inconsistentElementType =
isElementTypeInconsistent(store.getAddr(), store.getValue().getType());
if (!inconsistentElementType)
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index df64d561f46cb3e..ff2eb9ebfc42f5e 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -27,14 +27,19 @@ using namespace acc;
#include "mlir/Dialect/OpenACC/OpenACCTypeInterfaces.cpp.inc"
namespace {
-/// Model for pointer-like types that already provide a `getElementType` method.
-template <typename T>
-struct PointerLikeModel
- : public PointerLikeType::ExternalModel<PointerLikeModel<T>, T> {
+struct MemRefPointerLikeModel
+ : public PointerLikeType::ExternalModel<MemRefPointerLikeModel,
+ MemRefType> {
Type getElementType(Type pointer) const {
- return llvm::cast<T>(pointer).getElementType();
+ return llvm::cast<MemRefType>(pointer).getElementType();
}
};
+
+struct LLVMPointerPointerLikeModel
+ : public PointerLikeType::ExternalModel<LLVMPointerPointerLikeModel,
+ LLVM::LLVMPointerType> {
+ Type getElementType(Type pointer) const { return Type(); }
+};
} // namespace
//===----------------------------------------------------------------------===//
@@ -58,9 +63,9 @@ void OpenACCDialect::initialize() {
// 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.
- LLVM::LLVMPointerType::attachInterface<
- PointerLikeModel<LLVM::LLVMPointerType>>(*getContext());
- MemRefType::attachInterface<PointerLikeModel<MemRefType>>(*getContext());
+ MemRefType::attachInterface<MemRefPointerLikeModel>(*getContext());
+ LLVM::LLVMPointerType::attachInterface<LLVMPointerPointerLikeModel>(
+ *getContext());
}
//===----------------------------------------------------------------------===//
@@ -1023,17 +1028,13 @@ void EnterDataOp::getCanonicalizationPatterns(RewritePatternSet &results,
// AtomicReadOp
//===----------------------------------------------------------------------===//
-LogicalResult AtomicReadOp::verify() {
- return verifyCommon();
-}
+LogicalResult AtomicReadOp::verify() { return verifyCommon(); }
//===----------------------------------------------------------------------===//
// AtomicWriteOp
//===----------------------------------------------------------------------===//
-LogicalResult AtomicWriteOp::verify() {
- return verifyCommon();
-}
+LogicalResult AtomicWriteOp::verify() { return verifyCommon(); }
//===----------------------------------------------------------------------===//
// AtomicUpdateOp
@@ -1054,13 +1055,9 @@ LogicalResult AtomicUpdateOp::canonicalize(AtomicUpdateOp op,
return failure();
}
-LogicalResult AtomicUpdateOp::verify() {
- return verifyCommon();
-}
+LogicalResult AtomicUpdateOp::verify() { return verifyCommon(); }
-LogicalResult AtomicUpdateOp::verifyRegions() {
- return verifyRegionsCommon();
-}
+LogicalResult AtomicUpdateOp::verifyRegions() { return verifyRegionsCommon(); }
//===----------------------------------------------------------------------===//
// AtomicCaptureOp
@@ -1084,9 +1081,7 @@ AtomicUpdateOp AtomicCaptureOp::getAtomicUpdateOp() {
return dyn_cast<AtomicUpdateOp>(getSecondOp());
}
-LogicalResult AtomicCaptureOp::verifyRegions() {
- return verifyRegionsCommon();
-}
+LogicalResult AtomicCaptureOp::verifyRegions() { return verifyRegionsCommon(); }
//===----------------------------------------------------------------------===//
// DeclareEnterOp
diff --git a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
index 3b792a26d1823fe..f6757aba664f985 100644
--- a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
+++ b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
@@ -39,15 +39,20 @@ using namespace mlir;
using namespace mlir::omp;
namespace {
-/// Model for pointer-like types that already provide a `getElementType` method.
-template <typename T>
-struct PointerLikeModel
- : public PointerLikeType::ExternalModel<PointerLikeModel<T>, T> {
+struct MemRefPointerLikeModel
+ : public PointerLikeType::ExternalModel<MemRefPointerLikeModel,
+ MemRefType> {
Type getElementType(Type pointer) const {
- return llvm::cast<T>(pointer).getElementType();
+ return llvm::cast<MemRefType>(pointer).getElementType();
}
};
+struct LLVMPointerPointerLikeModel
+ : public PointerLikeType::ExternalModel<LLVMPointerPointerLikeModel,
+ LLVM::LLVMPointerType> {
+ Type getElementType(Type pointer) const { return Type(); }
+};
+
struct OpenMPDialectFoldInterface : public DialectFoldInterface {
using DialectFoldInterface::DialectFoldInterface;
@@ -73,11 +78,9 @@ void OpenMPDialect::initialize() {
>();
addInterface<OpenMPDialectFoldInterface>();
- LLVM::LLVMPointerType::attachInterface<
- PointerLikeModel<LLVM::LLVMPointerType>>(*getContext());
- MemRefType::attachInterface<PointerLikeModel<MemRefType>>(*getContext());
- LLVM::LLVMPointerType::attachInterface<
- PointerLikeModel<LLVM::LLVMPointerType>>(*getContext());
+ MemRefType::attachInterface<MemRefPointerLikeModel>(*getContext());
+ LLVM::LLVMPointerType::attachInterface<LLVMPointerPointerLikeModel>(
+ *getContext());
// Attach default offload module interface to module op to access
// offload functionality through
@@ -1342,9 +1345,7 @@ LogicalResult AtomicUpdateOp::verify() {
return verifySynchronizationHint(*this, getHintVal());
}
-LogicalResult AtomicUpdateOp::verifyRegions() {
- return verifyRegionsCommon();
-}
+LogicalResult AtomicUpdateOp::verifyRegions() { return verifyRegionsCommon(); }
//===----------------------------------------------------------------------===//
// Verifier for AtomicCaptureOp
diff --git a/mlir/test/Dialect/LLVMIR/global.mlir b/mlir/test/Dialect/LLVMIR/global.mlir
index a33fff3967e4d59..e5d7d6d8f8c40f8 100644
--- a/mlir/test/Dialect/LLVMIR/global.mlir
+++ b/mlir/test/Dialect/LLVMIR/global.mlir
@@ -117,11 +117,6 @@ llvm.mlir.global internal protected unnamed_addr @protected(42 : i32) : i32
// -----
-// expected-error @+1 {{expects type to be a valid element type for an LLVM pointer}}
-llvm.mlir.global internal constant @constant(37.0) : !llvm.label
-
-// -----
-
// expected-error @+1 {{'addr_space' failed to satisfy constraint: 32-bit signless integer attribute whose value is non-negative}}
"llvm.mlir.global"() ({}) {sym_name = "foo", global_type = i64, value = 42 : i64, addr_space = -1 : i32, linkage = #llvm.linkage<private>} : () -> ()
diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir
index fe2f94454561a08..1d51796abb03faf 100644
--- a/mlir/test/Dialect/LLVMIR/invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/invalid.mlir
@@ -90,30 +90,23 @@ func.func @alloca_non_integer_alignment() {
// -----
-func.func @alloca_opaque_ptr_no_type(%sz : i64) {
- // expected-error at below {{expected 'elem_type' attribute if opaque pointer type is used}}
- "llvm.alloca"(%sz) : (i64) -> !llvm.ptr
-}
-
-// -----
-
func.func @gep_missing_input_result_type(%pos : i64, %base : !llvm.ptr) {
// expected-error at +1 {{2 operands present, but expected 0}}
- llvm.getelementptr %base[%pos] : () -> ()
+ llvm.getelementptr %base[%pos] : () -> (), i64
}
// -----
func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr) {
// expected-error at +1 {{2 operands present, but expected 0}}
- llvm.getelementptr %base[%pos] : () -> (!llvm.ptr)
+ llvm.getelementptr %base[%pos] : () -> (!llvm.ptr), i64
}
// -----
func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr) {
// expected-error at +1 {{op requires one result}}
- llvm.getelementptr %base[%pos] : (!llvm.ptr, i64) -> ()
+ llvm.getelementptr %base[%pos] : (!llvm.ptr, i64) -> (), i64
}
// -----
@@ -132,20 +125,6 @@ func.func @gep_too_few_dynamic(%base : !llvm.ptr) {
// -----
-func.func @load_non_llvm_type(%foo : memref<f32>) {
- // expected-error at +1 {{expected LLVM pointer type}}
- llvm.load %foo : memref<f32>
-}
-
-// -----
-
-func.func @load_non_ptr_type(%foo : f32) {
- // expected-error at +1 {{expected LLVM pointer type}}
- llvm.load %foo : f32
-}
-
-// -----
-
func.func @load_syncscope(%ptr : !llvm.ptr) {
// expected-error at below {{expected syncscope to be null for non-atomic access}}
%1 = "llvm.load"(%ptr) {syncscope = "singlethread"} : (!llvm.ptr) -> (f32)
@@ -181,27 +160,6 @@ func.func @load_unaligned_atomic(%ptr : !llvm.ptr) {
// -----
-func.func @store_non_llvm_type(%foo : memref<f32>, %bar : f32) {
- // expected-error at +1 {{expected LLVM pointer type}}
- llvm.store %bar, %foo : memref<f32>
-}
-
-// -----
-
-func.func @store_non_ptr_type(%foo : f32, %bar : f32) {
- // expected-error at +1 {{expected LLVM pointer type}}
- llvm.store %bar, %foo : f32
-}
-
-// -----
-
-func.func @store_malformed_elem_type(%foo: !llvm.ptr, %bar: f32) {
- // expected-error at +1 {{expected non-function type}}
- llvm.store %bar, %foo : !llvm.ptr, "f32"
-}
-
-// -----
-
func.func @store_syncscope(%val : f32, %ptr : !llvm.ptr) {
// expected-error at below {{expected syncscope to be null for non-atomic access}}
"llvm.store"(%val, %ptr) {syncscope = "singlethread"} : (f32, !llvm.ptr) -> ()
@@ -632,14 +590,6 @@ func.func @nvvm_invalid_mma_8(%a0 : i32, %a1 : i32,
// -----
-func.func @atomicrmw_expected_ptr(%f32 : f32) {
- // expected-error at +1 {{operand #0 must be LLVM pointer to floating point LLVM type or LLVM pointer type or integer}}
- %0 = "llvm.atomicrmw"(%f32, %f32) {bin_op=11, ordering=1} : (f32, f32) -> f32
- llvm.return
-}
-
-// -----
-
func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %f32 : f32) {
// expected-error at +1 {{op failed to verify that result #0 and operand #1 have the same type}}
%0 = "llvm.atomicrmw"(%f32_ptr, %f32) {bin_op=11, ordering=1} : (!llvm.ptr, f32) -> i32
@@ -672,14 +622,6 @@ func.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr, %f32 : f32) {
// -----
-func.func @cmpxchg_expected_ptr(%f32 : f32) {
- // expected-error at +1 {{op operand #0 must be LLVM pointer to integer or LLVM pointer type}}
- %0 = "llvm.cmpxchg"(%f32, %f32, %f32) {success_ordering=2,failure_ordering=2} : (f32, f32, f32) -> !llvm.struct<(f32, i1)>
- llvm.return
-}
-
-// -----
-
func.func @cmpxchg_mismatched_value_operands(%ptr : !llvm.ptr, %i32 : i32, %i64 : i64) {
// expected-error at +1 {{op failed to verify that operand #1 and operand #2 have the same type}}
%0 = "llvm.cmpxchg"(%ptr, %i32, %i64) {success_ordering=2,failure_ordering=2} : (!llvm.ptr, i32, i64) -> !llvm.struct<(i32, i1)>
diff --git a/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp b/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp
index 2d9f8d0e6074990..083dec819a0e054 100644
--- a/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp
+++ b/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp
@@ -17,37 +17,3 @@ TEST_F(LLVMIRTest, IsStructTypeMutable) {
ASSERT_TRUE(bool(structTy));
ASSERT_TRUE(structTy.hasTrait<TypeTrait::IsMutable>());
}
-
-TEST_F(LLVMIRTest, MutualReferencedSubElementTypes) {
- auto fooStructTy = LLVMStructType::getIdentified(&context, "foo");
- ASSERT_TRUE(bool(fooStructTy));
- auto barStructTy = LLVMStructType::getIdentified(&context, "bar");
- ASSERT_TRUE(bool(barStructTy));
-
- // Created two structs that are referencing each other.
- Type fooBody[] = {LLVMPointerType::get(barStructTy)};
- ASSERT_TRUE(succeeded(fooStructTy.setBody(fooBody, /*isPacked=*/false)));
- Type barBody[] = {LLVMPointerType::get(fooStructTy)};
- ASSERT_TRUE(succeeded(barStructTy.setBody(barBody, /*isPacked=*/false)));
-
- // Test if walkSubElements goes into infinite loops.
- SmallVector<Type, 4> subElementTypes;
- fooStructTy.walk([&](Type type) { subElementTypes.push_back(type); });
- ASSERT_EQ(subElementTypes.size(), 4U);
-
- // !llvm.ptr
- ASSERT_TRUE(isa<LLVMPointerType>(subElementTypes[0]));
-
- // !llvm.struct<"bar",...>
- auto structType = dyn_cast<LLVMStructType>(subElementTypes[1]);
- ASSERT_TRUE(bool(structType));
- ASSERT_TRUE(structType.getName().equals("bar"));
-
- // !llvm.ptr
- ASSERT_TRUE(isa<LLVMPointerType>(subElementTypes[2]));
-
- // !llvm.struct<"foo",...>
- structType = dyn_cast<LLVMStructType>(subElementTypes[3]);
- ASSERT_TRUE(bool(structType));
- ASSERT_TRUE(structType.getName().equals("foo"));
-}
>From a1ee18fa4332f2e25092ef056424bcf43bb63579 Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christianulmann at gmail.com>
Date: Sat, 4 Nov 2023 19:46:02 +0100
Subject: [PATCH 2/5] address review comments
---
mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td | 9 +++++----
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 4 ++--
mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td | 2 +-
mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp | 7 +++++++
mlir/test/Dialect/LLVMIR/global.mlir | 4 ++++
mlir/test/Dialect/LLVMIR/invalid.mlir | 7 +++++++
6 files changed, 26 insertions(+), 7 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
index 0166fbb647b41b5..a459000581028ae 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
@@ -56,10 +56,11 @@ def LLVM_AnyPointer : Type<CPred<"::llvm::isa<::mlir::LLVM::LLVMPointerType>($_s
"LLVM pointer type", "::mlir::LLVM::LLVMPointerType">;
// Opaque pointer in a given address space.
-class LLVM_OpaquePointerInAddressSpace<int addressSpace> : Type<
- CPred<
- "::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace() == "
- # addressSpace>,
+class LLVM_PointerInAddressSpace<int addressSpace> : Type<
+ And<[LLVM_AnyPointer.predicate,
+ CPred<
+ "::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace() == "
+ # addressSpace>]>,
"Opaque LLVM pointer in address space " # addressSpace,
"::mlir::LLVM::LLVMPointerType"> {
let builderCall = "$_builder.getType<::mlir::LLVM::LLVMPointerType>("
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index cc5c1bd7e199304..b305738fa483c68 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -19,8 +19,8 @@ include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
include "mlir/Interfaces/SideEffectInterfaces.td"
include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td"
-def LLVM_ptr_global : LLVM_OpaquePointerInAddressSpace<1>;
-def LLVM_ptr_shared : LLVM_OpaquePointerInAddressSpace<3>;
+def LLVM_ptr_global : LLVM_PointerInAddressSpace<1>;
+def LLVM_ptr_shared : LLVM_PointerInAddressSpace<3>;
//===----------------------------------------------------------------------===//
// NVVM dialect definitions
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
index 6c6419bf238b457..48b830ae34f2922 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
@@ -275,7 +275,7 @@ def ROCDL_wmma_i32_16x16x16_iu4 : ROCDL_Wmma_IntrOp<"wmma.i32.16x16x16.iu4">;
// raw buffer mode).
//===---------------------------------------------------------------------===//
-def ROCDLBufferRsrc : LLVM_OpaquePointerInAddressSpace<8>;
+def ROCDLBufferRsrc : LLVM_PointerInAddressSpace<8>;
def ROCDL_MakeBufferRsrcOp :
ROCDL_IntrOp<"make.buffer.rsrc", [], [0], [Pure], 1>,
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
index 3ce7e9abccaa387..33af73d8a4ac5fb 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
@@ -1954,6 +1954,13 @@ static bool isZeroAttribute(Attribute value) {
}
LogicalResult GlobalOp::verify() {
+ bool validType = isCompatibleOuterType(getType())
+ ? !llvm::isa<LLVMVoidType, LLVMTokenType,
+ LLVMMetadataType, LLVMLabelType>(getType())
+ : llvm::isa<PointerElementTypeInterface>(getType());
+ if (!validType)
+ return emitOpError(
+ "expects type to be a valid element type for an LLVM global");
if ((*this)->getParentOp() && !satisfiesLLVMModule((*this)->getParentOp()))
return emitOpError("must appear at the module level");
diff --git a/mlir/test/Dialect/LLVMIR/global.mlir b/mlir/test/Dialect/LLVMIR/global.mlir
index e5d7d6d8f8c40f8..81178b2ef901f4f 100644
--- a/mlir/test/Dialect/LLVMIR/global.mlir
+++ b/mlir/test/Dialect/LLVMIR/global.mlir
@@ -117,6 +117,10 @@ llvm.mlir.global internal protected unnamed_addr @protected(42 : i32) : i32
// -----
+// expected-error @+1 {{expects type to be a valid element type for an LLVM global}}
+llvm.mlir.global internal constant @constant(37.0) : !llvm.label
+
+// -----
// expected-error @+1 {{'addr_space' failed to satisfy constraint: 32-bit signless integer attribute whose value is non-negative}}
"llvm.mlir.global"() ({}) {sym_name = "foo", global_type = i64, value = 42 : i64, addr_space = -1 : i32, linkage = #llvm.linkage<private>} : () -> ()
diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir
index 1d51796abb03faf..d6960f00f0428cb 100644
--- a/mlir/test/Dialect/LLVMIR/invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/invalid.mlir
@@ -125,6 +125,13 @@ func.func @gep_too_few_dynamic(%base : !llvm.ptr) {
// -----
+func.func @load_non_llvm_type(%foo : memref<f32>) {
+ // expected-error at +1 {{op operand #0 must be LLVM pointer type}}
+ llvm.load %foo : memref<f32> -> f32
+}
+
+// -----
+
func.func @load_syncscope(%ptr : !llvm.ptr) {
// expected-error at below {{expected syncscope to be null for non-atomic access}}
%1 = "llvm.load"(%ptr) {syncscope = "singlethread"} : (!llvm.ptr) -> (f32)
>From 95a551989c7ee307a126b1278010cc18b4825af7 Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christian.ulmann at nextsilicon.com>
Date: Mon, 6 Nov 2023 07:39:20 +0000
Subject: [PATCH 3/5] fix comments & remove identity type conversion
---
mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td | 4 ++--
mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp | 4 ----
2 files changed, 2 insertions(+), 6 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
index a459000581028ae..4e42a0e46d9bf9c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
@@ -55,13 +55,13 @@ def LLVM_AnyFloat : Type<
def LLVM_AnyPointer : Type<CPred<"::llvm::isa<::mlir::LLVM::LLVMPointerType>($_self)">,
"LLVM pointer type", "::mlir::LLVM::LLVMPointerType">;
-// Opaque pointer in a given address space.
+// Pointer in a given address space.
class LLVM_PointerInAddressSpace<int addressSpace> : Type<
And<[LLVM_AnyPointer.predicate,
CPred<
"::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace() == "
# addressSpace>]>,
- "Opaque LLVM pointer in address space " # addressSpace,
+ "LLVM pointer in address space " # addressSpace,
"::mlir::LLVM::LLVMPointerType"> {
let builderCall = "$_builder.getType<::mlir::LLVM::LLVMPointerType>("
# addressSpace # ")";
diff --git a/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp b/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp
index fdd1d993fdb8f62..04496d6b8f63449 100644
--- a/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp
+++ b/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp
@@ -76,10 +76,6 @@ LLVMTypeConverter::LLVMTypeConverter(MLIRContext *ctx,
: std::nullopt;
});
- // LLVM container types may (recursively) contain other types that must be
- // converted even when the outer type is compatible.
- addConversion([&](LLVM::LLVMPointerType type) { return type; });
-
addConversion([&](LLVM::LLVMStructType type, SmallVectorImpl<Type> &results)
-> std::optional<LogicalResult> {
// Fastpath for types that won't be converted by this callback anyway.
>From 62cfdef3dbbbc5546c89dde0f749e7cc8327d796 Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christian.ulmann at nextsilicon.com>
Date: Mon, 6 Nov 2023 08:22:37 +0000
Subject: [PATCH 4/5] remove useles GEP helper function
---
mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td | 6 +-----
mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp | 8 +++-----
mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp | 9 ++++-----
3 files changed, 8 insertions(+), 15 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
index ec196e1bd13673d..cb8a6e08886df90 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
@@ -281,7 +281,7 @@ def LLVM_GEPOp : LLVM_Op<"getelementptr", [Pure,
indices.push_back(
builder.getInt32(valueOrAttr.get<IntegerAttr>().getInt()));
}
- Type baseElementType = op.getSourceElementType();
+ Type baseElementType = op.getElemType();
llvm::Type *elementType = moduleTranslation.convertType(baseElementType);
$res = builder.CreateGEP(elementType, $base, indices, "", $inbounds);
}];
@@ -294,9 +294,6 @@ def LLVM_GEPOp : LLVM_Op<"getelementptr", [Pure,
let extraClassDeclaration = [{
constexpr static int32_t kDynamicIndex = std::numeric_limits<int32_t>::min();
- /// Returns the type pointed to by the pointer argument of this GEP.
- Type getSourceElementType();
-
GEPIndicesAdaptor<ValueRange> getIndices();
}];
let hasFolder = 1;
@@ -346,7 +343,6 @@ def LLVM_LoadOp : LLVM_MemAccessOpBase<"load",
(`volatile` $volatile_^)? $addr
(`atomic` (`syncscope` `(` $syncscope^ `)`)? $ordering^)?
attr-dict `:` qualified(type($addr)) `->` type($res)
-
}];
string llvmBuilder = [{
auto *inst = builder.CreateLoad($_resultType, $addr, $volatile_);
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
index 33af73d8a4ac5fb..c22cff4c1328a55 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
@@ -665,16 +665,14 @@ LogicalResult LLVM::GEPOp::verify() {
return emitOpError("expected as many dynamic indices as specified in '")
<< getRawConstantIndicesAttrName().getValue() << "'";
- return verifyStructIndices(getSourceElementType(), getIndices(),
+ return verifyStructIndices(getElemType(), getIndices(),
[&] { return emitOpError(); });
}
-Type LLVM::GEPOp::getSourceElementType() { return getElemType(); }
-
Type GEPOp::getResultPtrElementType() {
// Set the initial type currently being used for indexing. This will be
// updated as the indices get walked over.
- Type selectedType = getSourceElementType();
+ Type selectedType = getElemType();
// Follow the indexed elements in the gep.
auto indices = getIndices();
@@ -2803,7 +2801,7 @@ OpFoldResult LLVM::GEPOp::fold(FoldAdaptor adaptor) {
if (changed) {
SmallVector<int32_t> rawConstantIndices;
SmallVector<Value> dynamicIndices;
- destructureIndices(getSourceElementType(), gepArgs, rawConstantIndices,
+ destructureIndices(getElemType(), gepArgs, rawConstantIndices,
dynamicIndices);
getDynamicIndicesMutable().assign(dynamicIndices);
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
index ee491a6c558b888..b094c650ff1932c 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp
@@ -164,9 +164,9 @@ static std::optional<uint64_t> gepToByteOffset(DataLayout &layout, GEPOp gep) {
indices.push_back(indexInt.getInt());
}
- uint64_t offset = indices[0] * layout.getTypeSize(gep.getSourceElementType());
+ uint64_t offset = indices[0] * layout.getTypeSize(gep.getElemType());
- Type currentType = gep.getSourceElementType();
+ Type currentType = gep.getElemType();
for (uint32_t index : llvm::drop_begin(indices)) {
bool shouldCancel =
TypeSwitch<Type, bool>(currentType)
@@ -571,7 +571,7 @@ LogicalResult SplitStores::matchAndRewrite(StoreOp store,
return failure();
offset = *byteOffset;
- typeHint = gepOp.getSourceElementType();
+ typeHint = gepOp.getElemType();
address = gepOp.getBase();
}
}
@@ -653,8 +653,7 @@ LogicalResult SplitGEP::matchAndRewrite(GEPOp gepOp,
// Split of the first GEP using the first two indices.
auto subGepOp = rewriter.create<GEPOp>(
- gepOp.getLoc(), gepOp.getType(), gepOp.getSourceElementType(),
- gepOp.getBase(),
+ gepOp.getLoc(), gepOp.getType(), gepOp.getElemType(), gepOp.getBase(),
llvm::map_to_vector(llvm::make_range(indices.begin(), splitIter),
indexToGEPArg),
gepOp.getInbounds());
>From 1ecf826e64011b3647debe1e23eafcc4d14ddae8 Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christian.ulmann at nextsilicon.com>
Date: Mon, 6 Nov 2023 09:01:23 +0000
Subject: [PATCH 5/5] harmonize NNVM type constraints
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 30 ++++++++++-----------
1 file changed, 15 insertions(+), 15 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index b305738fa483c68..ffe6f25fcd944b6 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -19,8 +19,8 @@ include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
include "mlir/Interfaces/SideEffectInterfaces.td"
include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td"
-def LLVM_ptr_global : LLVM_PointerInAddressSpace<1>;
-def LLVM_ptr_shared : LLVM_PointerInAddressSpace<3>;
+def LLVM_PointerGlobal : LLVM_PointerInAddressSpace<1>;
+def LLVM_PointerShared : LLVM_PointerInAddressSpace<3>;
//===----------------------------------------------------------------------===//
// NVVM dialect definitions
@@ -226,7 +226,7 @@ def NVVM_MBarrierInitOp : NVVM_PTXBuilder_Op<"mbarrier.init">,
/// mbarrier.init instruction with shared pointer type
def NVVM_MBarrierInitSharedOp : NVVM_PTXBuilder_Op<"mbarrier.init.shared">,
- Arguments<(ins LLVM_ptr_shared:$addr, I32:$count, PtxPredicate:$predicate)> {
+ Arguments<(ins LLVM_PointerShared:$addr, I32:$count, PtxPredicate:$predicate)> {
string llvmBuilder = [{
createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_init_shared, {$addr, $count});
}];
@@ -246,7 +246,7 @@ def NVVM_MBarrierInvalOp : NVVM_Op<"mbarrier.inval">,
}
def NVVM_MBarrierInvalSharedOp : NVVM_Op<"mbarrier.inval.shared">,
- Arguments<(ins LLVM_ptr_shared:$addr)> {
+ Arguments<(ins LLVM_PointerShared:$addr)> {
string llvmBuilder = [{
createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_inval_shared, {$addr});
}];
@@ -264,7 +264,7 @@ def NVVM_MBarrierArriveOp : NVVM_Op<"mbarrier.arrive">,
def NVVM_MBarrierArriveSharedOp : NVVM_Op<"mbarrier.arrive.shared">,
Results<(outs LLVM_Type:$res)>,
- Arguments<(ins LLVM_ptr_shared:$addr)> {
+ Arguments<(ins LLVM_PointerShared:$addr)> {
string llvmBuilder = [{
$res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_shared, {$addr});
}];
@@ -282,7 +282,7 @@ def NVVM_MBarrierArriveNocompleteOp : NVVM_Op<"mbarrier.arrive.nocomplete">,
def NVVM_MBarrierArriveNocompleteSharedOp : NVVM_Op<"mbarrier.arrive.nocomplete.shared">,
Results<(outs LLVM_Type:$res)>,
- Arguments<(ins LLVM_ptr_shared:$addr, I32:$count)> {
+ Arguments<(ins LLVM_PointerShared:$addr, I32:$count)> {
string llvmBuilder = [{
$res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete_shared, {$addr, $count});
}];
@@ -298,7 +298,7 @@ def NVVM_MBarrierArriveExpectTxOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_t
}
def NVVM_MBarrierArriveExpectTxSharedOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_tx.shared">,
- Arguments<(ins LLVM_ptr_shared:$addr, I32:$txcount, PtxPredicate:$predicate)> {
+ Arguments<(ins LLVM_PointerShared:$addr, I32:$txcount, PtxPredicate:$predicate)> {
let assemblyFormat = "$addr `,` $txcount (`,` `predicate` `=` $predicate^)? attr-dict `:` type(operands)";
let extraClassDefinition = [{
std::string $cppClass::getPtx() { return std::string("mbarrier.arrive.expect_tx.shared.b64 _, [%0], %1;"); }
@@ -325,7 +325,7 @@ def NVVM_MBarrierTryWaitParityOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity"
}
def NVVM_MBarrierTryWaitParitySharedOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity.shared">,
- Arguments<(ins LLVM_ptr_shared:$addr, I32:$phase, I32:$ticks)> {
+ Arguments<(ins LLVM_PointerShared:$addr, I32:$phase, I32:$ticks)> {
let assemblyFormat = "$addr `,` $phase `,` $ticks attr-dict `:` type(operands)";
let extraClassDefinition = [{
std::string $cppClass::getPtx() {
@@ -354,7 +354,7 @@ def NVVM_MBarrierTestWaitOp : NVVM_Op<"mbarrier.test.wait">,
def NVVM_MBarrierTestWaitSharedOp : NVVM_Op<"mbarrier.test.wait.shared">,
Results<(outs LLVM_Type:$res)>,
- Arguments<(ins LLVM_ptr_shared:$addr, LLVM_Type:$state)> {
+ Arguments<(ins LLVM_PointerShared:$addr, LLVM_Type:$state)> {
string llvmBuilder = [{
$res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_test_wait_shared, {$addr, $state});
}];
@@ -499,8 +499,8 @@ def LoadCacheModifierKind : I32EnumAttr<"LoadCacheModifierKind",
def LoadCacheModifierAttr : EnumAttr<NVVM_Dialect, LoadCacheModifierKind, "load_cache_modifier">;
def NVVM_CpAsyncOp : NVVM_PTXBuilder_Op<"cp.async.shared.global">,
- Arguments<(ins LLVM_ptr_shared:$dst,
- LLVM_ptr_global:$src,
+ Arguments<(ins LLVM_PointerShared:$dst,
+ LLVM_PointerGlobal:$src,
I32Attr:$size,
LoadCacheModifierAttr:$modifier,
Optional<LLVM_Type>:$cpSize)> {
@@ -1185,7 +1185,7 @@ def NVVM_WMMAMmaOp : NVVM_Op<"wmma.mma">,
}
def NVVM_StMatrixOp: NVVM_PTXBuilder_Op<"stmatrix">,
- Arguments<(ins LLVM_ptr_shared:$ptr,
+ Arguments<(ins LLVM_PointerShared:$ptr,
Variadic<I32>:$sources,
MMALayoutAttr:$layout)> {
let summary = "cooperative matrix store";
@@ -1402,9 +1402,9 @@ def NVVM_CpAsyncBulkTensorGlobalToSharedClusterOp :
NVVM_Op<"cp.async.bulk.tensor.shared.cluster.global",
[DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>,
AttrSizedOperandSegments]>,
- Arguments<(ins LLVM_ptr_shared:$dstMem,
+ Arguments<(ins LLVM_PointerShared:$dstMem,
LLVM_AnyPointer:$tmaDescriptor,
- LLVM_ptr_shared:$mbar,
+ LLVM_PointerShared:$mbar,
Variadic<I32>:$coordinates,
PtxPredicate:$predicate)> {
let assemblyFormat = [{
@@ -1438,7 +1438,7 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
[DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>,
AttrSizedOperandSegments]>,
Arguments<(ins LLVM_AnyPointer:$tmaDescriptor,
- LLVM_ptr_shared:$srcMem,
+ LLVM_PointerShared:$srcMem,
Variadic<I32>:$coordinates,
PtxPredicate:$predicate)> {
let assemblyFormat = [{
More information about the Mlir-commits
mailing list