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