Index: libomptarget/test/offloading/target_data/test_target_data_if.c
===================================================================
--- /dev/null
+++ libomptarget/test/offloading/target_data/test_target_data_if.c
@@ -0,0 +1,55 @@
+// 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;
+}
Index: libomptarget/test/offloading/target_data/test_target_data_map.c
===================================================================
--- /dev/null
+++ libomptarget/test/offloading/target_data/test_target_data_map.c
@@ -0,0 +1,268 @@
+// 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;
+}
Index: libomptarget/test/offloading/target_data/test_target_data_map_array_sections.c
===================================================================
--- /dev/null
+++ libomptarget/test/offloading/target_data/test_target_data_map_array_sections.c
@@ -0,0 +1,441 @@
+// 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;
+    }
+}
Index: libomptarget/test/offloading/target_data/test_target_data_map_classes.cpp
===================================================================
--- /dev/null
+++ libomptarget/test/offloading/target_data/test_target_data_map_classes.cpp
@@ -0,0 +1,128 @@
+// 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;
+}
Index: libomptarget/test/offloading/target_data/test_target_data_map_devices.c
===================================================================
--- /dev/null
+++ libomptarget/test/offloading/target_data/test_target_data_map_devices.c
@@ -0,0 +1,177 @@
+// 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;
+}
Index: libomptarget/test/offloading/target_data/test_target_data_use_device_ptr.c
===================================================================
--- /dev/null
+++ libomptarget/test/offloading/target_data/test_target_data_use_device_ptr.c
@@ -0,0 +1,49 @@
+// 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;
+}