On 09/12/2025 16:46, Tobias Burnus wrote:
Thus, we are left with:* Two deferred items (unchanged, presumably we want to update the PR).
Ack.
* What to do about 'omp_target_alloc' and similar memory
I've added new device allocation management to the GCN support, and set the NVPTX context so it shouldn't be on the wrong device.
* RFC whether to pass a boolean to the plugin function to be future proof
I thought the idea was that the plugin would (at some point in future) report everything it knows. Is the point of the boolean to optimize the case where some of the checks are expensive?
Another option would be to add a new (optional) plugin API for the enhanced mode, which would be provided by plugins that support (or require?) extra.
I didn't do anything about that.
* Checking PAGEABLE_MEMORY_ACCESS for the passed device num plus using the right device (→ device context) for the check.
The current code seems to give the correct result in all cases. More API calls means more runtime overhead, so what is the point of this?
Note that USM is either on or off, right now. If any device present doesn't support USM then the plugin reports that *no* devices are available.
I've attached my v4 patch. Andrew
From aa2320222090244c9cada0914ceee59e8bd504d3 Mon Sep 17 00:00:00 2001 From: Andrew Stubbs <[email protected]> Date: Mon, 8 Dec 2025 16:18:02 +0000 Subject: [PATCH v4] libgomp, amdgcn, nvptx: Improve omp_target_is_accessible [PR121813] 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_accessible_ptr): New prototype. * libgomp.h (struct gomp_device_descr): Add GOMP_OFFLOAD_is_accessible_ptr. * libgomp.texi: Update omp_target_is_accessible docs. * plugin/cuda-lib.def (cuPointerGetAttribute): New entry. * plugin/plugin-gcn.c (struct hsa_runtime_fn_info): Add hsa_amd_svm_attributes_get_fn and hsa_amd_pointer_info_fn. (init_hsa_runtime_functions): Add hsa_amd_svm_attributes_get and hsa_amd_pointer_info. (enum accessible): New enum type. (host_memory_is_accessible): New function. (device_memory_is_accessible): New function. (GOMP_OFFLOAD_is_accessible_ptr): New function. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_is_accessible_ptr): Likewise. * target.c (omp_target_is_accessible): Call is_accessible_ptr_func. (gomp_load_plugin_for_device): Add is_accessible_ptr. * 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. * testsuite/libgomp.c-c++-common/target-is-accessible-5.c: New test. --- include/cuda/cuda.h | 9 ++ libgomp/libgomp-plugin.h | 1 + libgomp/libgomp.h | 1 + libgomp/libgomp.texi | 14 ++- libgomp/plugin/cuda-lib.def | 1 + libgomp/plugin/plugin-gcn.c | 111 ++++++++++++++++++ libgomp/plugin/plugin-nvptx.c | 69 +++++++++-- libgomp/target.c | 18 ++- .../target-is-accessible-1.c | 77 ++++++++++-- .../target-is-accessible-2.c | 5 + .../target-is-accessible-3.c | 4 + .../target-is-accessible-4.c | 28 +++++ .../target-is-accessible-5.c | 33 ++++++ .../target-is-accessible-1.f90 | 93 ++++++++++++--- 14 files changed, 423 insertions(+), 41 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 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-5.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..ad01d37ff35 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 int GOMP_OFFLOAD_is_accessible_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..0496c2a803e 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_accessible_ptr) *is_accessible_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/libgomp.texi b/libgomp/libgomp.texi index ac96d2f29c7..602b1d0469e 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -2171,13 +2171,17 @@ and extending @var{size} bytes, is accessibly on the device specified by @var{device_num}. If so, it returns a nonzero value and otherwise zero. The address given by @var{ptr} is interpreted to be in the address space of -the device and @var{size} must be positive. +the device and @var{size} must be positive. NULL pointers and zero-length +ranges always return zero. Note that GCC's current implementation assumes that @var{ptr} is a valid host -pointer. Therefore, all addresses given by @var{ptr} are assumed to be -accessible on the initial device. And, to err on the safe side, this memory -is only available on a non-host device that can access all host memory -([uniform] shared memory access). +pointer. Therefore, all non-NULL addresses given by @var{ptr} are assumed to be +accessible on the initial device. The address is only reported as accessible +on non-host devices if this is @emph{known} to be the case, or if the device +reports that all memory is accessible (i.e. [unified] shared memory access). +If the runtime is uncertain it may report accessible memory as inaccessible. +For a memory range to be reported accessible, the whole range must be known to +be accessible. Running this routine in a @code{target} region except on the initial device is not supported. 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 92de6fb1b64..47e1df6cf94 100644 --- a/libgomp/plugin/plugin-gcn.c +++ b/libgomp/plugin/plugin-gcn.c @@ -233,6 +233,12 @@ 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); + hsa_status_t (*hsa_amd_pointer_info_fn) + (const void *, hsa_amd_pointer_info_t *, void *(*)(size_t), + uint32_t *, hsa_agent_t **); }; /* As an HIP runtime is dlopened, following structure defines function @@ -1494,6 +1500,8 @@ 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) + DLSYM_OPT_FN (hsa_amd_pointer_info) return true; #undef DLSYM_OPT_FN #undef DLSYM_FN @@ -5258,6 +5266,109 @@ GOMP_OFFLOAD_managed_free (int device, void *ptr) return true; } +enum accessible { + UNKNOWN, + INACCESSIBLE, + ACCESSIBLE +}; + +/* Is a host memory address accessible on the given device? + Returns UNKNOWN if the memory isn't registered, or if it isn't a valid host + pointer. */ + +static enum accessible +host_memory_is_accessible (hsa_agent_t agent, const void *ptr, size_t size) +{ + if (!hsa_fns.hsa_amd_svm_attributes_get_fn) + return UNKNOWN; + + /* 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.handle + }; + hsa_status_t status = hsa_fns.hsa_amd_svm_attributes_get_fn ((void*)p, + remaining, + &attr, 1); + if (status != HSA_STATUS_SUCCESS) + /* This happens when the memory isn't registered with ROCr at all. */ + return UNKNOWN; + + 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 INACCESSIBLE; + } + + p = (void*)(((uintptr_t)p + 4096) & ~0xfffUL); + remaining = size - ((uintptr_t)p - (uintptr_t)ptr); + } while (p < ptr + size); + + /* All pages were accessible. */ + return ACCESSIBLE; +} + +/* Is a device memory address accessible on the given device? + Returns UNKNOWN if it isn't a valid device address. Returns INACCESSIBLE if + the pointer is valid, but not the whole range, or if it refers to the wrong + device. */ + +static enum accessible +device_memory_is_accessible (hsa_agent_t agent, const void *ptr, size_t size) +{ + if (!hsa_fns.hsa_amd_pointer_info_fn) + return UNKNOWN; + + hsa_amd_pointer_info_t info; + uint32_t nagents; + hsa_agent_t *agents; + info.size = sizeof (hsa_amd_pointer_info_t); + + hsa_status_t status = hsa_fns.hsa_amd_pointer_info_fn (ptr, &info, NULL, + &nagents, &agents); + if (status != HSA_STATUS_SUCCESS + || info.type == HSA_EXT_POINTER_TYPE_UNKNOWN) + return UNKNOWN; + + if (agent.handle == info.agentOwner.handle) + return (info.sizeInBytes >= size ? ACCESSIBLE : INACCESSIBLE); + + for (unsigned i = 0; i < nagents; i++) + { + if (agent.handle == agents[0].handle) + return (info.sizeInBytes >= size ? ACCESSIBLE : INACCESSIBLE); + } + + return INACCESSIBLE; +} + +/* Backend implementation for omp_target_is_accessible. */ + +int +GOMP_OFFLOAD_is_accessible_ptr (int device, const void *ptr, size_t size) +{ + if (!init_hsa_context (false) + || device < 0 || device > hsa_context.agent_count) + return 0; + + struct agent_info *agent = get_agent_info (device); + + enum accessible result; + result = host_memory_is_accessible (agent->id, ptr, size); + if (result == UNKNOWN) + result = device_memory_is_accessible (agent->id, ptr, size); + return result == ACCESSIBLE; +} + /* }}} */ /* {{{ OpenACC Plugin API */ diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c index dd8bcf9c507..83c0a5ce4dd 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -353,6 +353,8 @@ struct ptx_device static struct ptx_device **ptx_devices; +static bool using_usm = false; + /* "Native" GPU thread stack size. */ static unsigned native_gpu_thread_stack_size = 0; @@ -1343,15 +1345,20 @@ GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask) if (num_devices > 0 && (omp_requires_mask & (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY | GOMP_REQUIRES_SELF_MAPS))) - for (int dev = 0; dev < num_devices; dev++) - { - int pi; - CUresult r; - r = CUDA_CALL_NOCHECK (cuDeviceGetAttribute, &pi, - CU_DEVICE_ATTRIBUTE_PAGEABLE_MEMORY_ACCESS, dev); - if (r != CUDA_SUCCESS || pi == 0) - return -1; - } + { + for (int dev = 0; dev < num_devices; dev++) + { + int pi; + CUresult r; + r = CUDA_CALL_NOCHECK (cuDeviceGetAttribute, &pi, + CU_DEVICE_ATTRIBUTE_PAGEABLE_MEMORY_ACCESS, + dev); + if (r != CUDA_SUCCESS || pi == 0) + return -1; + } + + using_usm = true; + } return num_devices; } @@ -1906,6 +1913,50 @@ GOMP_OFFLOAD_managed_free (int ord, void *ptr) return GOMP_OFFLOAD_free (ord, ptr); } +int +GOMP_OFFLOAD_is_accessible_ptr (int ord, + const void *ptr, size_t size) +{ + /* USM implies access. */ + if (using_usm) + return 1; + + struct ptx_device *ptx_dev = ptx_devices[ord]; + CUcontext old_ctx; + CUDA_CALL_ERET (false, cuCtxPushCurrent, ptx_dev->ctx); + + /* 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; + int result = 1; /* All pages accessible. */ + 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. */ + result = 0; + + switch (mem_type) + { + case CU_MEMORYTYPE_HOST: + case CU_MEMORYTYPE_UNIFIED: + case CU_MEMORYTYPE_DEVICE: + break; + case CU_MEMORYTYPE_ARRAY: + default: + result = 0; /* This page isn't accessible. */ + } + + p = (void*)(((uintptr_t)p + 4096) & ~0xfffUL); + } while (result && p < ptr + size); + + CUDA_CALL_ASSERT (cuCtxPopCurrent, &old_ctx); + return result; +} + bool GOMP_OFFLOAD_page_locked_host_alloc (void **ptr, size_t size) { diff --git a/libgomp/target.c b/libgomp/target.c index af7c702d439..896bedfd172 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -5590,6 +5590,9 @@ omp_get_mapped_ptr (const void *ptr, int device_num) int omp_target_is_accessible (const void *ptr, size_t size, int device_num) { + if (ptr == NULL || size == 0) + return false; + if (device_num == omp_default_device) device_num = gomp_get_default_device (); @@ -5601,9 +5604,19 @@ 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. */ + /* Managed memory (or other device feature). + is_accessible_ptr may, in future, report more than simply true or false, + but we can assume that positive responses are accessible, and + zero/negative responses are inaccessible. */ + if (devicep->is_accessible_ptr_func) + return (devicep->is_accessible_ptr_func (devicep->target_id, ptr, size) + > 0); + + /* 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; + return false; } int @@ -6009,6 +6022,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_accessible_ptr, is_accessible_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/libgomp.c-c++-common/target-is-accessible-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c index 71d6b2a8360..fff590e02fd 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; @@ -26,24 +30,81 @@ main () 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++) { + if (omp_target_is_accessible (NULL, 1, d)) + __builtin_abort (); + + if (omp_target_is_accessible (p, 0, d)) + __builtin_abort (); + + /* 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.c-c++-common/target-is-accessible-5.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-5.c new file mode 100644 index 00000000000..d6f3df7d6af --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-5.c @@ -0,0 +1,33 @@ +#include <omp.h> + +void check (int dev) +{ + constexpr int N = 10; + constexpr int size = N*sizeof(int); + int A[N] = {}; + + void *ptr = omp_target_alloc (size, dev); + + if (ptr == nullptr || !omp_target_is_accessible (ptr, size, dev)) + __builtin_abort (); + + #pragma omp target device(dev) firstprivate(ptr) + for (int i = 0; i < N; i++) + ((int *)ptr)[i] = i + 1; + + if (omp_target_memcpy (A, ptr, size, 0, 0, omp_initial_device, dev) != 0) + __builtin_abort (); + + for (int i = 0; i < N; i++) + if (A[i] != i + 1) + __builtin_abort (); + + omp_target_free (ptr, dev); +} + +int main () +{ + check (omp_default_device); + for (int dev = 0; dev <= omp_get_num_devices(); dev++) + check (dev); +} diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 index 2c25dca22f2..4cb401b18e4 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), -6) /= 0) & ! -6 = omp_default_device - 1 + if (omp_target_is_accessible (c_loc(p), c_sizeof(p(1)), -6) /= 0) & ! -6 = omp_default_device - 1 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
