[openmp-target-tests] OpenMP 4.5 Target data test cases

[sergiop]

Summary

This patch contains test cases for offloading in openmp 4.5. Specifically, these test cases cover the 'target data' construct.

test_target_data_if.c
test_target_data_map_array_sections.c
test_target_data_use_device_ptr.c
test_target_data_map.c
test_target_data_map_classes.cpp
test_target_data_map_devices.c

[Hahnfeld]

Hi,

I don't see any new files...? I'm also not entirely sure that we should have Fortran tests? @jlpeyton

The changes to offloading_success.{c,cpp} are already in the tree, please rebase your changes to the latest commit.
Phabricator say "Context not available". Did you generate the patches with the -U option? See http://llvm.org/docs/Phabricator.html#requesting-a-review-via-the-web-interface

Regards,
Jonas

[sergiop]

Hi Jonas,

Thank you for the comments. I generated again the patch following your comments.

Best regards,

Sergio

[Hahnfeld]

In D34701#794365, @sergiop wrote:

Thank you for the comments. I generated again the patch following your comments.

Now the patch looks like the changes done by clang-format to already existing tests?

[sergiop]

Jonas,

I squashed all the commits into one. Now, I see that Diff has the creation of the new files.

Let me know if this patch works.

Best

[schandra]

Hi Jonas

Did you have further comments on this commit?

Thanks
Sunita

In D34701#797206, @sergiop wrote:

Jonas,

I squashed all the commits into one. Now, I see that Diff has the creation of the new files.

Let me know if this patch works.

Best

In D34701#797206, @sergiop wrote:

Jonas,

I squashed all the commits into one. Now, I see that Diff has the creation of the new files.

Let me know if this patch works.

Best

In D34701#797206, @sergiop wrote:

Jonas,

I squashed all the commits into one. Now, I see that Diff has the creation of the new files.

Let me know if this patch works.

Best

[Hahnfeld]

In D34701#823166, @schandra wrote:

Hi Jonas

Did you have further comments on this commit?

Thanks
Sunita

Sorry that it's taken so long! Still not sure about the Fortran test as clang will currently not be able to handle this.

Additionally, these tests currently won't pass because omp_is_initial_device() will return 1 when offloading to the host, correct? See D35719 why we had to disable libomptarget for 5.0.

Find some more comments inline. In general, the tests should be brought to a consistent coding style (braces, end-comments, comment style and so on).

libomptarget/test/offloading/target_data/test_target_data_map.c

28

Please, add braces here. This is really hard to read and the style must be consistent accross the tests.

(The same applies to all other places in this file where this occurrs)

71–74

This does not test to. Did you mean to use += 1?

75

s/data target/target data/

204

s/cheking/checking/

libomptarget/test/offloading/target_data/test_target_data_map_array_sections.c

384

If I remember correctly, this will not force an array of length N. To avoid this misleading information, would you mind changing the signature to int *?

libomptarget/test/offloading/target_data/test_target_data_map_classes.cpp

35

Do we need an object on the heap, does this also work with the stack? If it should, please add another test method to check that.

43–44

I don't really understand the purpose of this test: Wouldn't it be more intuitive if the user hadn't to worry about how the array got there? What would happen if mapping obj->h_array directly, would compiler and runtime perform this work?

libomptarget/test/offloading/target_data/test_target_data_map_devices.c

26–27

I know that this variable allocation on the stack works, but could you change this to a malloc to avoid this rather ugly part of the C language?

131

Does this really work with libomptarget? If so, is this mandated by the standard and do we want to test this?

libomptarget/test/offloading/target_data/test_target_data_use_device_ptr.c

31

array_device[i]?

[schandra]

In D34701#828682, @Hahnfeld wrote:

In D34701#823166, @schandra wrote:

Hi Jonas

Did you have further comments on this commit?

Thanks
Sunita

Sorry that it's taken so long! Still not sure about the Fortran test as clang will currently not be able to handle this.

Additionally, these tests currently won't pass because omp_is_initial_device() will return 1 when offloading to the host, correct? See D35719 why we had to disable libomptarget for 5.0.

Find some more comments inline. In general, the tests should be brought to a consistent coding style (braces, end-comments, comment style and so on).

Great!! Thanks Jonas, we will fix these shortly.
Sunita

[sergiop]

Hi Jonas,

We addressed your comments about the code + removed the fortran file (you're right, sorry about that). I'm submitting an update to the patch.

Thank you for the time. Best

libomptarget/test/offloading/target_data/test_target_data_map_classes.cpp

43–44

Unfortunately, the mapping is a shallow copy of the object. This means that the address of the array is not translated automatically.

libomptarget/test/offloading/target_data/test_target_data_map_devices.c

131

In the spec there is not restriction about assigning computation/data movement to a device that has an id higher than the number of actual devices.

[sergiop]

Updated the patch with respect to the comments. Special focus on: 1) not using omp_is_initial_device as a condition to fail or pass, 2) removing fortran files, 3) consistent style.

