diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_allocatable.F90 b/tests/5.0/requires/test_requires_unified_shared_memory_allocatable.F90 new file mode 100644 index 000000000..f1a07607f --- /dev/null +++ b/tests/5.0/requires/test_requires_unified_shared_memory_allocatable.F90 @@ -0,0 +1,76 @@ +!===--- test_requires_unified_shared_memory_allocatable.F90 ------------------------------===// +! +! OpenMP API Version 5.0 Nov 2018 +! +! This test checks for unified shared memory support of an allocatable array. +! +!===------------------------------------------------------------------------------===// + +#define OMPVV_MODULE_REQUIRES_LINE !$omp requires unified_shared_memory +#include "ompvv.F90" + +#define N 1024 + +PROGRAM test_requires_unified_shared_memory_allocatable + 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_allocatable() .NE. 0) + OMPVV_REPORT_AND_RETURN() + +CONTAINS + INTEGER FUNCTION unified_shared_memory_allocatable() + INTEGER:: errors, i + INTEGER, ALLOCATABLE:: anArray(:) + INTEGER, ALLOCATABLE:: anArrayCopy(:) + + OMPVV_INFOMSG("Unified shared memory testing - Array on allocatable") + + errors = 0 + + ALLOCATE(anArray(N), anArrayCopy(N)) + + OMPVV_ERROR_IF(.NOT. ALLOCATED(anArray), "Memory was not properly allocated") + OMPVV_ERROR_IF(.NOT. ALLOCATED(anArrayCopy), "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 + 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, anArrayCopy) + unified_shared_memory_allocatable = errors + END FUNCTION unified_shared_memory_allocatable +END PROGRAM test_requires_unified_shared_memory_allocatable diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_allocatable_map.F90 b/tests/5.0/requires/test_requires_unified_shared_memory_allocatable_map.F90 new file mode 100644 index 000000000..3fba7db85 --- /dev/null +++ b/tests/5.0/requires/test_requires_unified_shared_memory_allocatable_map.F90 @@ -0,0 +1,73 @@ +!===--- test_requires_unified_shared_memory_allocatable_map.F90 --------------------------===// +! +! OpenMP API Version 5.0 Nov 2018 +! +! This test checks for unified shared memory of an array that is allocated +! using allocatable 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. +! +!===------------------------------------------------------------------------------===// + +#define OMPVV_MODULE_REQUIRES_LINE !$omp requires unified_shared_memory +#include "ompvv.F90" + +#define N 1024 + +PROGRAM test_requires_unified_shared_memory_allocatable_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_allocatable_map() .NE. 0) + + OMPVV_REPORT_AND_RETURN() + +CONTAINS + INTEGER FUNCTION unified_shared_memory_allocatable_map() + INTEGER:: errors, i + INTEGER, ALLOCATABLE:: anArray(:) + INTEGER:: ERR + + OMPVV_INFOMSG("Unified shared memory testing - Array on allocatable") + + 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 + END DO + + ! Modify in the device + !$omp target map(anArray) + DO i = 1, N + anArray(i) = anArray(i) + 10 + END DO + !$omp end target + + DO i = 1, N + OMPVV_TEST_AND_SET_VERBOSE(errors, anArray(i) .NE. (i + 10)) + IF (errors .NE. 0) THEN + exit + END IF + END DO + + DEALLOCATE(anArray) + unified_shared_memory_allocatable_map = errors + END FUNCTION unified_shared_memory_allocatable_map +END PROGRAM test_requires_unified_shared_memory_allocatable_map diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_malloc.c b/tests/5.0/requires/test_requires_unified_shared_memory_malloc.c new file mode 100644 index 000000000..a2d20d9a6 --- /dev/null +++ b/tests/5.0/requires/test_requires_unified_shared_memory_malloc.c @@ -0,0 +1,84 @@ +//===---test_requires_unified_shared_memory_malloc.c --------------------------===// +// +// OpenMP API Version 5.0 Nov 2018 +// +// This test checks for unified shared memory of an array that is allocated using +// malloc on host is accessed from host and device with the same pointer. +// +// It uses the default mapping of pointers to access the array. +// +////===----------------------------------------------------------------------===// +#include +#include +#include +#include "ompvv.h" + +#define N 1024 + +#pragma omp requires unified_shared_memory + +int unified_shared_memory_malloc() { + OMPVV_INFOMSG("Unified shared memory testing - Malloced Array"); + int errors = 0; + + int *anArray; + int *anArrayCopy; + + anArray = (int*)malloc(sizeof(int)*N); + anArrayCopy = (int*)malloc(sizeof(int)*N); + + if( anArray == NULL ) { + OMPVV_ERROR("Memory for anArray was not allocated"); + OMPVV_RETURN(1); + } + + if( anArrayCopy == NULL ) { + OMPVV_ERROR("Memory for anArrayCopy was not allocated"); + OMPVV_RETURN(1); + } + + 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 + { + 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); + free(anArrayCopy); + 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_malloc()); + + OMPVV_REPORT_AND_RETURN(errors); +} diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_malloc_is_device_ptr.c b/tests/5.0/requires/test_requires_unified_shared_memory_malloc_is_device_ptr.c new file mode 100644 index 000000000..873d14cc6 --- /dev/null +++ b/tests/5.0/requires/test_requires_unified_shared_memory_malloc_is_device_ptr.c @@ -0,0 +1,84 @@ +//===---test_requires_unified_shared_memory_malloc_is_device_ptr.c -----------===// +// +// OpenMP API Version 5.0 Nov 2018 +// +// This test Checks for unified shared memory of an array that is allocated using +// malloc and that is accessed from host and device with the same pointer +// +// To guarantee the use of the device pointer, we use the is_device_ptr clause +// +////===----------------------------------------------------------------------===// +#include +#include +#include +#include "ompvv.h" + +#define N 1024 + +#pragma omp requires unified_shared_memory + +int unified_shared_memory_malloc() { + OMPVV_INFOMSG("Unified shared memory testing - Array on malloc"); + int errors = 0; + + int *anArray; + int *anArrayCopy; + + anArray = (int*)malloc(sizeof(int)*N); + anArrayCopy = (int*)malloc(sizeof(int)*N); + + if( anArray == NULL ) { + OMPVV_ERROR("Memory for anArray was not allocated"); + OMPVV_RETURN(1); + } + + if( anArrayCopy == NULL ) { + OMPVV_ERROR("Memory for anArrayCopy was not allocated"); + OMPVV_RETURN(1); + } + + for (int i = 0; i < N; i++) { + anArray[i] = i; + anArrayCopy[i] = 0; + } + // Modify in the device +#pragma omp target is_device_ptr(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 is_device_ptr(anArray) + { + 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); + free(anArrayCopy); + 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_malloc()); + + OMPVV_REPORT_AND_RETURN(errors); +} + diff --git a/tests/5.0/requires/test_requires_unified_shared_memory_malloc_map.c b/tests/5.0/requires/test_requires_unified_shared_memory_malloc_map.c new file mode 100644 index 000000000..cb629989e --- /dev/null +++ b/tests/5.0/requires/test_requires_unified_shared_memory_malloc_map.c @@ -0,0 +1,87 @@ +//===---test_requires_unified_shared_memory_malloc_map.c ----------------------===// +// +// OpenMP API Version 5.0 Nov 2018 +// +// This test checks for unified shared memory of an array that is allocated using +// malloc 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. +// +////===----------------------------------------------------------------------===// +#include +#include +#include +#include "ompvv.h" + +#define N 1024 + +#pragma omp requires unified_shared_memory + +int unified_shared_memory_malloc_map() { + OMPVV_INFOMSG("Unified shared memory testing - Array on malloc"); + int errors = 0; + + int *anArray; + int *anArrayCopy; + + anArray = (int*)malloc(sizeof(int)*N); + anArrayCopy = (int*)malloc(sizeof(int)*N); + + if( anArray == NULL ) { + OMPVV_ERROR("Memory for anArray was not allocated"); + OMPVV_RETURN(1); + } + + if( anArrayCopy == NULL ) { + OMPVV_ERROR("Memory for anArrayCopy was not 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) + { + 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); + free(anArrayCopy); + 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_malloc_map()); + + OMPVV_REPORT_AND_RETURN(errors); +} + +