[PATCH] D79972: [OpenMP5.0] map item can be non-contiguous for target update

[Chi Chun Chen via Phabricator via cfe-commits]

cchen added a comment.

@ABataev , I'm considering emitting an extra dimension for a non-contiguous descriptor to support stride in this patch (stride = 1 in array section is just a special case for computing stride, however, the formula computing stride do not change). Do you think I should do it in this patch?

Computing of stride after support stride in array section:

  int arr[5][5][5];
  #pragma omp target update to(arr[0:2:2][1:2:1][0:2:2]
  
  D0: { offset = 0, count = 1, stride = 4 }                                           // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
  D1: { offset = 0, count = 2, stride = 4 * 1 * 2 = 8 }                        // stride = unit size * (production of dimension size of D0) * D1.stride = 4 * 1 * 2 = 8
  D2: { offset = 0, count = 1, stride = 4 * (1 * 5) * 1 = 20  }             // stride = unit size * (production of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
  D3: { offset = 0, count = 2, stride = 4 * (1 * 5 * 5) * 2 = 200 }      // stride = unit size * (production of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 2 = 200

For the case in this patch (stride = 1), we can use the same formula for computing stride with extra dimension:

  int arr[5][5][5];
  #pragma omp target update to(arr[0:2][1:2][0:2]
  
  D0: { offset = 0, count = 1, stride = 4 }                                          // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
  D1: { offset = 0, count = 2, stride = 4 * 1 * 1 = 4 }                        // stride = unit size * (production of dimension size of D0) * D1.stride = 4 * 1 * 1 = 4
  D2: { offset = 0, count = 1, stride = 4 * (1 * 5) * 1 = 20  }            // stride = unit size * (production of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
  D3: { offset = 0, count = 2, stride = 4 * (1 * 5 * 5) * 1 = 100 }     // stride = unit size * (production of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 1 = 100

The extra dimension does not affect the runtime implementation at all since runtime will try to merge inner dimensions if they are contiguous. Take the above case for example (arr[0:2][1:2][0:2]):
The product of count and stride for D0 is 4 which is the same as the stride of D1 <https://reviews.llvm.org/D1>, therefore, runtime just ignores D0.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D79972/new/

https://reviews.llvm.org/D79972