test_target_data_if.c
test_target_data_map_array_sections.c
test_target_data_use_device_ptr.c
test_target_data_map.c
test_target_data_map_classes.cpp
test_target_data_map_devices.c

[Hahnfeld]

1. Is there as specific reason for using both puts and printf? There are also calls to printf without formatting arguments...
2. Do these tests actually pass with a current version of clang?

libomptarget/test/offloading/target_data/test_target_data_if.c

38–40

Again thinking about this, the check may be wrong: The OpenMP standard says that the target data and target constructs execute on the host, if the if clause evaluates to false. This check is assuming that it doesn't execute at all, right?

libomptarget/test/offloading/target_data/test_target_data_map_array_sections.c

428

Please change this signature as well.

435

Same

libomptarget/test/offloading/target_data/test_target_data_map_devices.c

131

"The device expression must evaluate to a non-negative integer value less than the value of omp_get_num_devices()."
OpenMP 4.5, p. 97

[AndreyChurbanov]

libomptarget/test/offloading/target_data/test_target_data_if.c

38–40

This check looks like a dead code. The size gets three values - 256 to execute on host, 512 and 758 to execute on device. Neither of these values is less than 256, so (size < 256) is the same as false.

Maybe the intention was to check all 1024 values, and zeros should hold for (i >= size).

libomptarget/test/offloading/target_data/test_target_data_if.c

+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+#include <omp.h>
+#include <stdio.h>
+
+/*Test for OpenMP 4.5 target data with if*/
+int main() {
+  int a[1024];
+  int b[1024];
+  int c[1024];
+  int size, i = 0, errors = 0, isHost = -1;
+
+  for (i = 0; i < 1024; i++) {
+    a[i] = 1;
+    b[i] = i;
+    c[i] = 0;
+  }
+
+  for (size = 256; size < 1024; size += 256) {
+#pragma omp target data if (size > 256) map(to : size)  \
+      map(tofrom : c[0 : size])                         \
+      map(to : a[0 : size], b[0 : size], i) 
+    {
+#pragma omp target if (size > 256) map(tofrom : isHost)
+      {
+        isHost = omp_is_initial_device();
+        for (i = 0; i < size; i++)
+          c[i] = a[i] + b[i];
+      } /*end target*/
+    } /*end target data*/
+
+    /*checking results*/
+    for (i = 0; i < size; i++) {
+      if (size < 256) {
+        if (c[i] != 0) {
+          errors = 1;
+        }
+      } else {
+        if (c[i] != i + 1) {
+          errors = 1;
+        }
+      } /*end-else*/
+    }
+  } /*end-for*/
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}

libomptarget/test/offloading/target_data/test_target_data_map.c

