[Mlir-commits] [mlir] [MLIR][LLVM] Remove typed pointers from the LLVM dialect (PR #71285)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Sat Nov 4 07:27:13 PDT 2023
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir
Author: Christian Ulmann (Dinistro)
<details>
<summary>Changes</summary>
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
---
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:
- (modified) mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h (+1-2)
- (modified) mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td (+14-14)
- (modified) mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td (+1-27)
- (modified) mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td (+14-42)
- (modified) mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td (+4-17)
- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+26-28)
- (modified) mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp (+1-7)
- (modified) mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp (+13-220)
- (modified) mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp (+5-13)
- (modified) mlir/lib/Dialect/LLVMIR/IR/LLVMTypes.cpp (+4-80)
- (modified) mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp (-8)
- (modified) mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp (+18-23)
- (modified) mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp (+14-13)
- (modified) mlir/test/Dialect/LLVMIR/global.mlir (-5)
- (modified) mlir/test/Dialect/LLVMIR/invalid.mlir (+3-61)
- (modified) mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp (-34)
``````````diff
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]
``````````
</details>
https://github.com/llvm/llvm-project/pull/71285
More information about the Mlir-commits
mailing list