[Mlir-commits] [mlir] [MLIR][Doc] Remove LLVM dialect typed pointer documentation (PR #71246)
Christian Ulmann
llvmlistbot at llvm.org
Fri Nov 3 15:39:45 PDT 2023
https://github.com/Dinistro created https://github.com/llvm/llvm-project/pull/71246
This commit removes all references to typed pointers. Typed pointers have been deprecated for a while now and they will be removed in a followup.
Related PSA: https://discourse.llvm.org/t/psa-removal-of-typed-pointers-from-the-llvm-dialect/74502
>From c30422c73aa2f4e3868c3fd8d2f052e49ae106b5 Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christianulmann at gmail.com>
Date: Fri, 3 Nov 2023 23:36:37 +0100
Subject: [PATCH] [MLIR][Doc] Remove LLVM dialect typed pointer documentation
This commit removes all references to typed pointers. Typed pointers
have been deprecated for a while now and they will be removed in a
followup.
Related PSA: https://discourse.llvm.org/t/psa-removal-of-typed-pointers-from-the-llvm-dialect/74502
---
mlir/docs/Dialects/LLVM.md | 56 +++-----
mlir/docs/SPIRVToLLVMDialectConversion.md | 18 +--
mlir/docs/TargetLLVMIR.md | 133 ++++++++----------
.../mlir/Conversion/LLVMCommon/Pattern.h | 6 +-
.../mlir/Dialect/LLVMIR/LLVMAttrDefs.td | 10 +-
mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td | 30 ++--
mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td | 6 +-
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 8 +-
.../mlir/Dialect/MemRef/IR/MemRefOps.td | 4 +-
.../mlir/Dialect/OpenACC/OpenACCOps.td | 30 ++--
.../Dialect/SparseTensor/Transforms/Passes.td | 4 +-
.../unittests/Dialect/LLVMIR/LLVMTypeTest.cpp | 4 +-
12 files changed, 134 insertions(+), 175 deletions(-)
diff --git a/mlir/docs/Dialects/LLVM.md b/mlir/docs/Dialects/LLVM.md
index 796c3a7b768443b..5bbccda7cf6bd22 100644
--- a/mlir/docs/Dialects/LLVM.md
+++ b/mlir/docs/Dialects/LLVM.md
@@ -210,9 +210,9 @@ style for types with nested angle brackets and keyword specifiers rather than
using different bracket styles to differentiate types. Types inside the angle
brackets may omit the `!llvm.` prefix for brevity: the parser first attempts to
find a type (starting with `!` or a built-in type) and falls back to accepting a
-keyword. For example, `!llvm.ptr<!llvm.ptr<i32>>` and `!llvm.ptr<ptr<i32>>` are
-equivalent, with the latter being the canonical form, and denote a pointer to a
-pointer to a 32-bit integer.
+keyword. For example, `!llvm.struct<(!llvm.ptr, f32)>` and
+`!llvm.struct<(ptr, f32)>` are equivalent, with the latter being the canonical
+form, and denote a struct containing a pointer and a float.
### Built-in Type Compatibility
@@ -232,8 +232,8 @@ compatibility check.
Each LLVM IR type corresponds to *exactly one* MLIR type, either built-in or
LLVM dialect type. For example, because `i32` is LLVM-compatible, there is no
-`!llvm.i32` type. However, `!llvm.ptr<T>` is defined in the LLVM dialect as
-there is no corresponding built-in type.
+`!llvm.i32` type. However, `!llvm.struct<(T, ...)>` is defined in the LLVM
+dialect as there is no corresponding built-in type.
### Additional Simple Types
@@ -263,24 +263,19 @@ the element type, which can be either compatible built-in or LLVM dialect types.
Pointer types specify an address in memory.
-Both opaque and type-parameterized pointer types are supported.
-[Opaque pointers](https://llvm.org/docs/OpaquePointers.html) do not indicate the
-type of the data pointed to, and are intended to simplify LLVM IR by encoding
-behavior relevant to the pointee type into operations rather than into types.
-Non-opaque pointer types carry the pointee type as a type parameter. Both kinds
-of pointers may be additionally parameterized by an address space. The address
-space is an integer, but this choice may be reconsidered if MLIR implements
-named address spaces. The syntax of pointer types is as follows:
+Pointers are [opaque](https://llvm.org/docs/OpaquePointers.html), i.e., do not
+indicate the type of the data pointed to, and are intended to simplify LLVM IR
+by encoding behavior relevant to the pointee type into operations rather than
+into types. Pointers can optionally be parametrized with an address space. The
+address space is an integer, but this choice may be reconsidered if MLIR
+implements named address spaces. The syntax of pointer types is as follows:
```
llvm-ptr-type ::= `!llvm.ptr` (`<` integer-literal `>`)?
- | `!llvm.ptr<` type (`,` integer-literal)? `>`
```
-where the former case is the opaque pointer type and the latter case is the
-non-opaque pointer type; the optional group containing the integer literal
-corresponds to the memory space. All cases are represented by `LLVMPointerType`
-internally.
+where the optional group containing the integer literal corresponds to the
+address space. All cases are represented by `LLVMPointerType` internally.
#### Array Types
@@ -346,7 +341,7 @@ syntax:
Note that the sets of element types supported by built-in and LLVM dialect
vector types are mutually exclusive, e.g., the built-in vector type does not
-accept `!llvm.ptr<i32>` and the LLVM dialect fixed-width vector type does not
+accept `!llvm.ptr` and the LLVM dialect fixed-width vector type does not
accept `i32`.
The following functions are provided to operate on any kind of the vector types
@@ -367,12 +362,11 @@ compatible with the LLVM dialect:
```mlir
vector<42 x i32> // Vector of 42 32-bit integers.
-!llvm.vec<42 x ptr<i32>> // Vector of 42 pointers to 32-bit integers.
+!llvm.vec<42 x ptr> // Vector of 42 pointers.
!llvm.vec<? x 4 x i32> // Scalable vector of 32-bit integers with
// size divisible by 4.
!llvm.array<2 x vector<2 x i32>> // Array of 2 vectors of 2 32-bit integers.
-!llvm.array<2 x vec<2 x ptr<i32>>> // Array of 2 vectors of 2 pointers to 32-bit
- // integers.
+!llvm.array<2 x vec<2 x ptr>> // Array of 2 vectors of 2 pointers.
```
### Structure Types
@@ -421,21 +415,6 @@ type-or-ref ::= <any compatible type with optional !llvm.>
| `!llvm.`? `struct<` string-literal `>`
```
-The body of the identified struct is printed in full unless the it is
-transitively contained in the same struct. In the latter case, only the
-identifier is printed. For example, the structure containing the pointer to
-itself is represented as `!llvm.struct<"A", (ptr<"A">)>`, and the structure `A`
-containing two pointers to the structure `B` containing a pointer to the
-structure `A` is represented as `!llvm.struct<"A", (ptr<"B", (ptr<"A">)>,
-ptr<"B", (ptr<"A">))>`. Note that the structure `B` is "unrolled" for both
-elements. _A structure with the same name but different body is a syntax error._
-**The user must ensure structure name uniqueness across all modules processed in
-a given MLIR context.** Structure names are arbitrary string literals and may
-include, e.g., spaces and keywords.
-
-Identified structs may be _opaque_. In this case, the body is unknown but the
-structure type is considered _initialized_ and is valid in the IR.
-
#### Literal Structure Types
Literal structures are uniqued according to the list of elements they contain,
@@ -460,11 +439,10 @@ elements provided.
!llvm.struct<packed (i8, i32)> // packed struct
!llvm.struct<"a"> // recursive reference, only allowed within
// another struct, NOT allowed at top level
-!llvm.struct<"a", ptr<struct<"a">>> // supported example of recursive reference
!llvm.struct<"a", ()> // empty, named (necessary to differentiate from
// recursive reference)
!llvm.struct<"a", opaque> // opaque, named
-!llvm.struct<"a", (i32)> // named
+!llvm.struct<"a", (i32, ptr)> // named
!llvm.struct<"a", packed (i8, i32)> // named, packed
```
diff --git a/mlir/docs/SPIRVToLLVMDialectConversion.md b/mlir/docs/SPIRVToLLVMDialectConversion.md
index 7a3bc7c62bd9a2c..0aae02cff26be1b 100644
--- a/mlir/docs/SPIRVToLLVMDialectConversion.md
+++ b/mlir/docs/SPIRVToLLVMDialectConversion.md
@@ -45,7 +45,7 @@ A SPIR-V pointer also takes a Storage Class. At the moment, conversion does
SPIR-V Dialect | LLVM Dialect
:-------------------------------------------: | :-------------------------:
-`!spirv.ptr< <element-type>, <storage-class> >` | `!llvm.ptr<<element-type>>`
+`!spirv.ptr< <element-type>, <storage-class> >` | `!llvm.ptr`
### Array types
@@ -443,7 +443,7 @@ order to go through the pointer.
%i = ...
%var = ...
%0 = llvm.mlir.constant(0 : i32) : i32
-%el = llvm.getelementptr %var[%0, %i, %i] : (!llvm.ptr<struct<packed (f32, array<4 x f32>)>>, i32, i32, i32)
+%el = llvm.getelementptr %var[%0, %i, %i] : (!llvm.ptr, i32, i32, i32), !llvm.struct<packed (f32, array<4 x f32>)>
```
#### `spirv.Load` and `spirv.Store`
@@ -453,16 +453,16 @@ These ops are converted to their LLVM counterparts: `llvm.load` and
following cases, based on the value of the attribute:
* **Aligned**: alignment is passed on to LLVM op builder, for example: `mlir
- // llvm.store %ptr, %val {alignment = 4 : i64} : !llvm.ptr<f32> spirv.Store
+ // llvm.store %ptr, %val {alignment = 4 : i64} : !llvm.ptr spirv.Store
"Function" %ptr, %val ["Aligned", 4] : f32`
* **None**: same case as if there is no memory access attribute.
* **Nontemporal**: set `nontemporal` flag, for example: `mlir // %res =
- llvm.load %ptr {nontemporal} : !llvm.ptr<f32> %res = spirv.Load "Function"
+ llvm.load %ptr {nontemporal} : !llvm.ptr %res = spirv.Load "Function"
%ptr ["Nontemporal"] : f32`
* **Volatile**: mark the op as `volatile`, for example: `mlir // %res =
- llvm.load volatile %ptr : !llvm.ptr<f32> %res = spirv.Load "Function" %ptr
+ llvm.load volatile %ptr : !llvm.ptr f32> %res = spirv.Load "Function" %ptr
["Volatile"] : f32` Otherwise the conversion fails as other cases
(`MakePointerAvailable`, `MakePointerVisible`, `NonPrivatePointer`) are not
supported yet.
@@ -491,7 +491,7 @@ spirv.module Logical GLSL450 {
module {
llvm.mlir.global private @struct() : !llvm.struct<packed (f32, [10 x f32])>
llvm.func @func() {
- %0 = llvm.mlir.addressof @struct : !llvm.ptr<struct<packed (f32, [10 x f32])>>
+ %0 = llvm.mlir.addressof @struct : !llvm.ptr
llvm.return
}
}
@@ -535,13 +535,13 @@ Also, at the moment initialization is only possible via `spirv.Constant`.
```mlir
// Conversion of VariableOp without initialization
%size = llvm.mlir.constant(1 : i32) : i32
-%res = spirv.Variable : !spirv.ptr<vector<3xf32>, Function> => %res = llvm.alloca %size x vector<3xf32> : (i32) -> !llvm.ptr<vec<3 x f32>>
+%res = spirv.Variable : !spirv.ptr<vector<3xf32>, Function> => %res = llvm.alloca %size x vector<3xf32> : (i32) -> !llvm.ptr
// Conversion of VariableOp with initialization
%c = llvm.mlir.constant(0 : i64) : i64
%c = spirv.Constant 0 : i64 %size = llvm.mlir.constant(1 : i32) : i32
-%res = spirv.Variable init(%c) : !spirv.ptr<i64, Function> => %res = llvm.alloca %[[SIZE]] x i64 : (i32) -> !llvm.ptr<i64>
- llvm.store %c, %res : !llvm.ptr<i64>
+%res = spirv.Variable init(%c) : !spirv.ptr<i64, Function> => %res = llvm.alloca %[[SIZE]] x i64 : (i32) -> !llvm.ptr
+ llvm.store %c, %res : i64, !llvm.ptr
```
Note that simple conversion to `alloca` may not be sufficient if the code has
diff --git a/mlir/docs/TargetLLVMIR.md b/mlir/docs/TargetLLVMIR.md
index 4553accc5b9ae2b..73a74c394f2af79 100644
--- a/mlir/docs/TargetLLVMIR.md
+++ b/mlir/docs/TargetLLVMIR.md
@@ -135,20 +135,19 @@ Examples:
```mlir
// Assuming index is converted to i64.
-memref<f32> -> !llvm.struct<(ptr<f32> , ptr<f32>, i64)>
-memref<1 x f32> -> !llvm.struct<(ptr<f32>, ptr<f32>, i64,
+memref<f32> -> !llvm.struct<(ptr , ptr, i64)>
+memref<1 x f32> -> !llvm.struct<(ptr, ptr, i64,
array<1 x i64>, array<1 x i64>)>
-memref<? x f32> -> !llvm.struct<(ptr<f32>, ptr<f32>, i64
+memref<? x f32> -> !llvm.struct<(ptr, ptr, i64
array<1 x i64>, array<1 x i64>)>
-memref<10x42x42x43x123 x f32> -> !llvm.struct<(ptr<f32>, ptr<f32>, i64
+memref<10x42x42x43x123 x f32> -> !llvm.struct<(ptr, ptr, i64
array<5 x i64>, array<5 x i64>)>
-memref<10x?x42x?x123 x f32> -> !llvm.struct<(ptr<f32>, ptr<f32>, i64
+memref<10x?x42x?x123 x f32> -> !llvm.struct<(ptr, ptr, i64
array<5 x i64>, array<5 x i64>)>
// Memref types can have vectors as element types
-memref<1x? x vector<4xf32>> -> !llvm.struct<(ptr<vector<4 x f32>>,
- ptr<vector<4 x f32>>, i64,
- array<2 x i64>, array<2 x i64>)>
+memref<1x? x vector<4xf32>> -> !llvm.struct<(ptr, ptr, i64, array<2 x i64>,
+ array<2 x i64>)>
```
#### Unranked MemRef Types
@@ -159,7 +158,7 @@ as *unranked descriptor*. It contains:
1. a converted `index`-typed integer representing the dynamic rank of the
memref;
-2. a type-erased pointer (`!llvm.ptr<i8>`) to a ranked memref descriptor with
+2. a type-erased pointer (`!llvm.ptr`) to a ranked memref descriptor with
the contents listed above.
This descriptor is primarily intended for interfacing with rank-polymorphic
@@ -219,49 +218,42 @@ Examples:
// Function-typed arguments or results in higher-order functions:
(() -> ()) -> (() -> ())
-// are converted into pointers to functions.
-!llvm.func<ptr<func<void ()>> (ptr<func<void ()>>)>
-
-// These rules apply recursively: a function type taking a function that takes
-// another function
-( ( (i32) -> (i64) ) -> () ) -> ()
-// is converted into a function type taking a pointer-to-function that takes
-// another point-to-function.
-!llvm.func<void (ptr<func<void (ptr<func<i64 (i32)>>)>>)>
+// are converted into opaque pointers.
+!llvm.func<ptr (ptr)>
// A memref descriptor appearing as function argument:
(memref<f32>) -> ()
// gets converted into a list of individual scalar components of a descriptor.
-!llvm.func<void (ptr<f32>, ptr<f32>, i64)>
+!llvm.func<void (ptr, ptr, i64)>
// The list of arguments is linearized and one can freely mix memref and other
// types in this list:
(memref<f32>, f32) -> ()
// which gets converted into a flat list.
-!llvm.func<void (ptr<f32>, ptr<f32>, i64, f32)>
+!llvm.func<void (ptr, ptr, i64, f32)>
// For nD ranked memref descriptors:
(memref<?x?xf32>) -> ()
// the converted signature will contain 2n+1 `index`-typed integer arguments,
// offset, n sizes and n strides, per memref argument type.
-!llvm.func<void (ptr<f32>, ptr<f32>, i64, i64, i64, i64, i64)>
+!llvm.func<void (ptr, ptr, i64, i64, i64, i64, i64)>
// Same rules apply to unranked descriptors:
(memref<*xf32>) -> ()
// which get converted into their components.
-!llvm.func<void (i64, ptr<i8>)>
+!llvm.func<void (i64, ptr)>
// However, returning a memref from a function is not affected:
() -> (memref<?xf32>)
// gets converted to a function returning a descriptor structure.
-!llvm.func<struct<(ptr<f32>, ptr<f32>, i64, array<1xi64>, array<1xi64>)> ()>
+!llvm.func<struct<(ptr, ptr, i64, array<1xi64>, array<1xi64>)> ()>
// If multiple memref-typed results are returned:
() -> (memref<f32>, memref<f64>)
// their descriptor structures are additionally packed into another structure,
// potentially with other non-memref typed results.
-!llvm.func<struct<(struct<(ptr<f32>, ptr<f32>, i64)>,
- struct<(ptr<double>, ptr<double>, i64)>)> ()>
+!llvm.func<struct<(struct<(ptr, ptr, i64)>,
+ struct<(ptr, ptr, i64)>)> ()>
// If "func.varargs" attribute is set:
(i32) -> () attributes { "func.varargs" = true }
@@ -290,8 +282,7 @@ vector<4x8 x f32>
memref<2 x vector<4x8 x f32>
// ->
-!llvm.struct<(ptr<array<4 x vector<8xf32>>>, ptr<array<4 x vector<8xf32>>>
- i64, array<1 x i64>, array<1 x i64>)>
+!llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
```
#### Tensor Types
@@ -382,10 +373,10 @@ func.func @foo(%arg0: memref<?xf32>) -> () {
// Gets converted to the following
// (using type alias for brevity):
-!llvm.memref_1d = !llvm.struct<(ptr<f32>, ptr<f32>, i64, array<1xi64>, array<1xi64>)>
+!llvm.memref_1d = !llvm.struct<(ptr, ptr, i64, array<1xi64>, array<1xi64>)>
-llvm.func @foo(%arg0: !llvm.ptr<f32>, // Allocated pointer.
- %arg1: !llvm.ptr<f32>, // Aligned pointer.
+llvm.func @foo(%arg0: !llvm.ptr, // Allocated pointer.
+ %arg1: !llvm.ptr, // Aligned pointer.
%arg2: i64, // Offset.
%arg3: i64, // Size in dim 0.
%arg4: i64) { // Stride in dim 0.
@@ -412,7 +403,7 @@ func.func @bar() {
// Gets converted to the following
// (using type alias for brevity):
-!llvm.memref_1d = !llvm.struct<(ptr<f32>, ptr<f32>, i64, array<1xi64>, array<1xi64>)>
+!llvm.memref_1d = !llvm.struct<(ptr, ptr, i64, array<1xi64>, array<1xi64>)>
llvm.func @bar() {
%0 = "get"() : () -> !llvm.memref_1d
@@ -434,7 +425,7 @@ llvm.func @bar() {
For unranked memrefs, the list of function arguments always contains two
elements, same as the unranked memref descriptor: an integer rank, and a
-type-erased (`!llvm<"i8*">`) pointer to the ranked memref descriptor. Note that
+type-erased (`!llvm.ptr`) pointer to the ranked memref descriptor. Note that
while the *calling convention* does not require allocation, *casting* to
unranked memref does since one cannot take an address of an SSA value containing
the ranked memref, which must be stored in some memory instead. The caller is in
@@ -452,13 +443,13 @@ llvm.func @foo(%arg0: memref<*xf32>) -> () {
// Gets converted to the following.
llvm.func @foo(%arg0: i64 // Rank.
- %arg1: !llvm.ptr<i8>) { // Type-erased pointer to descriptor.
+ %arg1: !llvm.ptr) { // Type-erased pointer to descriptor.
// Pack the unranked memref descriptor.
- %0 = llvm.mlir.undef : !llvm.struct<(i64, ptr<i8>)>
- %1 = llvm.insertvalue %arg0, %0[0] : !llvm.struct<(i64, ptr<i8>)>
- %2 = llvm.insertvalue %arg1, %1[1] : !llvm.struct<(i64, ptr<i8>)>
+ %0 = llvm.mlir.undef : !llvm.struct<(i64, ptr)>
+ %1 = llvm.insertvalue %arg0, %0[0] : !llvm.struct<(i64, ptr)>
+ %2 = llvm.insertvalue %arg1, %1[1] : !llvm.struct<(i64, ptr)>
- "use"(%2) : (!llvm.struct<(i64, ptr<i8>)>) -> ()
+ "use"(%2) : (!llvm.struct<(i64, ptr)>) -> ()
llvm.return
}
```
@@ -473,14 +464,14 @@ llvm.func @bar() {
// Gets converted to the following.
llvm.func @bar() {
- %0 = "get"() : () -> (!llvm.struct<(i64, ptr<i8>)>)
+ %0 = "get"() : () -> (!llvm.struct<(i64, ptr)>)
// Unpack the memref descriptor.
- %1 = llvm.extractvalue %0[0] : !llvm.struct<(i64, ptr<i8>)>
- %2 = llvm.extractvalue %0[1] : !llvm.struct<(i64, ptr<i8>)>
+ %1 = llvm.extractvalue %0[0] : !llvm.struct<(i64, ptr)>
+ %2 = llvm.extractvalue %0[1] : !llvm.struct<(i64, ptr)>
// Pass individual values to the callee.
- llvm.call @foo(%1, %2) : (i64, !llvm.ptr<i8>)
+ llvm.call @foo(%1, %2) : (i64, !llvm.ptr)
llvm.return
}
```
@@ -524,12 +515,12 @@ func.func @caller(%0 : memref<2x4xf32>) {
// ->
-!descriptor = !llvm.struct<(ptr<f32>, ptr<f32>, i64,
+!descriptor = !llvm.struct<(ptr, ptr, i64,
array<2xi64>, array<2xi64>)>
-llvm.func @callee(!llvm.ptr<f32>)
+llvm.func @callee(!llvm.ptr)
-llvm.func @caller(%arg0: !llvm.ptr<f32>) {
+llvm.func @caller(%arg0: !llvm.ptr) {
// A descriptor value is defined at the function entry point.
%0 = llvm.mlir.undef : !descriptor
@@ -552,7 +543,7 @@ llvm.func @caller(%arg0: !llvm.ptr<f32>) {
// The function call corresponds to extracting the aligned data pointer.
%12 = llvm.extractelement %11[1] : !descriptor
- llvm.call @callee(%12) : (!llvm.ptr<f32>) -> ()
+ llvm.call @callee(%12) : (!llvm.ptr) -> ()
}
```
@@ -644,10 +635,10 @@ func.func @qux(%arg0: memref<?x?xf32>)
// Gets converted into the following
// (using type alias for brevity):
-!llvm.memref_2d = !llvm.struct<(ptr<f32>, ptr<f32>, i64, array<2xi64>, array<2xi64>)>
+!llvm.memref_2d = !llvm.struct<(ptr, ptr, i64, array<2xi64>, array<2xi64>)>
// Function with unpacked arguments.
-llvm.func @qux(%arg0: !llvm.ptr<f32>, %arg1: !llvm.ptr<f32>,
+llvm.func @qux(%arg0: !llvm.ptr, %arg1: !llvm.ptr,
%arg2: i64, %arg3: i64, %arg4: i64,
%arg5: i64, %arg6: i64) {
// Populate memref descriptor (as per calling convention).
@@ -663,23 +654,18 @@ llvm.func @qux(%arg0: !llvm.ptr<f32>, %arg1: !llvm.ptr<f32>,
// Store the descriptor in a stack-allocated space.
%8 = llvm.mlir.constant(1 : index) : i64
%9 = llvm.alloca %8 x !llvm.memref_2d
- : (i64) -> !llvm.ptr<struct<(ptr<f32>, ptr<f32>, i64,
- array<2xi64>, array<2xi64>)>>
- llvm.store %7, %9 : !llvm.ptr<struct<(ptr<f32>, ptr<f32>, i64,
- array<2xi64>, array<2xi64>)>>
+ : (i64) -> !llvm.ptr
+ llvm.store %7, %9 : !llvm.memref_2d, !llvm.ptr
// Call the interface function.
- llvm.call @_mlir_ciface_qux(%9)
- : (!llvm.ptr<struct<(ptr<f32>, ptr<f32>, i64,
- array<2xi64>, array<2xi64>)>>) -> ()
+ llvm.call @_mlir_ciface_qux(%9) : (!llvm.ptr) -> ()
// The stored descriptor will be freed on return.
llvm.return
}
// Interface function.
-llvm.func @_mlir_ciface_qux(!llvm.ptr<struct<(ptr<f32>, ptr<f32>, i64,
- array<2xi64>, array<2xi64>)>>)
+llvm.func @_mlir_ciface_qux(!llvm.ptr)
```
```mlir
@@ -689,20 +675,19 @@ func.func @foo(%arg0: memref<?x?xf32>) {
// Gets converted into the following
// (using type alias for brevity):
-!llvm.memref_2d = !llvm.struct<(ptr<f32>, ptr<f32>, i64, array<2xi64>, array<2xi64>)>
-!llvm.memref_2d_ptr = !llvm.ptr<struct<(ptr<f32>, ptr<f32>, i64, array<2xi64>, array<2xi64>)>>
+!llvm.memref_2d = !llvm.struct<(ptr, ptr, i64, array<2xi64>, array<2xi64>)>
// Function with unpacked arguments.
-llvm.func @foo(%arg0: !llvm.ptr<f32>, %arg1: !llvm.ptr<f32>,
+llvm.func @foo(%arg0: !llvm.ptr, %arg1: !llvm.ptr,
%arg2: i64, %arg3: i64, %arg4: i64,
%arg5: i64, %arg6: i64) {
llvm.return
}
// Interface function callable from C.
-llvm.func @_mlir_ciface_foo(%arg0: !llvm.memref_2d_ptr) {
+llvm.func @_mlir_ciface_foo(%arg0: !llvm.ptr) {
// Load the descriptor.
- %0 = llvm.load %arg0 : !llvm.memref_2d_ptr
+ %0 = llvm.load %arg0 : !llvm.ptr -> !llvm.memref_2d
// Unpack the descriptor as per calling convention.
%1 = llvm.extractvalue %0[0] : !llvm.memref_2d
@@ -713,7 +698,7 @@ llvm.func @_mlir_ciface_foo(%arg0: !llvm.memref_2d_ptr) {
%6 = llvm.extractvalue %0[4, 0] : !llvm.memref_2d
%7 = llvm.extractvalue %0[4, 1] : !llvm.memref_2d
llvm.call @foo(%1, %2, %3, %4, %5, %6, %7)
- : (!llvm.ptr<f32>, !llvm.ptr<f32>, i64, i64, i64,
+ : (!llvm.ptr, !llvm.ptr, i64, i64, i64,
i64, i64) -> ()
llvm.return
}
@@ -726,11 +711,10 @@ func.func @foo(%arg0: memref<?x?xf32>) -> memref<?x?xf32> {
// Gets converted into the following
// (using type alias for brevity):
-!llvm.memref_2d = !llvm.struct<(ptr<f32>, ptr<f32>, i64, array<2xi64>, array<2xi64>)>
-!llvm.memref_2d_ptr = !llvm.ptr<struct<(ptr<f32>, ptr<f32>, i64, array<2xi64>, array<2xi64>)>>
+!llvm.memref_2d = !llvm.struct<(ptr, ptr, i64, array<2xi64>, array<2xi64>)>
// Function with unpacked arguments.
-llvm.func @foo(%arg0: !llvm.ptr<f32>, %arg1: !llvm.ptr<f32>, %arg2: i64,
+llvm.func @foo(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: i64,
%arg3: i64, %arg4: i64, %arg5: i64, %arg6: i64)
-> !llvm.memref_2d {
%0 = llvm.mlir.undef : !llvm.memref_2d
@@ -745,8 +729,8 @@ llvm.func @foo(%arg0: !llvm.ptr<f32>, %arg1: !llvm.ptr<f32>, %arg2: i64,
}
// Interface function callable from C.
-llvm.func @_mlir_ciface_foo(%arg0: !llvm.memref_2d_ptr, %arg1: !llvm.memref_2d_ptr) {
- %0 = llvm.load %arg1 : !llvm.memref_2d_ptr
+llvm.func @_mlir_ciface_foo(%arg0: !llvm.ptr, %arg1: !llvm.ptr) {
+ %0 = llvm.load %arg1 : !llvm.ptr
%1 = llvm.extractvalue %0[0] : !llvm.memref_2d
%2 = llvm.extractvalue %0[1] : !llvm.memref_2d
%3 = llvm.extractvalue %0[2] : !llvm.memref_2d
@@ -755,8 +739,8 @@ llvm.func @_mlir_ciface_foo(%arg0: !llvm.memref_2d_ptr, %arg1: !llvm.memref_2d_p
%6 = llvm.extractvalue %0[4, 0] : !llvm.memref_2d
%7 = llvm.extractvalue %0[4, 1] : !llvm.memref_2d
%8 = llvm.call @foo(%1, %2, %3, %4, %5, %6, %7)
- : (!llvm.ptr<f32>, !llvm.ptr<f32>, i64, i64, i64, i64, i64) -> !llvm.memref_2d
- llvm.store %8, %arg0 : !llvm.memref_2d_ptr
+ : (!llvm.ptr, !llvm.ptr, i64, i64, i64, i64, i64) -> !llvm.memref_2d
+ llvm.store %8, %arg0 : !llvm.memref_2d, !llvm.ptr
llvm.return
}
```
@@ -809,7 +793,7 @@ is transformed into the equivalent of the following code:
// Compute the linearized index from strides.
// When strides or, in absence of explicit strides, the corresponding sizes are
// dynamic, extract the stride value from the descriptor.
-%stride1 = llvm.extractvalue[4, 0] : !llvm.struct<(ptr<f32>, ptr<f32>, i64,
+%stride1 = llvm.extractvalue[4, 0] : !llvm.struct<(ptr, ptr, i64,
array<4xi64>, array<4xi64>)>
%addr1 = arith.muli %stride1, %1 : i64
@@ -829,21 +813,20 @@ is transformed into the equivalent of the following code:
// If the linear offset is known to be zero, it can also be omitted. If it is
// dynamic, it is extracted from the descriptor.
-%offset = llvm.extractvalue[2] : !llvm.struct<(ptr<f32>, ptr<f32>, i64,
+%offset = llvm.extractvalue[2] : !llvm.struct<(ptr, ptr, i64,
array<4xi64>, array<4xi64>)>
%addr7 = arith.addi %addr6, %offset : i64
// All accesses are based on the aligned pointer.
-%aligned = llvm.extractvalue[1] : !llvm.struct<(ptr<f32>, ptr<f32>, i64,
+%aligned = llvm.extractvalue[1] : !llvm.struct<(ptr, ptr, i64,
array<4xi64>, array<4xi64>)>
// Get the address of the data pointer.
%ptr = llvm.getelementptr %aligned[%addr7]
- : !llvm.struct<(ptr<f32>, ptr<f32>, i64, array<4xi64>, array<4xi64>)>
- -> !llvm.ptr<f32>
+ : !llvm.struct<(ptr, ptr, i64, array<4xi64>, array<4xi64>)> -> !llvm.ptr
// Perform the actual load.
-%0 = llvm.load %ptr : !llvm.ptr<f32>
+%0 = llvm.load %ptr : !llvm.ptr -> f32
```
For stores, the address computation code is identical and only the actual store
diff --git a/mlir/include/mlir/Conversion/LLVMCommon/Pattern.h b/mlir/include/mlir/Conversion/LLVMCommon/Pattern.h
index aea6c38a441d3a5..df9d9884cd6983b 100644
--- a/mlir/include/mlir/Conversion/LLVMCommon/Pattern.h
+++ b/mlir/include/mlir/Conversion/LLVMCommon/Pattern.h
@@ -89,10 +89,10 @@ class ConvertToLLVMPattern : public ConversionPattern {
/// `strides[1]` = llvm.mlir.constant(1 : index) : i64
/// `strides[0]` = `sizes[0]`
/// %size = llvm.mul `sizes[0]`, `sizes[1]` : i64
- /// %nullptr = llvm.mlir.zero : !llvm.ptr<f32>
+ /// %nullptr = llvm.mlir.zero : !llvm.ptr
/// %gep = llvm.getelementptr %nullptr[%size]
- /// : (!llvm.ptr<f32>, i64) -> !llvm.ptr<f32>
- /// `sizeBytes` = llvm.ptrtoint %gep : !llvm.ptr<f32> to i64
+ /// : (!llvm.ptr, i64) -> !llvm.ptr, f32
+ /// `sizeBytes` = llvm.ptrtoint %gep : !llvm.ptr, f32
///
/// If `sizeInBytes = false`, memref<4x?xf32> emits:
/// `sizes[0]` = llvm.mlir.constant(4 : index) : i64
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMAttrDefs.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMAttrDefs.td
index 4a9bb5355d1c7ea..0d9af88157a8ae1 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMAttrDefs.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMAttrDefs.td
@@ -643,14 +643,14 @@ def LLVM_AliasScopeAttr : LLVM_Attr<"AliasScope", "alias_scope"> {
#domain = #llvm.alias_scope_domain<id = distinct[1]<>, description = "Optional domain description">
#scope1 = #llvm.alias_scope<id = distinct[2]<>, domain = #domain>
#scope2 = #llvm.alias_scope<id = distinct[3]<>, domain = #domain, description = "Optional scope description">
- llvm.func @foo(%ptr1 : !llvm.ptr<i32>) {
+ llvm.func @foo(%ptr1 : !llvm.ptr) {
%c0 = llvm.mlir.constant(0 : i32) : i32
%c4 = llvm.mlir.constant(4 : i32) : i32
- %1 = llvm.ptrtoint %ptr1 : !llvm.ptr<i32> to i32
+ %1 = llvm.ptrtoint %ptr1 : !llvm.ptr to i32
%2 = llvm.add %1, %c1 : i32
- %ptr2 = llvm.inttoptr %2 : i32 to !llvm.ptr<i32>
- llvm.store %c0, %ptr1 { alias_scopes = [#scope1], llvm.noalias = [#scope2] } : !llvm.ptr<i32>
- llvm.store %c4, %ptr2 { alias_scopes = [#scope2], llvm.noalias = [#scope1] } : !llvm.ptr<i32>
+ %ptr2 = llvm.inttoptr %2 : i32 to !llvm.ptr
+ llvm.store %c0, %ptr1 { alias_scopes = [#scope1], llvm.noalias = [#scope2] } : i32, !llvm.ptr
+ llvm.store %c4, %ptr2 { alias_scopes = [#scope2], llvm.noalias = [#scope1] } : i32, !llvm.ptr
llvm.return
}
```
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
index 638c31b39682ea6..77ab1e4efaa0a91 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td
@@ -267,14 +267,14 @@ def LLVM_GEPOp : LLVM_Op<"getelementptr", [Pure,
```mlir
// GEP with an SSA value offset
- %0 = llvm.getelementptr %1[%2] : (!llvm.ptr<f32>, i64) -> !llvm.ptr<f32>
+ %0 = llvm.getelementptr %1[%2] : (!llvm.ptr, i64) -> !llvm.ptr, f32
// GEP with a constant offset and the inbounds attribute set
- %0 = llvm.getelementptr inbounds %1[3] : (!llvm.ptr<f32>) -> !llvm.ptr<f32>
+ %0 = llvm.getelementptr inbounds %1[3] : (!llvm.ptr) -> !llvm.ptr, f32
// GEP with constant offsets into a structure
%0 = llvm.getelementptr %1[0, 1]
- : (!llvm.ptr<struct(i32, f32)>) -> !llvm.ptr<f32>
+ : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(i32, f32)>
```
}];
@@ -1053,16 +1053,16 @@ def LLVM_AddressOfOp : LLVM_Op<"mlir.addressof",
```mlir
func @foo() {
// Get the address of a global variable.
- %0 = llvm.mlir.addressof @const : !llvm.ptr<i32>
+ %0 = llvm.mlir.addressof @const : !llvm.ptr
// Use it as a regular pointer.
- %1 = llvm.load %0 : !llvm.ptr<i32>
+ %1 = llvm.load %0 : !llvm.ptr -> i32
// Get the address of a function.
- %2 = llvm.mlir.addressof @foo : !llvm.ptr<func<void ()>>
+ %2 = llvm.mlir.addressof @foo : !llvm.ptr
// The function address can be used for indirect calls.
- llvm.call %2() : () -> ()
+ llvm.call %2() : !llvm.ptr, () -> ()
}
// Define the global.
@@ -1141,13 +1141,13 @@ def LLVM_GlobalOp : LLVM_Op<"mlir.global",
```mlir
// This global is initialized with the equivalent of:
// i32* getelementptr (i32* @g2, i32 2)
- llvm.mlir.global constant @int_gep() : !llvm.ptr<i32> {
- %0 = llvm.mlir.addressof @g2 : !llvm.ptr<i32>
+ llvm.mlir.global constant @int_gep() : !llvm.ptr {
+ %0 = llvm.mlir.addressof @g2 : !llvm.ptr
%1 = llvm.mlir.constant(2 : i32) : i32
%2 = llvm.getelementptr %0[%1]
- : (!llvm.ptr<i32>, i32) -> !llvm.ptr<i32>
+ : (!llvm.ptr, i32) -> !llvm.ptr, i32
// The initializer region must end with `llvm.return`.
- llvm.return %2 : !llvm.ptr<i32>
+ llvm.return %2 : !llvm.ptr
}
```
@@ -1174,12 +1174,12 @@ def LLVM_GlobalOp : LLVM_Op<"mlir.global",
llvm.mlir.global constant @no_trailing_type("foo bar")
// A complex initializer is constructed with an initializer region.
- llvm.mlir.global constant @int_gep() : !llvm.ptr<i32> {
- %0 = llvm.mlir.addressof @g2 : !llvm.ptr<i32>
+ llvm.mlir.global constant @int_gep() : !llvm.ptr {
+ %0 = llvm.mlir.addressof @g2 : !llvm.ptr
%1 = llvm.mlir.constant(2 : i32) : i32
%2 = llvm.getelementptr %0[%1]
- : (!llvm.ptr<i32>, i32) -> !llvm.ptr<i32>
- llvm.return %2 : !llvm.ptr<i32>
+ : (!llvm.ptr, i32) -> !llvm.ptr, i32
+ llvm.return %2 : !llvm.ptri32>
}
```
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
index 5f746e21e831b59..e31029bfed5a54a 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.td
@@ -127,14 +127,12 @@ def LLVMPointerType : LLVMType<"LLVMPointer", "ptr", [
let summary = "LLVM pointer type";
let description = [{
The `!llvm.ptr` type is an LLVM pointer type. This type typically represents
- a reference to an object in memory. Pointers may be opaque or parameterized
- by the element type. Both opaque and non-opaque pointers are additionally
- parameterized by the address space.
+ a reference to an object in memory. Pointers are optionally parameterized
+ by the address space.
Example:
```mlir
- !llvm.ptr<i8>
!llvm.ptr
```
}];
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 6947cf10e3600d4..c49decde1638b1c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -1066,7 +1066,7 @@ def NVVM_WMMALoadOp: NVVM_Op<"wmma.load">,
```mlir
%2 = nvvm.wmma.load %0, %1
{eltype = "f16", frag = "a", k = 16 : i32, layout = "row", m = 16 : i32, n = 16 : i32}
- : (!llvm.ptr<i32, 3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+ : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
```
}];
@@ -1121,7 +1121,7 @@ def NVVM_WMMAStoreOp : NVVM_Op<"wmma.store">,
```mlir
nvvm.wmma.store %0, %1, %2, %3, %4, %5
{eltype = "f16", k = 16 : i32, layout = "row", m = 16 : i32, n = 16 : i32}
- : !llvm.ptr<i32, 3>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>
+ : !llvm.ptr<3>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>
```
}];
@@ -1249,9 +1249,9 @@ def NVVM_LdMatrixOp: NVVM_Op<"ldmatrix">,
Example:
```mlir
%l1 = nvvm.ldmatrix %ptr {num = 1 : i32, layout = #nvvm.mma_layout<row>} :
- (!llvm.ptr<i32, 3>) -> i32
+ (!llvm.ptr<3>) -> i32
%l2 = nvvm.ldmatrix %ptr {num = 4 : i32, layout = #nvvm.mma_layout<row>} :
- (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
+ (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>
```
}];
diff --git a/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td b/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td
index 8fa41f4e4b659f5..1c69980dfedd856 100644
--- a/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td
+++ b/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td
@@ -892,8 +892,8 @@ def MemRef_ExtractAlignedPointerAsIndexOp :
```
%0 = memref.extract_aligned_pointer_as_index %arg : memref<4x4xf32> -> index
%1 = arith.index_cast %0 : index to i64
- %2 = llvm.inttoptr %1 : i64 to !llvm.ptr<f32>
- call @foo(%2) : (!llvm.ptr<f32>) ->()
+ %2 = llvm.inttoptr %1 : i64 to !llvm.ptr
+ call @foo(%2) : (!llvm.ptr) ->()
```
}];
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index 17cb9d6f1c2cb63..d0b52a0b4024172 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -632,7 +632,7 @@ def OpenACC_FirstprivateRecipeOp : OpenACC_Op<"firstprivate.recipe",
// init region contains a sequence of operations to create and
// initialize the copy if needed. It yields the create copy.
} copy {
- ^bb0(%0: f32, %1: !llvm.ptr<f32>):
+ ^bb0(%0: f32, %1: !llvm.ptr):
// copy region contains a sequence of operations to copy the initial value
// of the firstprivate value to the newly created value.
} destroy {
@@ -1131,8 +1131,8 @@ def OpenACC_HostDataOp : OpenACC_Op<"host_data", [AttrSizedOperandSegments]> {
Example:
```mlir
- %0 = acc.use_device varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.host_data dataOperands(%0 : !llvm.ptr<f32>) {
+ %0 = acc.use_device varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.host_data dataOperands(%0 : !llvm.ptr) {
}
```
@@ -1424,8 +1424,8 @@ def OpenACC_DeclareEnterOp : OpenACC_Op<"declare_enter", []> {
Example showing `acc declare create(a)`:
```mlir
- %0 = acc.create varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.declare_enter dataOperands(%0 : !llvm.ptr<f32>)
+ %0 = acc.create varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.declare_enter dataOperands(%0 : !llvm.ptr)
```
}];
@@ -1452,9 +1452,9 @@ def OpenACC_DeclareExitOp : OpenACC_Op<"declare_exit", []> {
Example showing `acc declare device_resident(a)`:
```mlir
- %0 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32> {dataClause = #acc<data_clause declare_device_resident>}
- acc.declare_exit dataOperands(%0 : !llvm.ptr<f32>)
- acc.delete accPtr(%0 : !llvm.ptr<f32>) {dataClause = #acc<data_clause declare_device_resident>}
+ %0 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause declare_device_resident>}
+ acc.declare_exit dataOperands(%0 : !llvm.ptr)
+ acc.delete accPtr(%0 : !llvm.ptr) {dataClause = #acc<data_clause declare_device_resident>}
```
}];
@@ -1487,9 +1487,9 @@ def OpenACC_GlobalConstructorOp : OpenACC_Op<"global_ctor",
llvm.return %0 : i32
}
acc.global_ctor @acc_constructor {
- %0 = llvm.mlir.addressof @globalvar : !llvm.ptr<i32>
- %1 = acc.create varPtr(%0 : !llvm.ptr<i32>) -> !llvm.ptr<i32>
- acc.declare_enter dataOperands(%1 : !llvm.ptr<i32>)
+ %0 = llvm.mlir.addressof @globalvar : !llvm.ptr
+ %1 = acc.create varPtr(%0 : !llvm.ptr) -> !llvm.ptr
+ acc.declare_enter dataOperands(%1 : !llvm.ptr)
}
```
}];
@@ -1521,10 +1521,10 @@ def OpenACC_GlobalDestructorOp : OpenACC_Op<"global_dtor",
llvm.return %0 : i32
}
acc.global_dtor @acc_destructor {
- %0 = llvm.mlir.addressof @globalvar : !llvm.ptr<i32>
- %1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr<i32>) -> !llvm.ptr<i32> {dataClause = #acc<data_clause create>}
- acc.declare_exit dataOperands(%1 : !llvm.ptr<i32>)
- acc.delete accPtr(%1 : !llvm.ptr<i32>) {dataClause = #acc<data_clause create>}
+ %0 = llvm.mlir.addressof @globalvar : !llvm.ptr
+ %1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause create>}
+ acc.declare_exit dataOperands(%1 : !llvm.ptr)
+ acc.delete accPtr(%1 : !llvm.ptr) {dataClause = #acc<data_clause create>}
}
```
}];
diff --git a/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.td b/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.td
index 995e842289035b1..485b44a1d6a86c5 100644
--- a/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.td
@@ -227,10 +227,10 @@ def SparseTensorConversionPass : Pass<"sparse-tensor-conversion", "ModuleOp"> {
}
After:
- func.func @foo(%arg0: !llvm.ptr<i8>) -> memref<?xindex> {
+ func.func @foo(%arg0: !llvm.ptr) -> memref<?xindex> {
%c1 = arith.constant 1 : index
%0 = call @sparsePointers0(%arg0, %c1)
- : (!llvm.ptr<i8>, index) -> memref<?xindex>
+ : (!llvm.ptr, index) -> memref<?xindex>
return %0 : memref<?xindex>
}
```
diff --git a/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp b/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp
index aa19b5c651f5594..2d9f8d0e6074990 100644
--- a/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp
+++ b/mlir/unittests/Dialect/LLVMIR/LLVMTypeTest.cpp
@@ -35,7 +35,7 @@ TEST_F(LLVMIRTest, MutualReferencedSubElementTypes) {
fooStructTy.walk([&](Type type) { subElementTypes.push_back(type); });
ASSERT_EQ(subElementTypes.size(), 4U);
- // !llvm.ptr<struct<"foo",...>>
+ // !llvm.ptr
ASSERT_TRUE(isa<LLVMPointerType>(subElementTypes[0]));
// !llvm.struct<"bar",...>
@@ -43,7 +43,7 @@ TEST_F(LLVMIRTest, MutualReferencedSubElementTypes) {
ASSERT_TRUE(bool(structType));
ASSERT_TRUE(structType.getName().equals("bar"));
- // !llvm.ptr<struct<"bar",...>>
+ // !llvm.ptr
ASSERT_TRUE(isa<LLVMPointerType>(subElementTypes[2]));
// !llvm.struct<"foo",...>
More information about the Mlir-commits
mailing list