+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 1000
+
+/* Test for OpenMP 4.5 target data map(from: ) */
+int test_map_from() {
+
+  printf("test_map_from\n");
+
+  int sum = 0, sum2 = 0, errors = 0, isHost = 0;
+
+  /* host arrays: heap and stack */
+  int *h_array_h = (int *)malloc(N * sizeof(int));
+  int h_array_s[N];
+
+#pragma omp target data map(from : h_array_h[0 : N])  \
+    map(from : h_array_s[0 : N])
+  {
+#pragma omp target map(from : isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 0; i < N; ++i) {
+        h_array_h[i] = 1;
+        h_array_s[i] = 2;
+      }
+    } /*end target*/
+  } /*end target data*/
+
+  /*checking results*/
+  for (int i = 0; i < N; ++i) {
+    sum += h_array_h[i];
+    sum2 += h_array_s[i];
+  }
+  
+  free(h_array_h);
+  errors = (N != sum) || (2 * N != sum2);
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s: sum=%d, sum2=%d, N=%d\n", (isHost ? "host" : "device"), sum, sum2, N);
+
+  return errors;
+}
+
+/* Test for OpenMP 4.5 target data map(tofrom: ) */
+int test_map_tofrom() {
+
+  printf("test_map_tofrom\n");
+
+  int sum = 0, sum2 = 0, isHost = 0, errors = 0;
+
+  /* host arrays: heap and stack */
+  int *h_array_h = (int *)malloc(N * sizeof(int));
+  int h_array_s[N];
+
+  for (int i = 0; i < N; ++i) {
+    h_array_h[i] = 0;
+    h_array_s[i] = 0;
+  }
+
+#pragma omp target data map(tofrom : h_array_h[0 : N])    \
+    map(tofrom : h_array_s[0 : N]) 
+  {
+#pragma omp target map(tofrom : isHost)
+    { 
+      isHost += omp_is_initial_device();
+      for (int i = 0; i < N; ++i) {
+        h_array_h[i] += 1;
+        h_array_s[i] += 1;
+      }
+    } /* end target */
+  } /* end target data */
+
+   /* checking errors */
+  for (int i = 0; i < N; ++i) {
+    sum += h_array_h[i];
+    sum2 += h_array_s[i];
+  }
+
+  free(h_array_h);
+  errors = (N != sum) || (N != sum2);
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s: sum=%d, sum2=%d, N=%d\n", (isHost ? "host" : "device"), sum, sum2, N);
+
+  return errors;
+}
+
+/* Test for OpenMP 4.5 target data map(to: ) */
+int test_map_to() {
+
+  printf("test_map_to\n");
+
+  int sum = 0, sum2 = 0, errors = 0, isHost = 0;
+  
+  /* host arrays: heap and stack */
+  int *h_array_h = (int *)malloc(N * sizeof(int));
+  int *h_array2_h = (int *)malloc(N * sizeof(int));
+  int h_array_s[N];
+  int h_array2_s[N];
+
+  /* initializing arrays */
+  for (int i = 0; i < N; ++i) {
+    h_array_h[i] = 1;
+    h_array_s[i] = 1;
+    h_array2_h[i] = 0;
+    h_array2_s[i] = 0;
+  }
+
+  /* device arrays to get the data from the device */
+  /* pointer arithmetic is not supported on the devices for*/
+  /* the device address returned by omp_target_alloc*/
+  /* section 3.5.1 omp_target_alloc. OpenMP API Version 4.5 Nov 2015*/
+  int *d_array =
+      (int *)omp_target_alloc(N * sizeof(int), omp_get_default_device());
+  int *d_array2 =
+      (int *)omp_target_alloc(N * sizeof(int), omp_get_default_device());
+  int *d_ishost = (int *)omp_target_alloc(sizeof(int), omp_get_default_device());
+
+#pragma omp target data map(to : h_array_h[0 : N])  \
+    map(to : h_array_s[0 : N]) 
+  {
+#pragma omp target is_device_ptr(d_array, d_array2, d_ishost) map(to: isHost)
+    {
+      d_ishost[0] = isHost + omp_is_initial_device();
+      for (int i = 0; i < N; ++i) {
+        d_array[i] = h_array_h[i];
+        d_array2[i] = h_array_s[i];
+      }
+    } /* end target */
+  } /* end target data*/
+
+  /* copy from d to h */
+  omp_target_memcpy(h_array2_h, d_array, N * sizeof(int), 0, 0,
+                    omp_get_initial_device(), omp_get_default_device());
+  omp_target_memcpy(h_array2_s, d_array2, N * sizeof(int), 0, 0,
+                    omp_get_initial_device(), omp_get_default_device());
+  omp_target_memcpy(&isHost, d_ishost, sizeof(int), 0, 0,
+                    omp_get_initial_device(), omp_get_default_device());
+  /* deallocating device arrays */
+  omp_target_free(d_array, omp_get_default_device());
+  omp_target_free(d_array2, omp_get_default_device());
+  omp_target_free(d_ishost, omp_get_default_device());
+
+  /* checking errors */
+  for (int i = 0; i < N; ++i) {
+    sum += h_array2_h[i];
+    sum2 += h_array2_s[i];
+  }
+
+  free(h_array_h);
+  free(h_array2_h);
+  errors = (N != sum) || (N != sum2);
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s: sum=%d, sum2=%d, N=%d\n", (isHost ? "host" : "device"), sum, sum2, N);
+
+  return errors;
+}
+
+/* Test for OpenMP 4.5 target data map(to: ) and map(from:) */
+int test_map_to_from() {
+
+  printf("test_map_to_from\n");
+
+  int sum = 0, errors = 0, isHost = 0, isHost2 = 0;
+  int *h_array_h = (int *)malloc(N * sizeof(int));
+  int *h_array2_h = (int *)malloc(N * sizeof(int));
+
+  for (int i = 0; i < N; ++i) {
+    h_array_h[i] = 1;
+    h_array2_h[i] = 0;
+  }
+
+#pragma omp target data map(to : h_array_h[0 : N])  map(from : h_array2_h[0 : N])  
+  {
+#pragma omp target map(from: isHost) map(to: isHost2)
+    {
+      isHost = isHost2 + omp_is_initial_device();
+      for (int i = 0; i < N; ++i)
+        h_array2_h[i] = h_array_h[i];
+    } /* end target */
+  } /*end target data*/
+
+  /* checking errors */
+  for (int i = 0; i < N; ++i)
+    sum += h_array2_h[i];
+
+  free(h_array_h);
+  free(h_array2_h);
+  errors = N - sum;
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s: sum=%d, N=%d\n",(isHost ? "host" : "device"), sum, N);
+
+  return errors;
+}
+
+/* Test for OpenMP 4.5 target data map(alloc:) */
+int test_map_alloc() {
+
+  puts("test_map_alloc");
+
+  int sum = 0, errors = 0, isHost = 0;
+  int *h_array_h = (int *)malloc(N * sizeof(int));
+
+  /* pointer arithmetic is not supported on the devices for*/
+  /* the device address returned by omp_target_alloc*/
+  /* section 3.5.1 omp_target_alloc. OpenMP API Version 4.5 Nov 2015*/
+  int *d_sum = (int *)omp_target_alloc(sizeof(int), omp_get_default_device());
+  int *d_ishost = (int *)omp_target_alloc(sizeof(int), omp_get_default_device());
+
+#pragma omp target data map(alloc : h_array_h[0 : N])
+  {
+#pragma omp target is_device_ptr(d_sum, d_ishost)
+    {
+      d_ishost[0] = omp_is_initial_device();
+      for (int i = 0; i < N; ++i) 
+        h_array_h[i] = 1;
+      
+      /* checking errors*/
+      d_sum[0] = 0; 
+      for (int i = 0; i < N; ++i)
+        d_sum[0] += h_array_h[i];
+    } /* end target*/
+    omp_target_memcpy(&sum, d_sum, sizeof(int), 0, 0,
+                                  omp_get_initial_device(),
+                                  omp_get_default_device());
+    omp_target_memcpy(&isHost, d_ishost, sizeof(int), 0, 0,
+                                  omp_get_initial_device(),
+                                  omp_get_default_device());
+  } /* end target data */
+  omp_target_free(d_sum, omp_get_default_device());
+  omp_target_free(d_ishost, omp_get_default_device());
+
+  free(h_array_h);
+  errors = N - sum;
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s: sum=%d, N=%d\n", (isHost ? "host" : "device"), sum, N);
+
+  return errors;
+}
+
+int main() {
+
+  int errors = 0;
+
+  errors += test_map_from();
+  errors += test_map_tofrom();
+  errors += test_map_to();
+  errors += test_map_to_from();
+  errors += test_map_alloc();
+
+  return errors;
+}

