1*330d8983SJohannes Doerfert // Check that omp tile (introduced in OpenMP 5.1) is permitted and behaves when
2*330d8983SJohannes Doerfert // strictly nested within omp target.
3*330d8983SJohannes Doerfert
4*330d8983SJohannes Doerfert // RUN: %libomptarget-compile-generic -fopenmp-version=51
5*330d8983SJohannes Doerfert // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
6*330d8983SJohannes Doerfert
7*330d8983SJohannes Doerfert #include <stdio.h>
8*330d8983SJohannes Doerfert
9*330d8983SJohannes Doerfert #define I_NTILES 8
10*330d8983SJohannes Doerfert #define J_NTILES 9
11*330d8983SJohannes Doerfert #define I_NELEMS 2
12*330d8983SJohannes Doerfert #define J_NELEMS 3
13*330d8983SJohannes Doerfert
main()14*330d8983SJohannes Doerfert int main() {
15*330d8983SJohannes Doerfert int order[I_NTILES][J_NTILES][I_NELEMS][J_NELEMS];
16*330d8983SJohannes Doerfert int i, j;
17*330d8983SJohannes Doerfert #pragma omp target map(tofrom: i, j)
18*330d8983SJohannes Doerfert {
19*330d8983SJohannes Doerfert int next = 0;
20*330d8983SJohannes Doerfert #pragma omp tile sizes(I_NELEMS, J_NELEMS)
21*330d8983SJohannes Doerfert for (i = 0; i < I_NTILES * I_NELEMS; ++i) {
22*330d8983SJohannes Doerfert for (j = 0; j < J_NTILES * J_NELEMS; ++j) {
23*330d8983SJohannes Doerfert int iTile = i / I_NELEMS;
24*330d8983SJohannes Doerfert int jTile = j / J_NELEMS;
25*330d8983SJohannes Doerfert int iElem = i % I_NELEMS;
26*330d8983SJohannes Doerfert int jElem = j % J_NELEMS;
27*330d8983SJohannes Doerfert order[iTile][jTile][iElem][jElem] = next++;
28*330d8983SJohannes Doerfert }
29*330d8983SJohannes Doerfert }
30*330d8983SJohannes Doerfert }
31*330d8983SJohannes Doerfert int expected = 0;
32*330d8983SJohannes Doerfert for (int iTile = 0; iTile < I_NTILES; ++iTile) {
33*330d8983SJohannes Doerfert for (int jTile = 0; jTile < J_NTILES; ++jTile) {
34*330d8983SJohannes Doerfert for (int iElem = 0; iElem < I_NELEMS; ++iElem) {
35*330d8983SJohannes Doerfert for (int jElem = 0; jElem < J_NELEMS; ++jElem) {
36*330d8983SJohannes Doerfert int actual = order[iTile][jTile][iElem][jElem];
37*330d8983SJohannes Doerfert if (expected != actual) {
38*330d8983SJohannes Doerfert printf("error: order[%d][%d][%d][%d] = %d, expected %d\n",
39*330d8983SJohannes Doerfert iTile, jTile, iElem, jElem, actual, expected);
40*330d8983SJohannes Doerfert return 1;
41*330d8983SJohannes Doerfert }
42*330d8983SJohannes Doerfert ++expected;
43*330d8983SJohannes Doerfert }
44*330d8983SJohannes Doerfert }
45*330d8983SJohannes Doerfert }
46*330d8983SJohannes Doerfert }
47*330d8983SJohannes Doerfert // Tiling leaves the loop variables with their values from the final iteration
48*330d8983SJohannes Doerfert // rather than with the usual +1.
49*330d8983SJohannes Doerfert expected = I_NTILES * I_NELEMS - 1;
50*330d8983SJohannes Doerfert if (i != expected) {
51*330d8983SJohannes Doerfert printf("error: i = %d, expected %d\n", i, expected);
52*330d8983SJohannes Doerfert return 1;
53*330d8983SJohannes Doerfert }
54*330d8983SJohannes Doerfert expected = J_NTILES * J_NELEMS - 1;
55*330d8983SJohannes Doerfert if (j != expected) {
56*330d8983SJohannes Doerfert printf("error: j = %d, expected %d\n", j, expected);
57*330d8983SJohannes Doerfert return 1;
58*330d8983SJohannes Doerfert }
59*330d8983SJohannes Doerfert // CHECK: success
60*330d8983SJohannes Doerfert printf("success\n");
61*330d8983SJohannes Doerfert return 0;
62*330d8983SJohannes Doerfert }
63