From de69c84f48a68fd17ba062f595869e63b86fb58c Mon Sep 17 00:00:00 2001 From: Christian Ulmann 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 bbed1ea5cf622..06df4a601b7a3 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 72c932ac07a2e..1123466b7a75e 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:$ptr, + let arguments = (ins LLVM_AnyPointer:$ptr, LLVM_AnyPointer:$annotation, LLVM_AnyPointer:$fileName, I32:$line, LLVM_AnyPointer:$attr); - let results = (outs LLVM_PointerTo:$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 503988abfc090..0166fbb647b41 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($_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< - 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 : 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 : Type< - And<[LLVM_PointerTo>.predicate, - CPred<"::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace()" - " == " # addressSpace>]>, - "LLVM pointer to " # I.summary>; - -def LLVM_i8Ptr : LLVM_IntPtrBase<8>; - // Type constraint accepting any LLVM structure type. def LLVM_AnyStruct : Type($_self)">, "LLVM structure type">; diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td index c8549f146d029..ec196e1bd1367 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:$alignment, - OptionalAttr:$elem_type, + TypeAttr:$elem_type, UnitAttr:$inalloca); let results = (outs Res]>:$res); string llvmInstName = "Alloca"; string llvmBuilder = [{ auto addrSpace = $_resultType->getPointerAddressSpace(); - llvm::Type *elementType = moduleTranslation.convertType( - $elem_type ? *$elem_type - : ::llvm::cast(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( $_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(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(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:$base, Variadic>:$dynamicIndices, DenseI32ArrayAttr:$rawConstantIndices, - OptionalAttr:$elem_type, + TypeAttr:$elem_type, UnitAttr:$inbounds); let results = (outs LLVM_ScalarOrVectorOf:$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", "{}">:$attributes)>, - DeprecatedOpBuilder<"the usage of typed pointers is deprecated", - (ins "Type":$resultType, "Value":$basePtr, - "ValueRange":$indices, CArg<"bool", "false">:$inbounds, - CArg<"ArrayRef", "{}">:$attributes)>, - DeprecatedOpBuilder<"the usage of typed pointers is deprecated", - (ins "Type":$resultType, "Value":$basePtr, - "ArrayRef":$indices, CArg<"bool", "false">:$inbounds, - CArg<"ArrayRef", "{}">:$attributes)>, OpBuilder<(ins "Type":$resultType, "Type":$basePtrType, "Value":$basePtr, "ArrayRef":$indices, CArg<"bool", "false">:$inbounds, CArg<"ArrayRef", "{}">:$attributes)>, @@ -313,7 +288,7 @@ def LLVM_GEPOp : LLVM_Op<"getelementptr", [Pure, let assemblyFormat = [{ (`inbounds` $inbounds^)? $base `[` custom($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, DeclareOpInterfaceMethods, DeclareOpInterfaceMethods]> { - dag args = (ins LLVM_PointerTo:$addr, + dag args = (ins LLVM_AnyPointer:$addr, OptionalAttr:$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(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, DeclareOpInterfaceMethods]> { dag args = (ins LLVM_LoadableType:$value, - LLVM_PointerTo:$addr, + LLVM_AnyPointer:$addr, OptionalAttr:$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(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:$ptr, + LLVM_AnyPointer:$ptr, LLVM_AtomicRMWType:$val, AtomicOrdering:$ordering, OptionalAttr:$syncscope, OptionalAttr:$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:$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 e31029bfed5a5..0bd068c1be7c9 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($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 c49decde1638b..cc5c1bd7e1993 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; 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:$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:$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, 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:$coordinates, PtxPredicate:$predicate)> { let assemblyFormat = [{ @@ -1439,8 +1437,8 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp : NVVM_Op<"cp.async.bulk.tensor.global.shared.cta", [DeclareOpInterfaceMethods, AttrSizedOperandSegments]>, - Arguments<(ins LLVM_i64ptr_any:$tmaDescriptor, - LLVM_i64ptr_shared:$srcMem, + Arguments<(ins LLVM_AnyPointer:$tmaDescriptor, + LLVM_ptr_shared:$srcMem, Variadic:$coordinates, PtxPredicate:$predicate)> { let assemblyFormat = [{ @@ -1469,7 +1467,7 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp : def NVVM_PrefetchTensorMapOp : NVVM_Op<"prefetch.tensormap", [DeclareOpInterfaceMethods]>, - 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 35b95d7a5ebe9..fdd1d993fdb8f 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 { - 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 &results) -> std::optional { diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp index 7f5681e7bdc05..3ce7e9abccaa3 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(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(resultType)) { - if (ptrResultType.isOpaque()) - result.addAttribute(kElemTypeAttrName, TypeAttr::get(elemType)); - } + if (auto ptrResultType = llvm::dyn_cast(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 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(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(elemTy); + if (auto targetExtType = dyn_cast(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 indices, bool inbounds, - ArrayRef attributes) { - auto ptrType = - llvm::cast(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(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 attributes) { - build(builder, result, resultType, basePtr, SmallVector(indices), - inbounds, attributes); -} - void GEPOp::build(OpBuilder &builder, OperationState &result, Type resultType, Type elementType, Value basePtr, ValueRange indices, bool inbounds, ArrayRef attributes) { @@ -712,12 +659,6 @@ verifyStructIndices(Type baseGEPType, GEPIndicesAdaptor indices, } LogicalResult LLVM::GEPOp::verify() { - if (failed(verifyOpaquePtr( - getOperation(), - llvm::cast(extractVectorElementType(getType())), - getElemType()))) - return failure(); - if (static_cast( llvm::count(getRawConstantIndices(), kDynamicIndex)) != getDynamicIndices().size()) @@ -728,14 +669,7 @@ LogicalResult LLVM::GEPOp::verify() { [&] { return emitOpError(); }); } -Type LLVM::GEPOp::getSourceElementType() { - if (std::optional elemType = getElemType()) - return *elemType; - - return llvm::cast( - 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(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 -getLoadStoreElementType(OpAsmParser &parser, Type type, SMLoc trailingTypeLoc) { - auto llvmTy = llvm::dyn_cast(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 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(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 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(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( - cast(callee.getType()).getElementType()); - SmallVector 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(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(paramType); - auto typeAttr = llvm::cast(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(loc, resultType, - global.getSymNameAttr()); - return builder.create(loc, charPtr, type, globalPtr, + Value globalPtr = + builder.create(loc, ptrType, global.getSymNameAttr()); + return builder.create(loc, ptrType, type, globalPtr, ArrayRef{0, 0}); } diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp index 2cdc5d9fff84e..cc2a917e58f63 100644 --- a/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp +++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMMemorySlot.cpp @@ -32,7 +32,7 @@ llvm::SmallVector 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 LLVM::AllocaOp::getDestructurableSlots() { if (!mlir::matchPattern(getArraySize(), m_One())) return {}; - Type elemType = getResultPtrElementType(); - auto destructurable = dyn_cast(elemType); + auto destructurable = dyn_cast(getElemType()); if (!destructurable) return {}; @@ -75,7 +74,8 @@ SmallVector 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 @@ -83,12 +83,9 @@ LLVM::AllocaOp::destructure(const DestructurableMemorySlot &slot, const SmallPtrSetImpl &usedIndices, RewriterBase &rewriter) { assert(slot.ptr == getResult()); - Type elemType = - getElemType() ? *getElemType() : getResult().getType().getElementType(); - rewriter.setInsertionPointAfter(*this); - auto destructurableType = cast(elemType); + auto destructurableType = cast(getElemType()); DenseMap 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 bc8300a8b7329..8841aa8362569 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 params, p << ')'; } -//===----------------------------------------------------------------------===// -// custom -//===----------------------------------------------------------------------===// - -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 //===----------------------------------------------------------------------===// @@ -285,33 +251,6 @@ LLVMFunctionType::verify(function_ref emitError, return success(); } -//===----------------------------------------------------------------------===// -// LLVMPointerType -//===----------------------------------------------------------------------===// - -bool LLVMPointerType::isValidElementType(Type type) { - if (!type) - return true; - return isCompatibleOuterType(type) - ? !llvm::isa(type) - : llvm::isa(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 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(entry.getKey().get()); auto values = llvm::dyn_cast(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 &compatibleTypes) { return vecType.getRank() == 1 && isCompatible(vecType.getElementType()); }) - .Case([&](auto pointerType) { - if (pointerType.isOpaque()) - return true; - return isCompatible(pointerType.getElementType()); - }) + .Case([&](auto pointerType) { return true; }) .Case([&](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 9731689e55176..ee491a6c558b8 100644 --- a/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp @@ -92,10 +92,6 @@ LogicalResult AddFieldGetterToStructDirectUse::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::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 df64d561f46cb..ff2eb9ebfc42f 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 -struct PointerLikeModel - : public PointerLikeType::ExternalModel, T> { +struct MemRefPointerLikeModel + : public PointerLikeType::ExternalModel { Type getElementType(Type pointer) const { - return llvm::cast(pointer).getElementType(); + return llvm::cast(pointer).getElementType(); } }; + +struct LLVMPointerPointerLikeModel + : public PointerLikeType::ExternalModel { + 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>(*getContext()); - MemRefType::attachInterface>(*getContext()); + MemRefType::attachInterface(*getContext()); + LLVM::LLVMPointerType::attachInterface( + *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(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 3b792a26d1823..f6757aba664f9 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 -struct PointerLikeModel - : public PointerLikeType::ExternalModel, T> { +struct MemRefPointerLikeModel + : public PointerLikeType::ExternalModel { Type getElementType(Type pointer) const { - return llvm::cast(pointer).getElementType(); + return llvm::cast(pointer).getElementType(); } }; +struct LLVMPointerPointerLikeModel + : public PointerLikeType::ExternalModel { + Type getElementType(Type pointer) const { return Type(); } +}; + struct OpenMPDialectFoldInterface : public DialectFoldInterface { using DialectFoldInterface::DialectFoldInterface; @@ -73,11 +78,9 @@ void OpenMPDialect::initialize() { >(); addInterface(); - LLVM::LLVMPointerType::attachInterface< - PointerLikeModel>(*getContext()); - MemRefType::attachInterface>(*getContext()); - LLVM::LLVMPointerType::attachInterface< - PointerLikeModel>(*getContext()); + MemRefType::attachInterface(*getContext()); + LLVM::LLVMPointerType::attachInterface( + *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 a33fff3967e4d..e5d7d6d8f8c40 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} : () -> () diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir index fe2f94454561a..1d51796abb03f 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@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@+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@+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@+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) { - // expected-error@+1 {{expected LLVM pointer type}} - llvm.load %foo : memref -} - -// ----- - -func.func @load_non_ptr_type(%foo : f32) { - // expected-error@+1 {{expected LLVM pointer type}} - llvm.load %foo : f32 -} - -// ----- - func.func @load_syncscope(%ptr : !llvm.ptr) { // expected-error@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, %bar : f32) { - // expected-error@+1 {{expected LLVM pointer type}} - llvm.store %bar, %foo : memref -} - -// ----- - -func.func @store_non_ptr_type(%foo : f32, %bar : f32) { - // expected-error@+1 {{expected LLVM pointer type}} - llvm.store %bar, %foo : f32 -} - -// ----- - -func.func @store_malformed_elem_type(%foo: !llvm.ptr, %bar: f32) { - // expected-error@+1 {{expected non-function type}} - llvm.store %bar, %foo : !llvm.ptr, "f32" -} - -// ----- - func.func @store_syncscope(%val : f32, %ptr : !llvm.ptr) { // expected-error@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@+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@+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@+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@+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 2d9f8d0e60749..083dec819a0e0 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()); } - -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 subElementTypes; - fooStructTy.walk([&](Type type) { subElementTypes.push_back(type); }); - ASSERT_EQ(subElementTypes.size(), 4U); - - // !llvm.ptr - ASSERT_TRUE(isa(subElementTypes[0])); - - // !llvm.struct<"bar",...> - auto structType = dyn_cast(subElementTypes[1]); - ASSERT_TRUE(bool(structType)); - ASSERT_TRUE(structType.getName().equals("bar")); - - // !llvm.ptr - ASSERT_TRUE(isa(subElementTypes[2])); - - // !llvm.struct<"foo",...> - structType = dyn_cast(subElementTypes[3]); - ASSERT_TRUE(bool(structType)); - ASSERT_TRUE(structType.getName().equals("foo")); -} From a1ee18fa4332f2e25092ef056424bcf43bb63579 Mon Sep 17 00:00:00 2001 From: Christian Ulmann 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 0166fbb647b41..a459000581028 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($_s "LLVM pointer type", "::mlir::LLVM::LLVMPointerType">; // Opaque pointer in a given address space. -class LLVM_OpaquePointerInAddressSpace : Type< - CPred< - "::llvm::cast<::mlir::LLVM::LLVMPointerType>($_self).getAddressSpace() == " - # addressSpace>, +class LLVM_PointerInAddressSpace : 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 cc5c1bd7e1993..b305738fa483c 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 6c6419bf238b4..48b830ae34f29 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 3ce7e9abccaa3..33af73d8a4ac5 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(getType()) + : llvm::isa(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 e5d7d6d8f8c40..81178b2ef901f 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} : () -> () diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir index 1d51796abb03f..d6960f00f0428 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) { + // expected-error@+1 {{op operand #0 must be LLVM pointer type}} + llvm.load %foo : memref -> f32 +} + +// ----- + func.func @load_syncscope(%ptr : !llvm.ptr) { // expected-error@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 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 a459000581028..4e42a0e46d9bf 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($_self)">, "LLVM pointer type", "::mlir::LLVM::LLVMPointerType">; -// Opaque pointer in a given address space. +// Pointer in a given address space. class LLVM_PointerInAddressSpace : 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 fdd1d993fdb8f..04496d6b8f634 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 &results) -> std::optional { // Fastpath for types that won't be converted by this callback anyway. From 62cfdef3dbbbc5546c89dde0f749e7cc8327d796 Mon Sep 17 00:00:00 2001 From: Christian Ulmann 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 ec196e1bd1367..cb8a6e08886df 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().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::min(); - /// Returns the type pointed to by the pointer argument of this GEP. - Type getSourceElementType(); - GEPIndicesAdaptor 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 33af73d8a4ac5..c22cff4c1328a 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 rawConstantIndices; SmallVector 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 ee491a6c558b8..b094c650ff193 100644 --- a/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/TypeConsistency.cpp @@ -164,9 +164,9 @@ static std::optional 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(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.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 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 b305738fa483c..ffe6f25fcd944 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; 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:$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:$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, AttrSizedOperandSegments]>, - Arguments<(ins LLVM_ptr_shared:$dstMem, + Arguments<(ins LLVM_PointerShared:$dstMem, LLVM_AnyPointer:$tmaDescriptor, - LLVM_ptr_shared:$mbar, + LLVM_PointerShared:$mbar, Variadic:$coordinates, PtxPredicate:$predicate)> { let assemblyFormat = [{ @@ -1438,7 +1438,7 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp : [DeclareOpInterfaceMethods, AttrSizedOperandSegments]>, Arguments<(ins LLVM_AnyPointer:$tmaDescriptor, - LLVM_ptr_shared:$srcMem, + LLVM_PointerShared:$srcMem, Variadic:$coordinates, PtxPredicate:$predicate)> { let assemblyFormat = [{