diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_heap.F90 b/tests/5.0/requires/test_requires_unified_shared_memory_heap.F90 index 3c4638c79..fef75a9ca 100644 --- a/tests/5.0/requires/test_requires_unified_shared_memory_heap.F90 +++ b/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 - - diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_heap.c b/tests/5.0/requires/test_requires_unified_shared_memory_heap.c index c96935245..e6ceeb1c5 100644 --- a/tests/5.0/requires/test_requires_unified_shared_memory_heap.c +++ b/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() { diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_heap_is_device_ptr.c b/tests/5.0/requires/test_requires_unified_shared_memory_heap_is_device_ptr.c index 8411ccc9b..1bd8dafd7 100644 --- a/tests/5.0/requires/test_requires_unified_shared_memory_heap_is_device_ptr.c +++ b/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() { diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_heap_map.F90 b/tests/5.0/requires/test_requires_unified_shared_memory_heap_map.F90 index 1857f77e8..d13b1ec92 100644 --- a/tests/5.0/requires/test_requires_unified_shared_memory_heap_map.F90 +++ b/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 diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_heap_map.c b/tests/5.0/requires/test_requires_unified_shared_memory_heap_map.c index 150074392..c9d35e652 100644 --- a/tests/5.0/requires/test_requires_unified_shared_memory_heap_map.c +++ b/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() { diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_heap_map_stack_map.F90 b/tests/5.0/requires/test_requires_unified_shared_memory_heap_map_stack_map.F90 new file mode 100644 index 000000000..d3a00f4a4 --- /dev/null +++ b/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 + + diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_heap_map_stack_map.c b/tests/5.0/requires/test_requires_unified_shared_memory_heap_map_stack_map.c new file mode 100644 index 000000000..f54a41502 --- /dev/null +++ b/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 +#include +#include +#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); +} diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_heap_stack_map.F90 b/tests/5.0/requires/test_requires_unified_shared_memory_heap_stack_map.F90 new file mode 100644 index 000000000..b19c60db0 --- /dev/null +++ b/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 diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_heap_stack_map.c b/tests/5.0/requires/test_requires_unified_shared_memory_heap_stack_map.c new file mode 100644 index 000000000..bd983cb75 --- /dev/null +++ b/tests/5.0/requires/test_requires_unified_shared_memory_heap_stack_map.c @@ -0,0 +1,74 @@ +//===---test_requires_unified_shared_memory_heap_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. +// +// It uses the default mapping of pointers to access the array. +// +// Variant of test_requires_unified_shared_memory_heap.c +// which uses an explicitly mapped stack variable for checking. +// +////===----------------------------------------------------------------------===// +#include +#include +#include +#include "ompvv.h" + +#define N 1024 + +#pragma omp requires unified_shared_memory + +int unified_shared_memory_heap() { + OMPVV_INFOMSG("Unified shared memory testing - Array on heap"); + int errors = 0; + + int *anArray; + int anArrayCopy[N]; + + anArray = (int*)malloc(sizeof(int)*N); + + for (int i = 0; i < N; i++) { + anArray[i] = i; + anArrayCopy[i] = 0; + } + // Modify in the device +#pragma omp target + { + 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(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()); + + OMPVV_REPORT_AND_RETURN(errors); +}