This patch extends omp_target_is_accessible to check the actual device status
for the memory region, on amdgcn and nvptx devices (rather than just checking
if shared memory is enabled).
In both cases, we check the status of each 4k region within the given memory
range (assuming 4k pages should be safe for all the currently supported hosts)
and returns true if all of the pages report accessible.
The testcases have been modified to check that allocations marked accessible
actually are accessible (inaccessibility can't be checked without invoking
memory faults), and to understand that some parts of an array can be accessible
but other parts not (I have observed this intermittently for the stack memory
on amdgcn using the Fortran testcase, which can have the allocation span pages).
There's also new testcases for the various other memory modes, and for managed
memory.
include/ChangeLog:
* cuda/cuda.h (CUpointer_attribute): New enum.
(cuPointerGetAttribute): New prototype.
libgomp/ChangeLog:
PR libgomp/121813
* libgomp-plugin.h (GOMP_OFFLOAD_is_shared_ptr): New prototype.
* libgomp.h (struct gomp_device_descr): Add GOMP_OFFLOAD_is_shared_ptr.
* plugin/cuda-lib.def (cuPointerGetAttribute): New entry.
* plugin/plugin-gcn.c (struct hsa_runtime_fn_info): Add
hsa_amd_svm_attributes_get_fn.
(init_hsa_runtime_functions): Add hsa_amd_svm_attributes_get.
(GOMP_OFFLOAD_is_shared_ptr): New function.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_is_shared_ptr): Likewise.
* target.c (omp_target_is_accessible): Call is_shared_ptr_func.
(gomp_load_plugin_for_device): Add is_shared_ptr.
* testsuite/lib/libgomp.exp (check_effective_target_omp_usm): New.
* testsuite/libgomp.c-c++-common/target-is-accessible-1.c: Rework
to match more details of the GPU implementation.
* testsuite/libgomp.fortran/target-is-accessible-1.f90: Likewise.
* testsuite/libgomp.c-c++-common/target-is-accessible-2.c: New test.
* testsuite/libgomp.c-c++-common/target-is-accessible-3.c: New test.
* testsuite/libgomp.c-c++-common/target-is-accessible-4.c: New test.
---
OK for mainline?
Andrew
include/cuda/cuda.h | 9 ++
libgomp/libgomp-plugin.h | 1 +
libgomp/libgomp.h | 1 +
libgomp/plugin/cuda-lib.def | 1 +
libgomp/plugin/plugin-gcn.c | 50 ++++++++++
libgomp/plugin/plugin-nvptx.c | 35 +++++++
libgomp/target.c | 12 ++-
libgomp/testsuite/lib/libgomp.exp | 26 ++++++
.../target-is-accessible-1.c | 73 +++++++++++++--
.../target-is-accessible-2.c | 5 +
.../target-is-accessible-3.c | 4 +
.../target-is-accessible-4.c | 28 ++++++
.../target-is-accessible-1.f90 | 93 +++++++++++++++----
13 files changed, 310 insertions(+), 28 deletions(-)
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-3.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-4.c
diff --git a/include/cuda/cuda.h b/include/cuda/cuda.h
index 28510a3150c..ef62f150a5d 100644
--- a/include/cuda/cuda.h
+++ b/include/cuda/cuda.h
@@ -143,6 +143,13 @@ typedef enum {
CU_MEMORYTYPE_UNIFIED = 0x04
} CUmemorytype;
+typedef enum {
+ CU_POINTER_ATTRIBUTE_CONTEXT = 0x01,
+ CU_POINTER_ATTRIBUTE_MEMORY_TYPE = 0x02,
+ CU_POINTER_ATTRIBUTE_DEVICE_POINTER = 0x03,
+ CU_POINTER_ATTRIBUTE_HOST_POINTER = 0x04
+} CUpointer_attribute;
+
typedef struct {
size_t srcXInBytes, srcY;
CUmemorytype srcMemoryType;
@@ -300,6 +307,8 @@ CUresult cuModuleGetGlobal (CUdeviceptr *, size_t *,
CUmodule, const char *);
CUresult cuModuleLoad (CUmodule *, const char *);
CUresult cuModuleLoadData (CUmodule *, const void *);
CUresult cuModuleUnload (CUmodule);
+CUresult cuPointerGetAttribute (CUmemorytype *, CUpointer_attribute,
+ CUdeviceptr);
CUresult cuOccupancyMaxPotentialBlockSize(int *, int *, CUfunction,
CUoccupancyB2DSize, size_t, int);
typedef void (*CUstreamCallback)(CUstream, CUresult, void *);
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 71e74527e71..cf247caf51e 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -173,6 +173,7 @@ extern void *GOMP_OFFLOAD_alloc (int, size_t);
extern bool GOMP_OFFLOAD_free (int, void *);
extern void *GOMP_OFFLOAD_managed_alloc (int, size_t);
extern bool GOMP_OFFLOAD_managed_free (int, void *);
+extern bool GOMP_OFFLOAD_is_shared_ptr (int, const void *, size_t);
extern bool GOMP_OFFLOAD_page_locked_host_alloc (void **, size_t);
extern bool GOMP_OFFLOAD_page_locked_host_free (void *);
extern bool GOMP_OFFLOAD_dev2host (int, void *, const void *, size_t);
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 46db7d41f32..25d8b593c41 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1425,6 +1425,7 @@ struct gomp_device_descr
__typeof (GOMP_OFFLOAD_free) *free_func;
__typeof (GOMP_OFFLOAD_managed_alloc) *managed_alloc_func;
__typeof (GOMP_OFFLOAD_managed_free) *managed_free_func;
+ __typeof (GOMP_OFFLOAD_is_shared_ptr) *is_shared_ptr_func;
__typeof (GOMP_OFFLOAD_page_locked_host_alloc) *page_locked_host_alloc_func;
__typeof (GOMP_OFFLOAD_page_locked_host_free) *page_locked_host_free_func;
__typeof (GOMP_OFFLOAD_dev2host) *dev2host_func;
diff --git a/libgomp/plugin/cuda-lib.def b/libgomp/plugin/cuda-lib.def
index 67c783d8566..f87987db85c 100644
--- a/libgomp/plugin/cuda-lib.def
+++ b/libgomp/plugin/cuda-lib.def
@@ -55,6 +55,7 @@ CUDA_ONE_CALL (cuModuleLoad)
CUDA_ONE_CALL (cuModuleLoadData)
CUDA_ONE_CALL (cuModuleUnload)
CUDA_ONE_CALL_MAYBE_NULL (cuOccupancyMaxPotentialBlockSize)
+CUDA_ONE_CALL (cuPointerGetAttribute)
CUDA_ONE_CALL (cuStreamAddCallback)
CUDA_ONE_CALL (cuStreamCreate)
CUDA_ONE_CALL (cuStreamDestroy)
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index ece41c59bbb..90cc165b3ec 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -233,6 +233,9 @@ struct hsa_runtime_fn_info
hsa_status_t (*hsa_amd_svm_attributes_set_fn)
(void* ptr, size_t size, hsa_amd_svm_attribute_pair_t* attribute_list,
size_t attribute_count);
+ hsa_status_t (*hsa_amd_svm_attributes_get_fn)
+ (void* ptr, size_t size, hsa_amd_svm_attribute_pair_t* attribute_list,
+ size_t attribute_count);
};
/* As an HIP runtime is dlopened, following structure defines function
@@ -1494,6 +1497,7 @@ init_hsa_runtime_functions (void)
DLSYM_OPT_FN (hsa_amd_memory_unlock)
DLSYM_OPT_FN (hsa_amd_memory_async_copy_rect)
DLSYM_OPT_FN (hsa_amd_svm_attributes_set)
+ DLSYM_OPT_FN (hsa_amd_svm_attributes_get)
return true;
#undef DLSYM_OPT_FN
#undef DLSYM_FN
@@ -5240,6 +5244,52 @@ GOMP_OFFLOAD_managed_free (int device, void *ptr)
return true;
}
+bool
+GOMP_OFFLOAD_is_shared_ptr (int device, const void *ptr, size_t size)
+{
+ if (!init_hsa_context (false)
+ || device < 0 || device > hsa_context.agent_count
+ || !hsa_fns.hsa_amd_svm_attributes_get_fn)
+ return false;
+
+ struct agent_info *agent = get_agent_info (device);
+
+ /* The HSA API doesn't seem to report for the whole range given, so we call
+ once for each page the range straddles. */
+ const void *p = ptr;
+ size_t remaining = size;
+ do
+ {
+ /* Note: the access query returns in the attribute field. */
+ struct hsa_amd_svm_attribute_pair_s attr = {
+ HSA_AMD_SVM_ATTRIB_ACCESS_QUERY, agent->id.handle
+ };
+ hsa_status_t status = hsa_fns.hsa_amd_svm_attributes_get_fn ((void*)p,
+ remaining,
+ &attr, 1);
+ if (status != HSA_STATUS_SUCCESS)
+ /* This shouldn't happen, but if it does we can't say the memory is
+ accessible. */
+ return false;
+
+ switch (attr.attribute)
+ {
+ case HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE:
+ case HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE_IN_PLACE:
+ break;
+ case HSA_AMD_SVM_ATTRIB_AGENT_NO_ACCESS:
+ default:
+ return false;
+ }
+
+ p = (void*)(((uintptr_t)p + 4096) & ~0xfffUL);
+ remaining = size - ((uintptr_t)p - (uintptr_t)ptr);
+ } while (p < ptr + size);
+
+ /* All pages were accessible. */
+ return true;
+}
+
/* }}} */
/* {{{ OpenACC Plugin API */
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index dd8bcf9c507..5d49c437c78 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1906,6 +1906,41 @@ GOMP_OFFLOAD_managed_free (int ord, void *ptr)
return GOMP_OFFLOAD_free (ord, ptr);
}
+bool
+GOMP_OFFLOAD_is_shared_ptr (int device __attribute__((unused)),
+ const void *ptr, size_t size)
+{
+ /* The Cuda API does not permit testing a whole range, so we test each
+ 4K page within the range. If any page is inaccessible return false. */
+ const void *p = ptr;
+ do
+ {
+ CUmemorytype mem_type;
+ CUresult res = CUDA_CALL_NOCHECK (cuPointerGetAttribute, &mem_type,
+ CU_POINTER_ATTRIBUTE_MEMORY_TYPE,
+ (CUdeviceptr)p);
+ if (res != CUDA_SUCCESS)
+ /* Memory is not registered, and therefore not accessible. */
+ return false;
+
+ switch (mem_type)
+ {
+ case CU_MEMORYTYPE_HOST:
+ case CU_MEMORYTYPE_UNIFIED:
+ case CU_MEMORYTYPE_DEVICE:
+ break;
+ case CU_MEMORYTYPE_ARRAY:
+ default:
+ return false;
+ }
+
+ p = (void*)(((uintptr_t)p + 4096) & ~0xfffUL);
+ } while (p < ptr + size);
+
+ /* All pages were accessible. */
+ return true;
+}
+
bool
GOMP_OFFLOAD_page_locked_host_alloc (void **ptr, size_t size)
{
diff --git a/libgomp/target.c b/libgomp/target.c
index 859513b13b2..d413a934ded 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -5568,9 +5568,16 @@ omp_target_is_accessible (const void *ptr, size_t size,
int device_num)
if (devicep == NULL)
return false;
- /* TODO: Unified shared memory must be handled when available. */
+ /* Unified shared memory (or true shared memory). */
+ if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return true;
- return devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM;
+ /* Managed memory (or other device feature). */
+ if (devicep->is_shared_ptr_func
+ && devicep->is_shared_ptr_func (devicep->target_id, ptr, size))
+ return true;
+
+ return false;
}
int
@@ -5971,6 +5978,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr
*device,
DLSYM (free);
DLSYM_OPT (managed_alloc, managed_alloc);
DLSYM_OPT (managed_free, managed_free);
+ DLSYM_OPT (is_shared_ptr, is_shared_ptr);
DLSYM_OPT (page_locked_host_alloc, page_locked_host_alloc);
DLSYM_OPT (page_locked_host_free, page_locked_host_free);
DLSYM (dev2host);
diff --git a/libgomp/testsuite/lib/libgomp.exp
b/libgomp/testsuite/lib/libgomp.exp
index f5683b50725..e9979392aae 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -723,6 +723,32 @@ int main() {
} } "-lhipblas" ]
}
+# return 1 if OpenMP Unified Shared Memory is supported
+
+proc check_effective_target_omp_usm { } {
+ if { [check_effective_target_offload_device_nvptx] } {
+ return 1
+ }
+
+ if { [libgomp_check_effective_target_offload_target "amdgcn"] } {
+ if [check_runtime_nocache usm_available_ {
+ #include <omp.h>
+ #pragma omp requires unified_shared_memory
+ int main ()
+ {
+ int a;
+ #pragma omp target map(from: a)
+ a = omp_is_initial_device ();
+ return a;
+ }
+ } ] {
+ return 1
+ }
+ }
+
+ return 0
+}
+
# return 1 if OpenMP Device Managed Memory is supported
proc check_effective_target_omp_managedmem { } {
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c
b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c
index 2e75c6300ae..5745ef4013d 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c
@@ -1,4 +1,8 @@
#include <omp.h>
+#include <stdint.h>
+#include <stdlib.h>
+
+//#define __builtin_abort() __builtin_printf ("fail at line %d\n", __LINE__)
int
main ()
@@ -6,7 +10,7 @@ main ()
int d = omp_get_default_device ();
int id = omp_get_initial_device ();
int n = omp_get_num_devices ();
- void *p;
+ int *p = (int*)malloc (sizeof (int));
if (d < 0 || d >= n)
d = id;
@@ -20,30 +24,81 @@ main ()
if (!omp_target_is_accessible (p, sizeof (int), omp_initial_device))
__builtin_abort ();
- if (omp_target_is_accessible (p, sizeof (int), -5))
+ if (omp_target_is_accessible (p, sizeof (int), -99))
__builtin_abort ();
if (omp_target_is_accessible (p, sizeof (int), n + 1))
__builtin_abort ();
- /* Currently, a host pointer is accessible if the device supports shared
- memory or omp_target_is_accessible is executed on the host. This
- test case must be adapted when unified shared memory is avialable. */
int a[128];
for (int d = 0; d <= omp_get_num_devices (); d++)
{
+ /* Check if libgomp is treating the device as a shared memory device. */
int shared_mem = 0;
#pragma omp target map (alloc: shared_mem) device (d)
shared_mem = 1;
+
+ int heap_accessible = shared_mem;
if (omp_target_is_accessible (p, sizeof (int), d) != shared_mem)
- __builtin_abort ();
+ {
+ if (shared_mem)
+ __builtin_abort ();
+
+ /* shared_mem is false, but the memory is reading as accessible,
+ so let's check that by reading it. We should not do so
+ unconditionally because if it's wrong then we'll probably get
+ a memory fault. */
+ *p = 123;
+ uintptr_t addr = (uintptr_t)p;
+
+ #pragma omp target is_device_ptr(p) map(from:heap_accessible) \
+ device(d)
+ {
+ if ((uintptr_t)p == addr && *p == 123)
+ heap_accessible = 1;
+ }
+
+ if (!heap_accessible)
+ __builtin_abort ();
+ }
+ int stack_accessible = shared_mem;
if (omp_target_is_accessible (a, 128 * sizeof (int), d) != shared_mem)
- __builtin_abort ();
+ {
+ if (shared_mem)
+ __builtin_abort ();
+
+ /* shared_mem is false, but the memory is reading as accessible,
+ so let's check that by reading it. We should not do so
+ unconditionally because if it's wrong then we'll probably get
+ a memory fault. */
+ int test_accessible = 123;
+ uintptr_t addr = (uintptr_t)&test_accessible;
+
+ #pragma omp target has_device_addr(test_accessible) \
+ map(from:stack_accessible) device(d)
+ {
+ if ((uintptr_t)&test_accessible == addr
+ && test_accessible == 123)
+ stack_accessible = 1;
+ }
+ if (!stack_accessible)
+ __builtin_abort ();
+ }
+ __builtin_printf ("device #%d: shared_mem=%d heap_accessible=%d "
+ "stack_accessible=%d\n",
+ d, shared_mem, heap_accessible, stack_accessible);
+
+ /* omp_target_is_accessible returns false if *any* of the array is
+ inaccessible, so we only check the aggregate result.
+ (Varying access observed on amdgcn without xnack.) */
+ bool accessible = true;
for (int i = 0; i < 128; i++)
- if (omp_target_is_accessible (&a[i], sizeof (int), d) != shared_mem)
- __builtin_abort ();
+ if (!omp_target_is_accessible (&a[i], sizeof (int), d))
+ accessible = false;
+ if (accessible != (shared_mem || stack_accessible))
+ __builtin_abort ();
}
return 0;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c
b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c
new file mode 100644
index 00000000000..205c0a41014
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c
@@ -0,0 +1,5 @@
+/* { dg-require-effective-target omp_usm } */
+
+#pragma omp requires unified_shared_memory
+
+#include "target-is-accessible-1.c"
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-3.c
b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-3.c
new file mode 100644
index 00000000000..4e601b930c2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-3.c
@@ -0,0 +1,4 @@
+/* { dg-require-effective-target offload_target_amdgcn_with_xnack } */
+/* { dg-additional-options "-foffload-options=amdgcn-amdhsa=-mxnack=on" } */
+
+#include "target-is-accessible-1.c"
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-4.c
b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-4.c
new file mode 100644
index 00000000000..0478fb0134f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-4.c
@@ -0,0 +1,28 @@
+/* { dg-require-effective-target omp_managedmem } */
+
+#include <omp.h>
+#include <stdint.h>
+#include <stdlib.h>
+
+//#define __builtin_abort() __builtin_printf ("fail at line %d\n", __LINE__)
+
+int
+main ()
+{
+ int *p = (int*)omp_alloc (sizeof (int), ompx_gnu_managed_mem_alloc);
+
+ *p = 42;
+ uintptr_t a_p = (uintptr_t)p;
+
+ #pragma omp target is_device_ptr(p)
+ {
+ if (*p != 42 || a_p != (uintptr_t)p)
+ __builtin_abort ();
+ }
+ if (!p
+ || !omp_target_is_accessible (p, sizeof (int),
+ omp_get_default_device ()))
+ __builtin_abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
index 150df6f8a4f..d33fe73292f 100644
--- a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
@@ -1,53 +1,112 @@
+! { dg-do run }
+
program main
use omp_lib
use iso_c_binding
implicit none (external, type)
- integer :: d, id, n, shared_mem, i
+ integer :: d, id, n, shared_mem, i, heap_accessible, stack_accessible
+ integer, target :: test_accessible
+ integer, allocatable, target :: p(:)
integer, target :: a(1:128)
- type(c_ptr) :: p
+ integer(c_intptr_t) :: addr
+ logical :: condition
d = omp_get_default_device ()
id = omp_get_initial_device ()
n = omp_get_num_devices ()
+ allocate (p(1))
if (d < 0 .or. d >= n) &
d = id
- if (omp_target_is_accessible (p, c_sizeof (d), n) /= 1) &
+ if (omp_target_is_accessible (c_loc(p), c_sizeof(p(1)), n) == 0) &
stop 1
- if (omp_target_is_accessible (p, c_sizeof (d), id) /= 1) &
+ if (omp_target_is_accessible (c_loc(p), c_sizeof(p(1)), id) == 0) &
stop 2
- if (omp_target_is_accessible (p, c_sizeof (d), omp_initial_device) /= 1) &
+ if (omp_target_is_accessible (c_loc(p), c_sizeof(p(1)), omp_initial_device)
== 0) &
stop 3
- if (omp_target_is_accessible (p, c_sizeof (d), -5) /= 0) &
+ if (omp_target_is_accessible (c_loc(p), c_sizeof(p(1)), -99) /= 0) &
stop 4
- if (omp_target_is_accessible (p, c_sizeof (d), n + 1) /= 0) &
+ if (omp_target_is_accessible (c_loc(p), c_sizeof(p(1)), n + 1) /= 0) &
stop 5
! Currently, a host pointer is accessible if the device supports shared
! memory or omp_target_is_accessible is executed on the host. This
- ! test case must be adapted when unified shared memory is avialable.
+ ! test case must be adapted when unified shared memory is available.
do d = 0, omp_get_num_devices ()
- shared_mem = 0;
+ ! Check if libgomp is treating the device as a shared memory device.
+ shared_mem = 0
!$omp target map (alloc: shared_mem) device (d)
- shared_mem = 1;
+ shared_mem = 1
!$omp end target
- if (omp_target_is_accessible (p, c_sizeof (d), d) /= shared_mem) &
- stop 6;
+ heap_accessible = shared_mem
+ condition = omp_target_is_accessible (c_loc(p), c_sizeof(p(1)), d) /=
shared_mem
+ if (condition) then
+ if (shared_mem /= 0) &
+ stop 6
+
+ ! shared_mem is false, but the memory is reading as accessible,
+ ! so let's check that by reading it. We should not do so
+ ! unconditionally because if it's wrong then we'll probably get
+ ! a memory fault.
+ p(1) = 123
+ addr = transfer(c_loc(p), addr)
+
+ !$omp target is_device_ptr(p) map(from:heap_accessible) device(d)
+ if (transfer(c_loc(p), addr) == addr .and. p(1) == 123) &
+ heap_accessible = 1
+ !$omp end target
+
+ if (heap_accessible == 0) &
+ stop 7
+ end if
+
+ stack_accessible = shared_mem
+ condition = omp_target_is_accessible (c_loc(a), 128 * c_sizeof(a(1)), d)
/= shared_mem
+ if (condition) then
+ if (shared_mem /= 0) &
+ stop 8
- if (omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) /=
shared_mem) &
- stop 7;
+ ! shared_mem is false, but the memory is reading as accessible,
+ ! so let's check that by reading it. We should not do so
+ ! unconditionally because if it's wrong then we'll probably get
+ ! a memory fault.
+ test_accessible = 123
+ addr = transfer(c_loc(test_accessible), addr)
+ !$omp target has_device_addr(test_accessible) map(from:stack_accessible)
device(d)
+ if (transfer(c_loc(test_accessible), addr) == addr &
+ .and. test_accessible == 123) &
+ stack_accessible = 1
+ !$omp end target
+
+ if (stack_accessible == 0) &
+ stop 9
+ end if
+
+ print '(A,I0,A,I0,A,I0,A,I0)', &
+ 'device #', d, &
+ ': shared_mem=', shared_mem, &
+ ' heap_accessible=', heap_accessible, &
+ ' stack_accessible=', stack_accessible
+
+ ! omp_target_is_accessible returns false if *any* of the array is
+ ! inaccessible, so we only check the aggregate result.
+ ! (Varying access observed on amdgcn without xnack.)
+ condition = .true.
do i = 1, 128
- if (omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) /=
shared_mem) &
- stop 8;
+ if (omp_target_is_accessible (c_loc(a(i)), c_sizeof(a(i)), d) == 0) &
+ condition = .false.
end do
-
+ if (condition .neqv. stack_accessible /= 0) &
+ stop 10
end do
+ deallocate (p)
+
end program main
--
2.51.0
