[MLIR][LLVM] Remove typed pointers from the LLVM dialect #71285
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
@llvm/pr-subscribers-mlir-openmp @llvm/pr-subscribers-mlir-llvm Author: Christian Ulmann (Dinistro) ChangesThis 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: This change includes:
Patch is 66.79 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/71285.diff 16 Files Affected:
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 638c31b39682ea6..e12bd0ee8d9d471 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 5f746e21e831b59..f6ed05640a5fff1 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
@@ -139,30 +139,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 6947cf10e3600d4..406519d3eaa7272 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)> {
+ Ar...
[truncated]
|
|
@llvm/pr-subscribers-mlir Author: Christian Ulmann (Dinistro) ChangesThis 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: This change includes:
Patch is 66.79 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/71285.diff 16 Files Affected:
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 638c31b39682ea6..e12bd0ee8d9d471 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 5f746e21e831b59..f6ed05640a5fff1 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
@@ -139,30 +139,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 6947cf10e3600d4..406519d3eaa7272 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)> {
+ Ar...
[truncated]
|
Dinistro
commented
Nov 4, 2023
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
Dinistro
force-pushed
the
remove-typed-pointers
branch
from
f78ccce to
a1ee18f
Compare
|
Thanks for the review. I'll keep this PR open for a while, such that others can also take a look. |
ftynse
approved these changes
Nov 6, 2023
| @@ -445,7 +414,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)) | |||
There was a problem hiding this comment.
Without qualified, this just prints <1>, instead of !llvm.ptr<1>. Tbh, I do not fully understand the logic behind this, but I've observed similar issues in other contexts.
There was a problem hiding this comment.
@ftynse are you fine with this change or should we get to the bottom of this issue?
There was a problem hiding this comment.
I'd like to understand better why and when is this necessary (the feature looks poorly documented), but this should not block this from landing.
There was a problem hiding this comment.
This behaviour was introduced by ee09087. As far as I can tell it essentially always omits the dialect prefix if the attribute/type class being parsed is known. Since in this Op it is known to be a LLVMPointerType it'll elide the !llvm.ptr by default when printing.
63f0c00 then apparently introduced the qualified directive to be able to restore the behaviour to before where the dialect prefix is always printed.
Personally I think this is non-intuative and that the default should be printing the full type by default instead, but that is another discussion entirely. Using qualified seems to be the intended way to fully qualify the type in the assembly format.
There was a problem hiding this comment.
This is WAI: qualified is a poor default IMO.
For example why would !llvm be a better choice here? And even if it is, why shouldn't we express the choice?
Should we even print the type when the address space is the default one? Lot of possible simplifications...
There was a problem hiding this comment.
The problem is that not only !llvm is omitted, but the ptr prefix as well. It therefore prints just <1> when using an address space rather than either ptr<1> or !llvm.ptr<1>. This IMO makes it non-obvious what the type is, nor what the meaning of the <1> is supposed to be.
There was a problem hiding this comment.
Of course! But the answer is contextual...
For example it's not obvious to me that qualified is just the best answer here, what about the following for example?
let assemblyFormat = ` .... attr-dict `:` type($value) `,` `ptr` `` type($addr)`
I see the design of the assembly format as... a design! It needs to be intentional, and from this point of view we need the most basic blocks to construct a nice format.
There was a problem hiding this comment.
The main reason I added qualified in this revision is to ensure as few breaking changes as possible, considering that the pointer type changes already forced me to change hundreds of lines.
I agree that qualified might not be the best idea here, as the "llvm." part brings no benefit. The format you propose would in my opinion be a much better fit. The question remains if we should always print the address space or not, and if we should even print the pointer type if the address space is zero.
Note that we should also consider all other operations where I used qualified in this discussion, i.e., LLVM_StoreOp, LLVM_AtomicRMWOp, LLVM_AtomicCmpXchgOp, NVVM_MBarrierArriveSharedOp.
There was a problem hiding this comment.
To be clear I wasn't criticizing this PR: I was providing a rational for "building new dialects".
rmoyard
added a commit
to PennyLaneAI/catalyst
that referenced
this pull request
Feb 22, 2024
**Description of the Change:** NullOp -> ZeroOp llvm/llvm-project#67183 CodeGenOpt -> CodeGenOptLevel https://github.com/llvm/llvm-project/pull/66295/files Converters are const llvm/llvm-project@ce25459 Remove typed pointer llvm/llvm-project#71285 thlo gml_st on ice tensorflow/mlir-hlo@7077cec ByteCode is a new interface llvm/llvm-project@837d1ce getArgsOperandMutable was added CallOpInterface llvm/llvm-project@d790a21 Remove redundant memref.tensor_store op https://github.com/llvm/llvm-project/pull/71010/files Lowering of memref.realloc was extracted to its own pass https://reviews.llvm.org/D159430 **Benefits:** **Possible Drawbacks:** **Related GitHub Issues:** --------- Co-authored-by: erick-xanadu <110487834+erick-xanadu@users.noreply.github.com> Co-authored-by: David Ittah <dime10@users.noreply.github.com>
rauletorresc
pushed a commit
to PennyLaneAI/catalyst
that referenced
this pull request
Feb 26, 2024
**Description of the Change:** NullOp -> ZeroOp llvm/llvm-project#67183 CodeGenOpt -> CodeGenOptLevel https://github.com/llvm/llvm-project/pull/66295/files Converters are const llvm/llvm-project@ce25459 Remove typed pointer llvm/llvm-project#71285 thlo gml_st on ice tensorflow/mlir-hlo@7077cec ByteCode is a new interface llvm/llvm-project@837d1ce getArgsOperandMutable was added CallOpInterface llvm/llvm-project@d790a21 Remove redundant memref.tensor_store op https://github.com/llvm/llvm-project/pull/71010/files Lowering of memref.realloc was extracted to its own pass https://reviews.llvm.org/D159430 **Benefits:** **Possible Drawbacks:** **Related GitHub Issues:** --------- Co-authored-by: erick-xanadu <110487834+erick-xanadu@users.noreply.github.com> Co-authored-by: David Ittah <dime10@users.noreply.github.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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:
LLVMPointerType