blob: 808c5f9ceae6cfaf5dd2f85111703826273535c3 [file]
// { 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;
}