mirror of
https://github.com/capstone-engine/llvm-capstone.git
synced 2024-11-23 13:50:11 +00:00
[MLIR][Doc] Remove LLVM dialect typed pointer documentation (#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
This commit is contained in:
Christian Ulmann
GitHub
parent
3d870434b8
commit
b3eac1ac1e
@ -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
|
||||
```
|
||||
|
||||
|
||||
@ -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
|
||||
|
||||
@ -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
|
||||
|
||||
@ -89,10 +89,10 @@ protected:
|
||||
/// `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 to i64
|
||||
///
|
||||
/// If `sizeInBytes = false`, memref<4x?xf32> emits:
|
||||
/// `sizes[0]` = llvm.mlir.constant(4 : index) : i64
|
||||
|
||||
@ -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
|
||||
}
|
||||
```
|
||||
|
||||
@ -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.ptr
|
||||
}
|
||||
```
|
||||
|
||||
|
||||
@ -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
|
||||
```
|
||||
}];
|
||||
|
||||
@ -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)>
|
||||
```
|
||||
}];
|
||||
|
||||
|
||||
@ -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) ->()
|
||||
```
|
||||
}];
|
||||
|
||||
|
||||
@ -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>}
|
||||
}
|
||||
```
|
||||
}];
|
||||
|
||||
@ -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>
|
||||
}
|
||||
```
|
||||
|
||||
@ -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",...>
|
||||
|
||||
Loading…
Reference in New Issue
Block a user