Avoid stack allocation for test_requires_unified_shared_memory_heap*

tests/5.0/requires/test_requires_unified_shared_memory_heap.F90

@@ -32,13 +32,13 @@ PROGRAM test_requires_unified_shared_memory_heap
   INTEGER FUNCTION unified_shared_memory_heap()
     INTEGER:: errors, i
     INTEGER, ALLOCATABLE:: anArray(:)
-    INTEGER, DIMENSION(N):: anArrayCopy
+    INTEGER, ALLOCATABLE:: anArrayCopy(:)
 
     OMPVV_INFOMSG("Unified shared memory testing - Array on heap")
 
     errors = 0
 
-    ALLOCATE(anArray(N))
+    ALLOCATE(anArray(N), anArrayCopy(N))
 
     OMPVV_ERROR_IF(.NOT. ALLOCATED(anArray), "Memory was not properly allocated")
 
@@ -74,9 +74,7 @@ INTEGER FUNCTION unified_shared_memory_heap()
       END IF
     END DO
 
-    DEALLOCATE(anArray)
+    DEALLOCATE(anArray, anArrayCopy)
     unified_shared_memory_heap = errors
   END FUNCTION unified_shared_memory_heap
 END PROGRAM test_requires_unified_shared_memory_heap
-
-

tests/5.0/requires/test_requires_unified_shared_memory_heap.c

