| // { dg-do run { target offload_device_nonshared_as } } |
| |
| #include <string.h> |
| #include <assert.h> |
| #include <stdlib.h> |
| |
| volatile int yy = 4, zz = 2, str_str = 2; |
| |
| int main() |
| { |
| int *arr; |
| int x = 5; |
| int arr2d[10][10]; |
| |
| arr = calloc (100, sizeof (int)); |
| |
| /* Update whole reshaped array. */ |
| |
| #pragma omp target enter data map(to: arr[:100]) |
| |
| for (int j = 0; j < x; j++) |
| for (int i = 0; i < 10; i++) |
| arr[j * 10 + i] = i ^ j; |
| |
| #pragma omp target update to(([10][x]) arr) |
| |
| #pragma omp target exit data map(from: arr[:100]) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| if (j < x) |
| assert (arr[j * 10 + i] == i ^ j); |
| else |
| assert (arr[j * 10 + i] == 0); |
| |
| |
| /* Strided update. */ |
| |
| memset (arr, 0, 100 * sizeof (int)); |
| |
| #pragma omp target enter data map(to: arr[:100]) |
| |
| for (int j = 0; j < 20; j++) |
| for (int i = 0; i < 5; i++) |
| arr[j * 5 + i] = i + j; |
| |
| #pragma omp target update to(([5][5]) arr[0:3][0:3:2]) |
| |
| #pragma omp target exit data map(from: arr[:100]) |
| |
| for (int j = 0; j < 20; j++) |
| for (int i = 0; i < 5; i++) |
| if (j < 3 && (i & 1) == 0 && i < 6) |
| assert (arr[j * 5 + i] == i + j); |
| else |
| assert (arr[j * 5 + i] == 0); |
| |
| |
| /* Reshaped update, contiguous. */ |
| |
| memset (arr, 0, 100 * sizeof (int)); |
| |
| #pragma omp target enter data map(to: arr[:100]) |
| |
| for (int j = 0; j < 20; j++) |
| for (int i = 0; i < 5; i++) |
| arr[j * 5 + i] = 2 * j + i; |
| |
| #pragma omp target update to(([5][5]) arr[0:5][0:5]) |
| |
| #pragma omp target exit data map(from: arr[:100]) |
| |
| for (int j = 0; j < 20; j++) |
| for (int i = 0; i < 5; i++) |
| if (j < 5 && i < 5) |
| assert (arr[j * 5 + i] == 2 * j + i); |
| else |
| assert (arr[j * 5 + i] == 0); |
| |
| |
| /* Strided update on actual array. */ |
| |
| memset (arr2d, 0, 100 * sizeof (int)); |
| |
| #pragma omp target enter data map(to: arr2d) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| arr2d[j][i] = j + 2 * i; |
| |
| #pragma omp target update to(arr2d[0:5:2][5:2]) |
| |
| #pragma omp target exit data map(from: arr2d) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| if ((j & 1) == 0 && i >= 5 && i < 7) |
| assert (arr2d[j][i] == j + 2 * i); |
| else |
| assert (arr2d[j][i] == 0); |
| |
| |
| /* Update with non-constant bounds. */ |
| |
| memset (arr, 0, 100 * sizeof (int)); |
| |
| #pragma omp target enter data map(to: arr[:100]) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| arr[j * 10 + i] = (2 * j) ^ i; |
| |
| x = 3; |
| int y = yy, z = zz, str = str_str; |
| /* This is actually [0:3:2] [4:2:2]. */ |
| #pragma omp target update to(([10][10]) arr[0:x:2][y:z:str]) |
| |
| #pragma omp target exit data map(from: arr[:100]) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| if ((j & 1) == 0 && j < 6 && (i & 1) == 0 && i >= 4 && i < 8) |
| assert (arr[j * 10 + i] == (2 * j) ^ i); |
| else |
| assert (arr[j * 10 + i] == 0); |
| |
| |
| /* Update with full "major" dimension. */ |
| |
| memset (arr, 0, 100 * sizeof (int)); |
| |
| #pragma omp target enter data map(to: arr[:100]) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| arr[j * 10 + i] = i + j; |
| |
| #pragma omp target update to(([10][10]) arr[0:10][3:1]) |
| |
| #pragma omp target exit data map(from: arr[:100]) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| if (i == 3) |
| assert (arr[j * 10 + i] == i + j); |
| else |
| assert (arr[j * 10 + i] == 0); |
| |
| |
| /* Update with full "minor" dimension. */ |
| |
| memset (arr, 0, 100 * sizeof (int)); |
| |
| #pragma omp target enter data map(to: arr[:100]) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| arr[j * 10 + i] = 3 * (i + j); |
| |
| #pragma omp target update to(([10][10]) arr[3:2][0:10]) |
| |
| #pragma omp target exit data map(from: arr[:100]) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| if (j >= 3 && j < 5) |
| assert (arr[j * 10 + i] == 3 * (i + j)); |
| else |
| assert (arr[j * 10 + i] == 0); |
| |
| |
| /* Rectangle update. */ |
| |
| memset (arr, 0, 100 * sizeof (int)); |
| |
| #pragma omp target enter data map(to: arr[:100]) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| arr[j * 10 + i] = 5 * (i + j); |
| |
| #pragma omp target update to(([10][10]) arr[3:2][0:9]) |
| |
| #pragma omp target exit data map(from: arr[:100]) |
| |
| for (int j = 0; j < 10; j++) |
| for (int i = 0; i < 10; i++) |
| if (j >= 3 && j < 5 && i < 9) |
| assert (arr[j * 10 + i] == 5 * (i + j)); |
| else |
| assert (arr[j * 10 + i] == 0); |
| |
| |
| /* One-dimensional strided update. */ |
| |
| memset (arr, 0, 100 * sizeof (int)); |
| |
| #pragma omp target enter data map(to: arr[:100]) |
| |
| for (int i = 0; i < 100; i++) |
| arr[i] = i + 99; |
| |
| #pragma omp target update to(([100]) arr[3:33:3]) |
| |
| #pragma omp target exit data map(from: arr[:100]) |
| |
| for (int i = 0; i < 100; i++) |
| if (i >= 3 && ((i - 3) % 3) == 0) |
| assert (arr[i] == i + 99); |
| else |
| assert (arr[i] == 0); |
| |
| |
| /* One-dimensional strided update without explicit array shape. */ |
| |
| memset (arr, 0, 100 * sizeof (int)); |
| |
| #pragma omp target enter data map(to: arr[:100]) |
| |
| for (int i = 0; i < 100; i++) |
| arr[i] = i + 121; |
| |
| #pragma omp target update to(arr[3:33:3]) |
| |
| #pragma omp target exit data map(from: arr[:100]) |
| |
| for (int i = 0; i < 100; i++) |
| if (i >= 3 && ((i - 3) % 3) == 0) |
| assert (arr[i] == i + 121); |
| else |
| assert (arr[i] == 0); |
| |
| free (arr); |
| |
| return 0; |
| } |