[flang-commits] [flang] [OpenMP][Flang] Add "IsolatedFromAbove" trait to omp.target (PR #67164)

[Akash Banerjee via flang-commits]

TIFitis wrote:

> > ```
> >    integer :: n
> >    integer :: a(n, 1024)
> >    !$omp target 
> >    a(11, 22) = 33
> >    !$omp end target
> > ```
> 
> This map operation is indeed missing bounds and this is not correct/expected.

I've added the bounds information for implicit map ops. Here's what the generated FIR looks like now for the same test.
```
func.func @_QPomp_target_implicit_bounds(%arg0: !fir.ref<i32> {fir.bindc_name = "n"}) {
  %0 = fir.load %arg0 : !fir.ref<i32>
  %1 = fir.convert %0 : (i32) -> i64
  %2 = fir.convert %1 : (i64) -> index
  %c0 = arith.constant 0 : index
  %3 = arith.cmpi sgt, %2, %c0 : index
  %4 = arith.select %3, %2, %c0 : index
  %c1024_i64 = arith.constant 1024 : i64
  %5 = fir.convert %c1024_i64 : (i64) -> index
  %c0_0 = arith.constant 0 : index
  %6 = arith.cmpi sgt, %5, %c0_0 : index
  %7 = arith.select %6, %5, %c0_0 : index
  %8 = fir.alloca !fir.array<?x1024xi32>, %4 {bindc_name = "a", uniq_name = "_QFomp_target_implicit_boundsEa"}
  %c1 = arith.constant 1 : index
  %c0_1 = arith.constant 0 : index
  %9 = arith.subi %4, %c1 : index
  %10 = omp.bounds lower_bound(%c0_1 : index) upper_bound(%9 : index) extent(%4 : index) stride(%c1 : index) start_idx(%c1 : index)
  %c0_2 = arith.constant 0 : index
  %11 = arith.subi %7, %c1 : index
  %12 = omp.bounds lower_bound(%c0_2 : index) upper_bound(%11 : index) extent(%7 : index) stride(%c1 : index) start_idx(%c1 : index)
  %13 = omp.map_info var_ptr(%8 : !fir.ref<!fir.array<?x1024xi32>>) map_clauses(literal, implicit, exit_release_or_enter_alloc) capture(ByCopy) bounds(%10, %12) -> !fir.ref<!fir.array<?x1024xi32>> {name = "a"}
  omp.target map_entries(%13 -> %arg1 : !fir.ref<!fir.array<?x1024xi32>>) block_args(%4, %7 : index, index) {
  ^bb0(%arg1: !fir.ref<!fir.array<?x1024xi32>>, %arg2: index, %arg3: index):
    %c33_i32 = arith.constant 33 : i32
    %14 = fir.convert %arg1 : (!fir.ref<!fir.array<?x1024xi32>>) -> !fir.ref<!fir.array<?xi32>>
    %c1_3 = arith.constant 1 : index
    %c0_4 = arith.constant 0 : index
    %c11_i64 = arith.constant 11 : i64
    %15 = fir.convert %c11_i64 : (i64) -> index
    %16 = arith.subi %15, %c1_3 : index
    %17 = arith.muli %c1_3, %16 : index
    %18 = arith.addi %17, %c0_4 : index
    %19 = arith.muli %c1_3, %arg2 : index
    %c22_i64 = arith.constant 22 : i64
    %20 = fir.convert %c22_i64 : (i64) -> index
    %21 = arith.subi %20, %c1_3 : index
    %22 = arith.muli %19, %21 : index
    %23 = arith.addi %22, %18 : index
    %24 = fir.coordinate_of %14, %23 : (!fir.ref<!fir.array<?xi32>>, index) -> !fir.ref<i32>
    fir.store %c33_i32 to %24 : !fir.ref<i32>
    omp.terminator
  }
  return
}
```

> This example also shows that it is concerning making all mapping decisions so early in order to satisfy the IsolatedFromAbove property. Namely, this case would ideally be map(tofrom:a(11,22)) not a map(tofrom:a). Maybe it is indeed possible for this case - but consider a simple loop modifying `a(11,22+ii)`. Wouldn't it be great to just map what is needed? Hard to get this info at lowering. https://www.openmp.org/spec-html/5.2/openmpsu60.html
> 
> I wonder if it should be considered to add the trait IsolatedFromAbove sometime after lowering and after running the passes that will decide values passed by reference/value and those that decide the mapping actions for implicit data mapping attributes.

Please correct me if I'm wrong but what I can tell Clang doesn't map this as map(tofrom:a(11,22)).

The following C codes generate identical code apart from the offload_maptypes flags.

Implicit-C
```
void foo(int n){
  int a[n][1024];
  #pragma omp target map(tofrom:a)
  {
    a[11][22] = 33;
  }
  return;
}
```

Explicit-C
```
void foo(int n){
  int a[n][1024];
  #pragma omp target
  {
    a[11][22] = 33;
  }
  return;
}
```

Only when we change the pragma to #pragma omp target map(alloc:a[11:22]) do I see notable changes.

https://github.com/llvm/llvm-project/pull/67164