@@ -22,9 +22,10 @@ int unified_shared_memory_heap() {
   int errors = 0;
 
   int *anArray;
-  int anArrayCopy[N];
+  int *anArrayCopy;
 
   anArray = (int*)malloc(sizeof(int)*N);
+  anArrayCopy = (int*)malloc(sizeof(int)*N);
 
   for (int i = 0; i < N; i++) {
     anArray[i] = i;
@@ -57,6 +58,7 @@ int unified_shared_memory_heap() {
   }
 
   free(anArray);
+  free(anArrayCopy);
   return errors;
 }
 int main() {

tests/5.0/requires/test_requires_unified_shared_memory_heap_is_device_ptr.c

@@ -22,9 +22,10 @@ int unified_shared_memory_heap() {
   int errors = 0;
 
   int *anArray;
-  int anArrayCopy[N];
+  int *anArrayCopy;
 
   anArray = (int*)malloc(sizeof(int)*N);
+  anArrayCopy = (int*)malloc(sizeof(int)*N);
 
   for (int i = 0; i < N; i++) {
     anArray[i] = i;
@@ -57,6 +58,7 @@ int unified_shared_memory_heap() {
   }
 
   free(anArray);
+  free(anArrayCopy);
   return errors;
 }
 int main() {

tests/5.0/requires/test_requires_unified_shared_memory_heap_map.F90

@@ -36,16 +36,16 @@ PROGRAM test_requires_unified_shared_memory_heap_map
   INTEGER FUNCTION unified_shared_memory_heap_map()
     INTEGER:: errors, i
     INTEGER, ALLOCATABLE:: anArray(:)
-    INTEGER, DIMENSION(N):: anArrayCopy
+    INTEGER, ALLOCATABLE:: anArrayCopy(:)
     INTEGER:: ERR
 
     OMPVV_INFOMSG("Unified shared memory testing - Array on heap")
 
     errors = 0
 
-    ALLOCATE(anArray(N), STAT=ERR)
+    ALLOCATE(anArray(N), anArrayCopy(N), STAT=ERR)
 
-    IF( .NOT. ALLOCATED(anArray) ) THEN
+    IF( ERR /= 0 ) THEN
       OMPVV_ERROR("Memory was not properly allocated")
       OMPVV_RETURN(ERR)
     END IF
@@ -82,7 +82,7 @@ INTEGER FUNCTION unified_shared_memory_heap_map()
       END IF
     END DO
 
-    DEALLOCATE(anArray)
+    DEALLOCATE(anArray, anArrayCopy)
     unified_shared_memory_heap_map = errors
   END FUNCTION unified_shared_memory_heap_map
 END PROGRAM test_requires_unified_shared_memory_heap_map

tests/5.0/requires/test_requires_unified_shared_memory_heap_map.c

@@ -23,9 +23,10 @@ int unified_shared_memory_heap_map() {
   int errors = 0;
 
   int *anArray;
-  int anArrayCopy[N];
+  int *anArrayCopy;
 
   anArray = (int*)malloc(sizeof(int)*N);
+  anArrayCopy = (int*)malloc(sizeof(int)*N);
   if( anArray == NULL ) {
     OMPVV_ERROR("Memory was not properly allocated");
     OMPVV_RETURN(1);
@@ -62,6 +63,7 @@ int unified_shared_memory_heap_map() {
   }
 
   free(anArray);
+  free(anArrayCopy);
   return errors;
 }
 int main() {

tests/5.0/requires/test_requires_unified_shared_memory_heap_map_stack_map.F90

@@ -0,0 +1,93 @@
+!===--- test_requires_unified_shared_memory_heap_map_stack_map.F90 ---------------===//
+!
+! OpenMP API Version 5.0 Nov 2018
+!
+! This test checks for unified shared memory of an array that is allocated on 
+! the heap and that is accessed from host and device with the same pointer.
+!
+! The mapping of a pointer under shared memory should allow for the pointer to
+! have the same value as in the host.
+!
+! Variant of test_requires_unified_shared_memory_heap_map.F90
+! which uses an explicitly mapped stack variable for checking.
+!
+!===------------------------------------------------------------------------------===//
+
+#define OMPVV_MODULE_REQUIRES_LINE !$omp requires unified_shared_memory
+#include "ompvv.F90"
+
+#define N 1024
+
+PROGRAM test_requires_unified_shared_memory_heap_map
+   USE iso_fortran_env
+   USE ompvv_lib
+   USE omp_lib
+   implicit none
+   LOGICAL:: isOffloading
+
+!$omp requires unified_shared_memory
+
+  OMPVV_TEST_AND_SET_OFFLOADING(isOffloading)
+
+  OMPVV_WARNING_IF(.NOT. isOffloading, "With no offloading, unified shared memory is guaranteed due to host execution")
+
+   OMPVV_TEST_VERBOSE(unified_shared_memory_heap_map() .NE. 0)
+
+   OMPVV_REPORT_AND_RETURN()
+
+CONTAINS
+  INTEGER FUNCTION unified_shared_memory_heap_map()
+    INTEGER:: errors, i
+    INTEGER, ALLOCATABLE:: anArray(:)
+    INTEGER, DIMENSION(N):: anArrayCopy
+    INTEGER:: ERR
+
+    OMPVV_INFOMSG("Unified shared memory testing - Array on heap")
+
+    errors = 0
+
+    ALLOCATE(anArray(N), STAT=ERR)
+
+    IF( ERR /= 0 ) THEN
+      OMPVV_ERROR("Memory was not properly allocated")
+      OMPVV_RETURN(ERR)
+    END IF
+
+    DO i = 1, N
+      anArray(i) = i
+      anArrayCopy(i) = 0
+    END DO
+
+    ! Modify in the device
+    !$omp target map(anArray)
+    DO i = 1, N
+      anArray(i) = anArray(i) + 10
+    END DO
+    !$omp end target
+
+    ! Modify again on the host
+    DO i = 1, N
+      anArray(i) = anArray(i) + 10
+    END DO
+
+    ! Get the value the device is seeing
+    !$omp target map(anArray, anArrayCopy)
+    DO i = 1, N
+      anArrayCopy(i) = anArray(i)
+    END DO
+    !$omp end target
+
+    DO i = 1, N
+      OMPVV_TEST_AND_SET_VERBOSE(errors, anArray(i) .NE. (i + 20))
+      OMPVV_TEST_AND_SET_VERBOSE(errors, anArrayCopy(i) .NE. (i + 20))
+      IF (errors .NE. 0) THEN
+        exit
+      END IF
+    END DO
+
+    DEALLOCATE(anArray)
+    unified_shared_memory_heap_map = errors
+  END FUNCTION unified_shared_memory_heap_map
+END PROGRAM test_requires_unified_shared_memory_heap_map
+
+

tests/5.0/requires/test_requires_unified_shared_memory_heap_map_stack_map.c

@@ -0,0 +1,79 @@
+//===---test_requires_unified_shared_memory_heap_map_stack_map.c ------------===//
+//
+// OpenMP API Version 5.0 Nov 2018
+// 
+// This test checks for unified shared memory of an array that is allocated on 
+// the heap and that is accessed from host and device with the same pointer.
+//
+// The mapping of a pointer under shared memory should allow for the pointer to
+// have the same value as in the host.
+//
+// Variant of test_requires_unified_shared_memory_heap_map.c
+// which uses an explicitly mapped stack variable for checking.
+//
+////===----------------------------------------------------------------------===//
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include "ompvv.h"
+
+#define N 1024
+
+#pragma omp requires unified_shared_memory
+
+int unified_shared_memory_heap_map() {
+  OMPVV_INFOMSG("Unified shared memory testing - Array on heap");
+  int errors = 0;
+
+  int *anArray;
+  int anArrayCopy[N];
+
+  anArray = (int*)malloc(sizeof(int)*N);
+  if( anArray == NULL ) {
+    OMPVV_ERROR("Memory was not properly allocated");
+    OMPVV_RETURN(1);
+  }
+
+  for (int i = 0; i < N; i++) {
+    anArray[i] = i;
+    anArrayCopy[i] = 0;
+  }
+  // Modify in the device
+#pragma omp target map(anArray)
+  {
+    for (int i = 0; i < N; i++) {
+      anArray[i] += 10;
+    }
+  }
+  // Modify again on the host
+  for (int i = 0; i < N; i++) {
+    anArray[i] += 10;
+  }
+
+  // Get the value the device is seeing
+#pragma omp target map(anArray, anArrayCopy)
+  {
+    for (int i = 0; i < N; i++) {
+      anArrayCopy[i] = anArray[i];
+    }
+  }
+
+  for (int i = 0; i < N; i++) {
+    OMPVV_TEST_AND_SET_VERBOSE(errors, anArray[i] != i + 20);
+    OMPVV_TEST_AND_SET_VERBOSE(errors, anArrayCopy[i] != i + 20);
+    if (errors) break;
+  }
+
+  free(anArray);
+  return errors;
+}
+int main() {
+  int isOffloading;
+  OMPVV_TEST_AND_SET_OFFLOADING(isOffloading);
+  OMPVV_WARNING_IF(!isOffloading, "With no offloading, unified shared memory is guaranteed due to host execution");
+  int errors = 0;
+
+  OMPVV_TEST_AND_SET_VERBOSE(errors, unified_shared_memory_heap_map());
+
+  OMPVV_REPORT_AND_RETURN(errors);
+}

tests/5.0/requires/test_requires_unified_shared_memory_heap_stack_map.F90

@@ -0,0 +1,83 @@
+!===--- test_requires_unified_shared_memory_heap_stack_map.F90 -------------------===//
+!
+! OpenMP API Version 5.0 Nov 2018
+!
+! This test checks for unified shared memory of an array that is allocated on 
+! the heap and that is accessed from host and device with the same pointer.
+!
+! It uses the default mapping of pointers to access the array.
+!
+! Variant of test_requires_unified_shared_memory_heap_map.c
+! which uses an explicitly mapped stack variable for checking.
+!
+!===------------------------------------------------------------------------------===//
+
+#define OMPVV_MODULE_REQUIRES_LINE !$omp requires unified_shared_memory
+#include "ompvv.F90"
+
+#define N 1024
+
+PROGRAM test_requires_unified_shared_memory_heap
+   USE iso_fortran_env
+   USE ompvv_lib
+   USE omp_lib
+   implicit none
+
+!$omp requires unified_shared_memory
+
+   OMPVV_TEST_OFFLOADING
+
+   OMPVV_TEST_VERBOSE(unified_shared_memory_heap() .NE. 0)
+
+   OMPVV_REPORT_AND_RETURN()
+
+CONTAINS
+  INTEGER FUNCTION unified_shared_memory_heap()
+    INTEGER:: errors, i
+    INTEGER, ALLOCATABLE:: anArray(:)
+    INTEGER, DIMENSION(N):: anArrayCopy
+
+    OMPVV_INFOMSG("Unified shared memory testing - Array on heap")
+
+    errors = 0
+
+    ALLOCATE(anArray(N))
+
+    OMPVV_ERROR_IF(.NOT. ALLOCATED(anArray), "Memory was not properly allocated")
+
+    DO i = 1, N
+      anArray(i) = i
+      anArrayCopy(i) = 0
+    END DO
+
+    ! Modify in the device
+    !$omp target
+    DO i = 1, N
+      anArray(i) = anArray(i) + 10
+    END DO
+    !$omp end target
+
+    ! Modify again on the host
+    DO i = 1, N
+      anArray(i) = anArray(i) + 10
+    END DO
+
+    ! Get the value the device is seeing
+    !$omp target map(anArrayCopy)
+    DO i = 1, N
+      anArrayCopy(i) = anArray(i)
+    END DO
+    !$omp end target
+
+    DO i = 1, N
+      OMPVV_TEST_AND_SET_VERBOSE(errors, anArray(i) .NE. (i + 20))
+      OMPVV_TEST_AND_SET_VERBOSE(errors, anArrayCopy(i) .NE. (i + 20))
+      IF (errors .NE. 0) THEN
+        exit
+      END IF
+    END DO
+
+    DEALLOCATE(anArray)
+    unified_shared_memory_heap = errors
+  END FUNCTION unified_shared_memory_heap
+END PROGRAM test_requires_unified_shared_memory_heap