[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.

In D79972#2049838 <https://reviews.llvm.org/D79972#2049838>, @ABataev wrote:

> How are you going to pass this non-contiguous data in the runtime? Are you going to map it in a loop or convert this non-contiguous data into the contiguous and map it as a contiguous chunk of data? Your presentation provides interface only interface changes but has nothing about implementation in the runtime.


Hi Alexey, thanks for asking. The runtime implementation I'm thinking of is to convert the non-contiguous data into several chunks of contiguous.

For example:

int arr[3][3][3];

#pragma omp target update to (arr[1:2][1:2][0:2])

We can visualize the noncontiguous data as below (X is the data we want to transfer, O is the data want don't bother with):

Dim 0 = {Offset: 0, Count: 1, Stride: 4bytes (int)}
XXO

Dim 1 = {Offset: 1, Count: 2, Stride: 12bytes (4 * 3 - since Dim 0 has 3 elements)
OOO
XXO
XXO

Dim 2 = {Offset: 1, Count: 2, Stride: 36 bytes (12 * 3 since Dim 1 has 3 elements)
OOO
OOO
OOO
\\\\\
OOO
XXO
XXO
\\\\\
OOO
XXO
XXO

For the visualization, we know that we want to transfer 4 contiguous chunks and the runtime code could be something similar to:

  // we expect this loop to transfer 4 contiguous chunks:
  // arr[1][1][0:2]
  // arr[1][2][0:2]
  // arr[2][1][0:2]
  // arr[2][2][0:2]
  for (int i = Dim[2].offset; i < Dim[2].count; i++) {
    for (int j = Dim[1].offset; j < Dim[1].count; j++) {
      ptr = bast_ptr + Dim[2].stride * i + Dim[1].stride * j + Dim[2].stride * Dim[0].offset;
      size = Dim[0].count * Dim[0].stride;  // we can hoist it I think
      transfer(ptr, size, /*flag or some other stuff...*/);
    }
  }


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D79972