libomptarget/test/offloading/target_data/test_target_data_map_array_sections.c

+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 1000
+
+void init_1d(int* a);
+void init_2d(int a[N][2]);
+void init_3d(int a[N][2][2]);
+
+/*Test for OpenMP 4.5 target data map with array section [lower:length]*/
+int test_lower_length_1d() {
+  /* array sections of the form a[lower:length] */
+  puts("test_lower_length_1d");
+
+  int errors = 0, isHost = 0;
+
+  int a1d[N];
+  init_1d(a1d);
+
+#pragma omp target data map(from : a1d[1 : N - 2])
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 1; i < N - 1; ++i)
+        a1d[i] = 1;
+    } /*end target*/
+  } /*end target data*/
+
+  /* checking errors */
+  for (int i = 0; i < N; ++i) {
+    if (i == 0 || i == N - 1)
+      errors += a1d[i] == 0 ? 0 : 1;
+    else
+      errors += a1d[i] == 1 ? 0 : 1;
+  }
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data map with array 2d section [lower:length]*/
+int test_lower_length_2d() {
+  /* array sections of the form a[lower:length]*/
+  puts("test_lower_length_2d");
+
+  int errors = 0, isHost = 0;
+
+  // stack
+  int a2d[N][2];
+  init_2d(a2d);
+
+  /* OpenMP API v 4.5 Nov 2015. 2.15.5.1 map Clause, page 218 line 17: If a list item is an array section, it must specify contiguous storage.*/
+#pragma omp target data map(from : a2d[1 : N - 2][0 : 2])
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 1; i < N - 1; ++i) {
+        a2d[i][0] = 1;
+        a2d[i][1] = 1;
+      }
+    } /*end target*/
+  } /*end target data*/
+
+  /* checking errors */
+  for (int i = 0; i < N; ++i) {
+    if (i == 0 || i == N - 1)
+      errors += a2d[i][0] == 0 && a2d[i][1] == 0 ? 0 : 1;
+    else
+      errors += a2d[i][0] == 1 && a2d[i][1] == 1 ? 0 : 1;
+  }
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data map with array 3d section [lower:length]*/
+int test_lower_length_3d() {
+  /* array sections of the form a[lower:length] */
+  puts("test_lower_length_3d");
+  /* If a list item is an array section, it must specify contiguous storage. */
+
+  int errors = 0, isHost = 0;
+
+  // stack
+  int a3d[N][2][2];
+  init_3d(a3d);
+  int a3d2[N][2][2];
+  init_3d(a3d2);
+
+  /* OpenMP API v 4.5 Nov 2015. 2.15.5.1 map Clause, page 218 line 17: If a list item is an array section, it must specify contiguous storage.*/
+#pragma omp target data map(from : a3d[1 : N - 2][0 : 2][0 : 2])  \
+    map(from: a3d2[0 : N][0 : 2][0 : 2])
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 0; i < N; ++i) {
+        for (int j = 0; j < 2; ++j) {
+          if (i > 0 && i < N - 1) {
+            a3d[i][j][0] = 1;
+            a3d[i][j][1] = 1;
+          }
+          a3d2[i][j][0] = 1;
+          a3d2[i][j][1] = 1;
+        }
+      }
+    } /*end target*/
+  } /*end target data*/
+
+  /* checking errors*/
+  for (int i = 0; i < N; ++i) {
+    for (int j = 0; j < 2; ++j) {
+      /* a3d */
+      if (i == 0 || i == N - 1)
+        errors += a3d[i][j][0] == 0 && a3d[i][j][1] == 0 ? 0 : 1;
+      else
+        errors += a3d[i][j][0] == 1 && a3d[i][j][1] == 1 ? 0 : 1;
+      /* a3d2 */
+      errors += a3d2[i][j][0] == 1 && a3d2[i][j][1] == 1 ? 0 : 1;
+    }
+  }
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data map with array 1d section [:length]*/
+int test_length_1d() {
+  /* array sections of the form a[:length]*/
+  puts("test_length_1d");
+
+  int errors = 0, isHost = 0;
+
+  int a1d[N];
+  init_1d(a1d);
+
+  /* OpenMP API - V4.5 Nov2015. 2.4. Array sections, page 45 line 14: When the lower-bound is absent it defaults to 0.*/
+#pragma omp target data map(from : a1d[ : N - 2]) 
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 0; i < N - 2; ++i)
+        a1d[i] = 1;
+    } /* end target*/
+  } /*end target data*/
+
+  /* checking errors */
+  for (int i = 0; i < N - 2; ++i)
+    errors += a1d[i] == 1 ? 0 : 1;
+  /* N-2 */
+  errors += a1d[N - 2] == 0 ? 0 : 1;
+  /* N-1 */
+  errors += a1d[N - 1] == 0 ? 0 : 1;
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data map with array 2d section [:length]*/
+int test_length_2d() {
+  /* array sections of the form a[:length] */
+  puts("test_length_2d");
+
+  int errors = 0, isHost = 0;
+
+  int a2d[N][2];
+  init_2d(a2d);
+
+  /* OpenMP API - V4.5 Nov2015. 2.4. Array sections, page 45 line 14: When the lower-bound is absent it defaults to 0.*/
+  /* OpenMP API v 4.5 Nov 2015. 2.15.5.1 map Clause, page 218 line 17: If a list item is an array section, it must specify contiguous storage.*/
+#pragma omp target data map(from : a2d[ : N - 2][ : 2])
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 0; i < N - 2; ++i) {
+        a2d[i][0] = 1;
+        a2d[i][1] = 1;
+      }
+    } /*end target*/
+  } /*end target data*/
+
+  /* checking errors*/
+  for (int i = 0; i < N - 2; ++i)
+    errors += a2d[i][0] == 1 && a2d[i][1] == 1 ? 0 : 1;
+  errors += a2d[N - 2][0] == 0 && a2d[N - 2][1] == 0 ? 0 : 1;
+  errors += a2d[N - 1][0] == 0 && a2d[N - 1][1] == 0 ? 0 : 1;
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data map with array 3d section [:length]*/
+int test_length_3d() {
+  /* array sections of the form a[:length] */
+  puts("test_length_3d");
+
+  int errors = 0, isHost = 0;
+
+  int a3d[N][2][2];
+  init_3d(a3d);
+  int a3d2[N][2][2];
+  init_3d(a3d2);
+
+  /* OpenMP API - V4.5 Nov2015. 2.4. Array sections, page 45 line 14: When the lower-bound is absent it defaults to 0.*/
+  /* OpenMP API v 4.5 Nov 2015. 2.15.5.1 map Clause, page 218 line 17: If a list item is an array section, it must specify contiguous storage.*/
+#pragma omp target data map(from : a3d[ : N - 2][ : 2][ : 2])   \
+    map(from: a3d2[ : N][ : 2][ : 2])
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 0; i < N; ++i) {
+        for (int j = 0; j < 2; ++j) {
+          if (i < N - 2) {
+            a3d[i][j][0] = 1;
+            a3d[i][j][1] = 1;
+          }
+          a3d2[i][j][0] = 1;
+          a3d2[i][j][1] = 1;
+        }
+      }
+    } /*end target*/
+  } /*end target data*/
+
+  /* checking errors*/
+  for (int i = 0; i < N; ++i) {
+    for (int j = 0; j < 2; ++j) {
+      if (i >= N - 2)
+        errors += a3d[i][j][0] == 0 && a3d[i][j][1] == 0 ? 0 : 1;
+      else
+        errors += a3d[i][j][0] == 1 && a3d[i][j][1] == 1 ? 0 : 1;
+      /* a3d2 */
+      errors += a3d2[i][j][0] == 1 && a3d2[i][j][1] == 1 ? 0 : 1;
+    }
+  }
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data map with array 1d section [lower:]*/
+int test_lower_1d() {
+  /* array sections of the form a[lower:] */
+  puts("test_lower_1d");
+
+  int errors = 0, isHost = 0;
+
+  int a1d[N];
+  init_1d(a1d);
+
+  /* OpenMP API - V4.5 Nov2015. 2.4. Array sections, page 45 line 13: When the length is absent, it defaults to the size of the array dimension minus the lower-bound.*/
+#pragma omp target data map(from : a1d[1 : ])
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 1; i < N; ++i)
+        a1d[i] = 1;
+    } /* end target*/
+  } /* end target data*/
+
+  /* checking errors */
+  for (int i = 0; i < N; ++i) {
+    if (i == 0)
+      errors += a1d[i] == 0 ? 0 : 1;
+    else
+      errors += a1d[i] == 1 ? 0 : 1;
+  }
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data map with array 2d section [lower:]*/
+int test_lower_2d() {
+  /* array sections of the form a[lower:] */
+  puts("test_lower_2d");
+
+  int errors = 0, isHost = 0;
+
+  int a2d[N][2];
+  init_2d(a2d);
+
+  /* OpenMP API - V4.5 Nov2015. 2.4. Array sections, page 45 line 13: When the length is absent, it defaults to the size of the array dimension minus the lower-bound.*/
+  /* OpenMP API v 4.5 Nov 2015. 2.15.5.1 map Clause, page 218 line 17: If a list item is an array section, it must specify contiguous storage.*/
+#pragma omp target data map(from : a2d[1 : ][0 : ])
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 1; i < N; ++i) {
+        a2d[i][0] = 1;
+        a2d[i][1] = 1;
+      }
+    } /* end target*/
+  } /* end target data*/
+
+  /* checking errors */
+  for (int i = 0; i < N; ++i) {
+    if (i == 0)
+      errors += a2d[i][0] == 0 && a2d[i][1] == 0 ? 0 : 1;
+    else
+      errors += a2d[i][0] == 1 && a2d[i][1] == 1 ? 0 : 1;
+  }
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data map with array 3d section [lower:]*/
+int test_lower_3d() {
+  /* array sections of the form a[lower:] */
+  puts("test_lower_3d");
+
+  int errors = 0, isHost = 0;
+
+  int a3d[N][2][2];
+  init_3d(a3d);
+  int a3d2[N][2][2];
+  init_3d(a3d2);
+
+  /* OpenMP API - V4.5 Nov2015. 2.4. Array sections, page 45 line 13: When the length is absent, it defaults to the size of the array dimension minus the lower-bound.*/
+  /* OpenMP API v 4.5 Nov 2015. 2.15.5.1 map Clause, page 218 line 17: If a list item is an array section, it must specify contiguous storage.*/
+#pragma omp target data map(from : a3d[1 : ][0 : ][0 : ])   \
+    map(from : a3d2[0 : ][0 : ][0 : ])
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 0; i < N; ++i) {
+        for (int j = 0; j < 2; ++j) {
+          if (i > 0) {
+            a3d[i][j][0] = 1;
+            a3d[i][j][1] = 1;
+          }
+          a3d2[i][j][0] = 1;
+          a3d2[i][j][1] = 1;
+        }
+      }
+    } /* end target*/
+  } /* end target data*/
+
+  /* checking errors */
+  for (int i = 0; i < N; ++i) {
+    for (int j = 0; j < 2; ++j) {
+      /* a3d*/
+      if (i == 0)
+        errors += a3d[i][j][0] == 0 && a3d[i][j][1] == 0 ? 0 : 1;
+      else
+        errors += a3d[i][j][0] == 1 && a3d[i][j][1] == 1 ? 0 : 1;
+      /* a3d2*/
+      errors += a3d2[i][j][0] == 1 && a3d2[i][j][1] == 1 ? 0 : 1;
+    }
+  }
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}
+
+int main() {
+
+  int errors = 0;
+
+  errors += test_lower_length_1d();
+  errors += test_lower_length_2d();
+  errors += test_lower_length_3d();
+  errors += test_length_1d();
+  errors += test_length_2d();
+  errors += test_length_3d();
+  errors += test_lower_1d();
+  errors += test_lower_2d();
+  errors += test_lower_3d();
+
+  return errors;
+}
+
+void init_1d(int* a) {
+  for (int i = 0; i < N; ++i)
+    a[i] = 0;
+}
+
+void init_2d(int a[N][2]) {
+  for (int i = 0; i < N; ++i) {
+    a[i][0] = 0;
+    a[i][1] = 0;
+  }
+}
+
+void init_3d(int a[N][2][2]) {
+  for (int i = 0; i < N; ++i)
+    for (int j = 0; j < 2; ++j) {
+      a[i][j][0] = 0;
+      a[i][j][1] = 0;
+    }
+}

