[Mlir-commits] [mlir] 482c1df - [mlir][openacc] Add descriptions for new acc data operations

[Razvan Lupusoru]

Author: Razvan Lupusoru
Date: 2023-04-21T10:21:29-07:00
New Revision: 482c1dfed314df844c3de68f02bb7c5708dea4bf

URL: https://github.com/llvm/llvm-project/commit/482c1dfed314df844c3de68f02bb7c5708dea4bf
DIFF: https://github.com/llvm/llvm-project/commit/482c1dfed314df844c3de68f02bb7c5708dea4bf.diff

LOG: [mlir][openacc] Add descriptions for new acc data operations

The acc.bounds operation now clarifies that it is meant to be
zero-based. I added example for both C++ and Fortran to
show the normalized representation.

I also added description for each operand in the data entry
and data exit operations.

Reviewed By: clementval

Differential Revision: https://reviews.llvm.org/D148860

Added: 
    

Modified: 
    mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td

Removed: 
    


################################################################################
diff  --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index 535e10ee13c02..36e8b3a5b5782 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -101,7 +101,43 @@ def OpenACC_DataClauseEnum : I64EnumAttr<"DataClause",
 // Either (or both) extent and upperbound must be specified.
 def OpenACC_DataBoundsOp : OpenACC_Op<"bounds",
     [AttrSizedOperandSegments, NoMemoryEffect]> {
-  let summary = "Represents bounds information for acc data clause.";
+  let summary = "Represents normalized bounds information for acc data clause.";
+
+  let description = [{
+    This operation is used to record bounds used in acc data clause in a
+    normalized fashion (zero-based). This works well with the `PointerLikeType`
+    requirement in data clauses - since a `lowerbound` of 0 means looking
+    at data at the zero offset from pointer.
+
+    The operation must have an `upperbound` or `extent` (or both are allowed
+    - but not checked for consistency). When the source language's arrays are
+    not zero-based, the `startIdx` must specify the zero-position index.
+
+    Examples below show copying a slice of 10-element array except first element.
+    Note that the examples use extent in data clause for C++ and upperbound
+    for Fortran (as per 2.7.1). To simplify examples, the constants are used
+    directly in the acc.bounds operands - this is not the syntax of operation.
+
+    C++:
+    ```
+    int array[10];
+    #pragma acc copy(array[1:9])
+    ```
+    =>
+    ```mlir
+    acc.bounds lb(1) ub(9) extent(9) startIdx(0)
+    ```
+
+    Fortran:
+    ```
+    integer :: array(1:10)
+    !$acc copy(array(2:10))
+    ```
+    =>
+    ```mlir
+    acc.bounds lb(1) ub(9) extent(9) startIdx(1)
+    ```
+  }];
 
   let arguments = (ins Optional<IntOrIndex>:$lowerbound,
                        Optional<IntOrIndex>:$upperbound,
@@ -144,6 +180,26 @@ class OpenACC_DataEntryOp<string mnemonic, string clause, list<Trait> traits = [
                        OptionalAttr<StrAttr>:$name);
   let results = (outs OpenACC_PointerLikeTypeInterface:$accPtr);
 
+  let description = [{
+    - `varPtr`: The address of variable to copy.
+    - `varPtrPtr`: Specifies the address of varPtr - only used when the variable
+    copied is a field in a struct. This is important for OpenACC due to implicit
+    attach semantics on data clauses (2.6.4).
+    - `bounds`: Used when copying just slice of array or array's bounds are not
+    encoded in type. They are in rank order where rank 0 is inner-most dimension.
+    - `dataClause`: Keeps track of the data clause the user used. This is because
+    the acc operations are decomposed. So a 'copy' clause is decomposed to both 
+    `acc.copyin` and `acc.copyout` operations, but both have dataClause that
+    specifies `acc_copy` in this field.
+    - `structured`: Flag to note whether this is associated with structured region
+    (parallel, kernels, data) or unstructured (enter data, exit data). This is
+    important due to spec specifically calling out structured and dynamic reference
+    counters (2.6.7).
+    - `implicit`: Whether this is an implicitly generated operation, such as copies
+    done to satisfy "Variables with Implicitly Determined Data Attributes" in 2.6.2.
+    - `name`: Holds the name of variable as specified in user clause (including bounds).
+  }];
+
   let assemblyFormat = [{
     `varPtr` `(` $varPtr `:` type($varPtr) `)`
     oilist(
@@ -245,6 +301,26 @@ class OpenACC_DataExitOp<string mnemonic, string clause, list<Trait> traits = []
                        DefaultValuedAttr<BoolAttr, "false">:$implicit,
                        OptionalAttr<StrAttr>:$name);
 
+  let description = [{
+    - `varPtr`: The address of variable to copy back to. This only applies to
+    `acc.copyout`
+    - `accPtr`: The acc address of variable. This is the link from the data-entry
+    operation used.
+    - `bounds`: Used when copying just slice of array or array's bounds are not
+    encoded in type. They are in rank order where rank 0 is inner-most dimension.
+    - `dataClause`: Keeps track of the data clause the user used. This is because
+    the acc operations are decomposed. So a 'copy' clause is decomposed to both 
+    `acc.copyin` and `acc.copyout` operations, but both have dataClause that
+    specifies `acc_copy` in this field.
+    - `structured`: Flag to note whether this is associated with structured region
+    (parallel, kernels, data) or unstructured (enter data, exit data). This is
+    important due to spec specifically calling out structured and dynamic reference
+    counters (2.6.7).
+    - `implicit`: Whether this is an implicitly generated operation, such as copies
+    done to satisfy "Variables with Implicitly Determined Data Attributes" in 2.6.2.
+    - `name`: Holds the name of variable as specified in user clause (including bounds).
+  }];
+
   let assemblyFormat = [{
     `accPtr` `(` $accPtr `:` type($accPtr) `)`
     oilist(