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

[Alexey Bataev via Phabricator via cfe-commits]

ABataev added a comment.

In D79972#2124108 <https://reviews.llvm.org/D79972#2124108>, @cchen wrote:

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


You can do this patch. But at first, you need to commit the runtime part of the patch that supports it, and the part that introduces stride support.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D79972