libomptarget/test/offloading/target_data/test_target_data_map_classes.cpp

+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+#include <iostream>
+#include <omp.h>
+
+using namespace std;
+
+#define N 1000
+
+class A {
+
+public:
+  /* breaks encapsulation*/
+  int *h_array;
+  int size;
+  int sum;
+
+  A(int *array, const int s) : h_array(array), size(s) { sum = 0; }
+
+  /* TODO: Add virtual once supported */
+  ~A() { h_array = NULL; }
+};
+
+/*Test for OpenMP 4.5 target data mapping objects in the heap*/
+int test_map_tofrom_class_heap() {
+
+  cout << "test_map_tofrom_class_heap" << endl;
+
+  int sum = 0, errors = 0, isHost = 0;
+
+  int *array = new int[N];
+  A *obj = new A(array, N);
+
+  /* mapping an object + array: it is shallow copy thus pointers are not translated automatically*/
+#pragma omp target data map(from : array[0 : N]) map(tofrom : obj[0 : 1])
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      
+      /* assign device array ptr to device obj */
+      int *tmp_h_array = obj->h_array;
+      obj->h_array = array;
+      int tmp = 0;
+      for (int i = 0; i < N; ++i) {
+        obj->h_array[i] = 1;
+        tmp += 1;
+      }
+      /* swap array device ptr to host ptr */
+      obj->h_array = tmp_h_array;
+
+      obj->sum = tmp;
+    } /* end target*/
+  } /* end target data*/
+
+  /* checking results */
+  for (int i = 0; i < N; ++i)
+    sum += obj->h_array[i];
+
+  errors = (N != sum) || (N != obj->sum);
+  if (!errors)
+    cout << "Test passed on " << (isHost ? "host" : "device") << endl;
+  else
+    cout << "Test failed on " << (isHost ? "host" : "device") << ": sum=" << sum << ", N=" << N << endl;
+
+  delete obj;
+  delete[] array;
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data mapping objects on the stack*/
+int test_map_tofrom_class_stack() {
+
+  cout << "test_map_tofrom_class_stack" << endl;
+
+  int sum = 0, errors = 0, isHost = 0;
+
+  int array[N];
+  A obj(array, N);
+
+  /* mapping an object + array: it is shallow copy thus pointers are not translated automatically*/
+#pragma omp target data map(from : array[0 : N]) map(tofrom : obj)
+  {
+#pragma omp target map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      
+      /* assign device array ptr to device obj */
+      int *tmp_h_array = obj.h_array;
+      obj.h_array = array;
+      int tmp = 0;
+      for (int i = 0; i < N; ++i) {
+        obj.h_array[i] = 1;
+        tmp += 1;
+      }
+      /* swap array device ptr to host ptr */
+      obj.h_array = tmp_h_array;
+
+      obj.sum = tmp;
+    } /* end target*/
+  } /* end target data*/
+
+  /* checking results */
+  for (int i = 0; i < N; ++i)
+    sum += obj.h_array[i];
+
+  errors = (N != sum) || (N != obj.sum);
+  if (!errors)
+    cout << "Test passed on " << (isHost ? "host" : "device") << endl;
+  else
+    cout << "Test failed on " << (isHost ? "host" : "device") << ": sum=" << sum << ", N=" << N << endl;
+
+  return errors;
+}
+
+int main() {
+
+  int errors = 0;
+
+  errors += test_map_tofrom_class_heap();
+  errors += test_map_tofrom_class_stack();
+
+  return errors;
+}

