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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
