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