libomptarget/test/offloading/target_data/test_target_data_map_devices.c

+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 1000
+
+/*Test for OpenMP 4.5 target data to multiple devices using API*/
+int test_map_set_default_dev() {
+
+  puts("test_map_set_default_dev");
+
+  /* Get number of devices */
+  int num_dev = omp_get_num_devices();
+  printf("num_devices: %d\n", num_dev);
+
+  int def_dev = omp_get_default_device();
+  printf("initial device: %d\n", omp_get_initial_device());
+  printf("default device: %d\n", def_dev);
+
+  int sum[num_dev], errors = 0, isHost = 0;
+  int* h_matrix = (int*) malloc(num_dev * N * sizeof(int));
+
+  for (int dev = 0; dev < num_dev; ++dev) {
+    omp_set_default_device(dev);
+#pragma omp target data map(from : h_matrix[dev*N : N]) map(tofrom: errors)
+    {
+      errors = dev == omp_get_default_device();
+#pragma omp target map(from : h_matrix[dev*N : N]) map(tofrom: isHost)
+      {
+        isHost = omp_is_initial_device();
+        for (int i = 0; i < N; ++i)
+          h_matrix[dev*N + i] = dev;
+      } /* end target*/
+    } /* end target data*/
+  }
+
+  /* checking results */
+  errors = 0;
+  for (int dev = 0; dev < num_dev; ++dev) {
+    sum[dev] = h_matrix[dev*N + 0];
+    for (int i = 1; i < N; ++i)
+      sum[dev] += h_matrix[dev*N + i];
+    errors |= (dev * N != sum[dev]);
+  }
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s: num_devices = %d\n", (isHost ? "host" : "device"), num_dev);
+
+  omp_set_default_device(def_dev);
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data to multiple devices using directives*/
+int test_map_device() {
+
+  puts("test_map_device");
+
+  /* Get number of devices */
+  int num_dev = omp_get_num_devices();
+  printf("num_devices: %d\n", num_dev);
+
+  printf("initial device: %d\n", omp_get_initial_device());
+  printf("default device: %d\n", omp_get_default_device());
+
+  int sum[num_dev], errors = 0, isHost = 0;
+  int* h_matrix = (int*) malloc(num_dev * N * sizeof(int));
+
+  for (int dev = 0; dev < num_dev; ++dev) {
+#pragma omp target data map(from : h_matrix[dev*N : N]) device(dev)
+    {
+#pragma omp target map(from : h_matrix[dev*N : N]) device(dev) map(tofrom: isHost)
+      {
+        isHost = omp_is_initial_device();
+        for (int i = 0; i < N; ++i)
+          h_matrix[dev*N + i] = dev;
+      } /* end target*/
+    } /* end target data*/
+  }
+
+  /* checking results */
+  errors = 0;
+  for (int dev = 0; dev < num_dev; ++dev) {
+    sum[dev] = h_matrix[dev*N + 0];
+    for (int i = 1; i < N; ++i)
+      sum[dev] += h_matrix[dev*N + i];
+    errors |= (dev * N != sum[dev]);
+  }
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s: num_devices = %d\n", (isHost ? "host" : "device"), num_dev);
+
+  return errors;
+}
+
+/*Test for OpenMP 4.5 target data to more than available devices using API*/
+int test_more_devices() {
+
+  puts("test_more_devices");
+
+  /* Get number of devices */
+  int real_num_dev = omp_get_num_devices();
+  int num_dev = real_num_dev + 2;
+  printf("num_devices: %d, real_num_devices: %d\n", num_dev, real_num_dev);
+
+  int def_dev = omp_get_default_device();
+  printf("initial device: %d\n", omp_get_initial_device());
+  printf("default device: %d\n", def_dev);
+
+  int sum[num_dev], errors = 0, extra_on_host = 0, isHost = 0;
+  int* h_matrix = (int*) malloc(num_dev * N * sizeof(int));
+
+  /* omp_set_default_device is implementation dependent */
+  for (int dev = 0; dev < num_dev; ++dev) {
+    omp_set_default_device(dev);
+    int dev_tmp = omp_get_default_device();
+#pragma omp target data map(from : h_matrix[dev_tmp*N : N])
+    {
+      errors = dev == dev_tmp;
+#pragma omp target map(from : h_matrix[dev_tmp*N : N])    \
+        map(tofrom : extra_on_host) map(tofrom: isHost)
+      {
+        isHost = omp_is_initial_device();
+        if (!isHost) {
+          for (int i = 0; i < N; ++i)
+            h_matrix[dev_tmp*N + i] = dev;
+        } else {
+          /* allows implementations that map extra devices to the host */
+          extra_on_host = 1;
+          for (int i = 0; i < N; ++i)
+            h_matrix[dev_tmp*N + i] = dev;
+        }
+      } /* end target*/
+    } /* end target data*/
+  }
+
+  /* checking results */
+  errors = 0;
+  for (int dev = 0; dev < num_dev; ++dev) {
+    sum[dev] = h_matrix[dev*N + 0];
+    for (int i = 1; i < N; ++i)
+      sum[dev] += h_matrix[dev*N + i];
+    errors |= (dev * N != sum[dev]);
+    if (errors)
+      printf("Error at dev=%d\n", dev);
+  }
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s: num_devices = %d, real_num_dev = %d, extra_on_host = %d\n", (isHost ? "host" : "device"), num_dev, real_num_dev, extra_on_host);
+
+  omp_set_default_device(def_dev);
+
+  return errors;
+}
+
+int main() {
+
+  int errors = 0;
+
+  errors += test_map_set_default_dev();
+  errors += test_map_device();
+  errors += test_more_devices();
+
+  return errors;
+}

libomptarget/test/offloading/target_data/test_target_data_use_device_ptr.c

+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+#include <omp.h>
+#include <stdio.h>
+
+/*Test for OpenMP 4.5 target data with use_device_ptr*/
+int main() {
+  int errors = 0, len = 10000, isHost = 0;
+  int *array_device = NULL;
+  int *array_host = NULL;
+
+  array_device = (int *)malloc(len * sizeof(int));
+  array_host = (int *)malloc(len * sizeof(int));
+
+  for (int i = 0; i < len; ++i)
+    array_host[i] = i;
+
+#pragma omp target data map(tofrom : array_device[0 : len])   \
+    use_device_ptr(array_device)
+  {
+#pragma omp target is_device_ptr(array_device) map(tofrom : array_host[0 : len]) map(tofrom: isHost)
+    {
+      isHost = omp_is_initial_device();
+      for (int i = 0; i < len; ++i) {
+        array_device[i] = i;
+        array_host[i] += array_device[i];
+      } 
+    } /* end target*/
+  } /* end target data*/
+
+  /* checking results*/
+  for (int i = 0; i < len; ++i) {
+    if (array_host[i] != 2 * i)
+      errors = 1;
+  }
+
+  free(array_device);
+  free(array_host);
+
+  if (!errors)
+    printf("Test passed on %s\n", (isHost ? "host" : "device"));
+  else
+    printf("Test failed on %s\n", (isHost ? "host" : "device"));
+
+  return errors;
+}