diff options
author | Martin Liska <mliska@suse.cz> | 2022-10-26 12:59:00 +0200 |
---|---|---|
committer | Martin Liska <mliska@suse.cz> | 2022-10-26 12:59:00 +0200 |
commit | 62e475bad0d668c432bb97113cbf73fa281b8b55 (patch) | |
tree | 1c8993afe363ddbf4ae80795e47042df1e33d181 /libgomp | |
parent | 5776a5ffab3b92d6ccac87ccf32c580ee2742d5a (diff) | |
parent | d80b7744c6ae6f6d8ca1f56982a50d1888b8279f (diff) | |
download | gcc-62e475bad0d668c432bb97113cbf73fa281b8b55.zip gcc-62e475bad0d668c432bb97113cbf73fa281b8b55.tar.gz gcc-62e475bad0d668c432bb97113cbf73fa281b8b55.tar.bz2 |
Merge branch 'master' into devel/sphinx
Diffstat (limited to 'libgomp')
-rw-r--r-- | libgomp/ChangeLog | 39 | ||||
-rw-r--r-- | libgomp/config/nvptx/icv-device.c | 2 | ||||
-rw-r--r-- | libgomp/config/nvptx/libgomp-nvptx.h | 51 | ||||
-rw-r--r-- | libgomp/config/nvptx/target.c | 54 | ||||
-rw-r--r-- | libgomp/libgomp-plugin.c | 12 | ||||
-rw-r--r-- | libgomp/libgomp-plugin.h | 7 | ||||
-rw-r--r-- | libgomp/libgomp.h | 5 | ||||
-rw-r--r-- | libgomp/libgomp.map | 5 | ||||
-rw-r--r-- | libgomp/plugin/cuda-lib.def | 1 | ||||
-rw-r--r-- | libgomp/plugin/plugin-nvptx.c | 109 | ||||
-rw-r--r-- | libgomp/target.c | 19 |
11 files changed, 286 insertions, 18 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 5ebda8f..d61a806 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,42 @@ +2022-10-24 Thomas Schwinge <thomas@codesourcery.com> + + * plugin/plugin-nvptx.c (nvptx_open_device): Initialize + 'ptx_dev->rev_data'. + +2022-10-24 Tobias Burnus <tobias@codesourcery.com> + + * config/nvptx/icv-device.c (GOMP_DEVICE_NUM_VAR): Remove + 'static' for this variable. + * config/nvptx/libgomp-nvptx.h: New file. + * config/nvptx/target.c: Include it. + (GOMP_ADDITIONAL_ICVS): Declare extern var. + (GOMP_REV_OFFLOAD_VAR): Declare var. + (GOMP_target_ext): Handle reverse offload. + * libgomp-plugin.h (GOMP_PLUGIN_target_rev): New prototype. + * libgomp-plugin.c (GOMP_PLUGIN_target_rev): New, call ... + * target.c (gomp_target_rev): ... this new stub function. + * libgomp.h (gomp_target_rev): Declare. + * libgomp.map (GOMP_PLUGIN_1.4): New; add GOMP_PLUGIN_target_rev. + * plugin/cuda-lib.def (cuMemHostAlloc): Add. + * plugin/plugin-nvptx.c: Include libgomp-nvptx.h. + (struct ptx_device): Add rev_data member. + (nvptx_open_device): Remove async_engines query, last used in + r10-304-g1f4c5b9b; add unified-address assert check. + (GOMP_OFFLOAD_get_num_devices): Claim unified address + support. + (GOMP_OFFLOAD_load_image): Free rev_fn_table if no + offload functions exist. Make offload var available + on host and device. + (rev_off_dev_to_host_cpy, rev_off_host_to_dev_cpy): New. + (GOMP_OFFLOAD_run): Handle reverse offload. + +2022-10-21 Thomas Schwinge <thomas@codesourcery.com> + + PR tree-optimization/107195 + PR target/107344 + * testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Restore SESE + regions checking. + 2022-10-20 Tobias Burnus <tobias@codesourcery.com> * testsuite/lib/libgomp.exp (check_effective_target_offload_device_gcn): diff --git a/libgomp/config/nvptx/icv-device.c b/libgomp/config/nvptx/icv-device.c index 6f869be..eef151c 100644 --- a/libgomp/config/nvptx/icv-device.c +++ b/libgomp/config/nvptx/icv-device.c @@ -30,7 +30,7 @@ /* This is set to the ICV values of current GPU during device initialization, when the offload image containing this libgomp portion is loaded. */ -static volatile struct gomp_offload_icvs GOMP_ADDITIONAL_ICVS; +volatile struct gomp_offload_icvs GOMP_ADDITIONAL_ICVS; void omp_set_default_device (int device_num __attribute__((unused))) diff --git a/libgomp/config/nvptx/libgomp-nvptx.h b/libgomp/config/nvptx/libgomp-nvptx.h new file mode 100644 index 0000000..5da9aae --- /dev/null +++ b/libgomp/config/nvptx/libgomp-nvptx.h @@ -0,0 +1,51 @@ +/* Copyright (C) 2022 Free Software Foundation, Inc. + Contributed by Tobias Burnus <tobias@codesourcery.com>. + + This file is part of the GNU Offloading and Multi Processing Library + (libgomp). + + Libgomp is free software; you can redistribute it and/or modify it + under the terms of the GNU General Public License as published by + the Free Software Foundation; either version 3, or (at your option) + any later version. + + Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY + WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS + FOR A PARTICULAR PURPOSE. See the GNU General Public License for + more details. + + Under Section 7 of GPL version 3, you are granted additional + permissions described in the GCC Runtime Library Exception, version + 3.1, as published by the Free Software Foundation. + + You should have received a copy of the GNU General Public License and + a copy of the GCC Runtime Library Exception along with this program; + see the files COPYING3 and COPYING.RUNTIME respectively. If not, see + <http://www.gnu.org/licenses/>. */ + +/* This file contains defines and type definitions shared between the + nvptx target's libgomp.a and the plugin-nvptx.c, but that is only + needef for this target. */ + +#ifndef LIBGOMP_NVPTX_H +#define LIBGOMP_NVPTX_H 1 + +#define GOMP_REV_OFFLOAD_VAR __gomp_rev_offload_var + +struct rev_offload { + uint64_t fn; + uint64_t mapnum; + uint64_t addrs; + uint64_t sizes; + uint64_t kinds; + int32_t dev_num; +}; + +#if (__SIZEOF_SHORT__ != 2 \ + || __SIZEOF_SIZE_T__ != 8 \ + || __SIZEOF_POINTER__ != 8) +#error "Data-type conversion required for rev_offload" +#endif + +#endif /* LIBGOMP_NVPTX_H */ + diff --git a/libgomp/config/nvptx/target.c b/libgomp/config/nvptx/target.c index 11108d2..0e79388 100644 --- a/libgomp/config/nvptx/target.c +++ b/libgomp/config/nvptx/target.c @@ -24,9 +24,12 @@ <http://www.gnu.org/licenses/>. */ #include "libgomp.h" +#include "libgomp-nvptx.h" /* For struct rev_offload + GOMP_REV_OFFLOAD_VAR. */ #include <limits.h> extern int __gomp_team_num __attribute__((shared)); +extern volatile struct gomp_offload_icvs GOMP_ADDITIONAL_ICVS; +volatile struct rev_offload *GOMP_REV_OFFLOAD_VAR; bool GOMP_teams4 (unsigned int num_teams_lower, unsigned int num_teams_upper, @@ -88,16 +91,53 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, void **hostaddrs, size_t *sizes, unsigned short *kinds, unsigned int flags, void **depend, void **args) { - (void) device; - (void) fn; - (void) mapnum; - (void) hostaddrs; - (void) sizes; - (void) kinds; + static int lock = 0; /* == gomp_mutex_t lock; gomp_mutex_init (&lock); */ (void) flags; (void) depend; (void) args; - __builtin_unreachable (); + + if (device != GOMP_DEVICE_HOST_FALLBACK + || fn == NULL + || GOMP_REV_OFFLOAD_VAR == NULL) + return; + + gomp_mutex_lock (&lock); + + GOMP_REV_OFFLOAD_VAR->mapnum = mapnum; + GOMP_REV_OFFLOAD_VAR->addrs = (uint64_t) hostaddrs; + GOMP_REV_OFFLOAD_VAR->sizes = (uint64_t) sizes; + GOMP_REV_OFFLOAD_VAR->kinds = (uint64_t) kinds; + GOMP_REV_OFFLOAD_VAR->dev_num = GOMP_ADDITIONAL_ICVS.device_num; + + /* Set 'fn' to trigger processing on the host; wait for completion, + which is flagged by setting 'fn' back to 0 on the host. */ + uint64_t addr_struct_fn = (uint64_t) &GOMP_REV_OFFLOAD_VAR->fn; +#if __PTX_SM__ >= 700 + asm volatile ("st.global.release.sys.u64 [%0], %1;" + : : "r"(addr_struct_fn), "r" (fn) : "memory"); +#else + __sync_synchronize (); /* membar.sys */ + asm volatile ("st.volatile.global.u64 [%0], %1;" + : : "r"(addr_struct_fn), "r" (fn) : "memory"); +#endif + +#if __PTX_SM__ >= 700 + uint64_t fn2; + do + { + asm volatile ("ld.acquire.sys.global.u64 %0, [%1];" + : "=r" (fn2) : "r" (addr_struct_fn) : "memory"); + } + while (fn2 != 0); +#else + /* ld.global.u64 %r64,[__gomp_rev_offload_var]; + ld.u64 %r36,[%r64]; + membar.sys; */ + while (__atomic_load_n (&GOMP_REV_OFFLOAD_VAR->fn, __ATOMIC_ACQUIRE) != 0) + ; /* spin */ +#endif + + gomp_mutex_unlock (&lock); } void diff --git a/libgomp/libgomp-plugin.c b/libgomp/libgomp-plugin.c index 9d4cc62..316de74 100644 --- a/libgomp/libgomp-plugin.c +++ b/libgomp/libgomp-plugin.c @@ -78,3 +78,15 @@ GOMP_PLUGIN_fatal (const char *msg, ...) gomp_vfatal (msg, ap); va_end (ap); } + +void +GOMP_PLUGIN_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, + uint64_t sizes_ptr, uint64_t kinds_ptr, int dev_num, + void (*dev_to_host_cpy) (void *, const void *, size_t, + void *), + void (*host_to_dev_cpy) (void *, const void *, size_t, + void *), void *token) +{ + gomp_target_rev (fn_ptr, mapnum, devaddrs_ptr, sizes_ptr, kinds_ptr, dev_num, + dev_to_host_cpy, host_to_dev_cpy, token); +} diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h index 6ab5ac6..875f967 100644 --- a/libgomp/libgomp-plugin.h +++ b/libgomp/libgomp-plugin.h @@ -121,6 +121,13 @@ extern void GOMP_PLUGIN_error (const char *, ...) extern void GOMP_PLUGIN_fatal (const char *, ...) __attribute__ ((noreturn, format (printf, 1, 2))); +extern void GOMP_PLUGIN_target_rev (uint64_t, uint64_t, uint64_t, uint64_t, + uint64_t, int, + void (*) (void *, const void *, size_t, + void *), + void (*) (void *, const void *, size_t, + void *), void *); + /* Prototypes for functions implemented by libgomp plugins. */ extern const char *GOMP_OFFLOAD_get_name (void); extern unsigned int GOMP_OFFLOAD_get_caps (void); diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 7519274..5803683 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1128,6 +1128,11 @@ extern int gomp_pause_host (void); extern void gomp_init_targets_once (void); extern int gomp_get_num_devices (void); extern bool gomp_target_task_fn (void *); +extern void gomp_target_rev (uint64_t, uint64_t, uint64_t, uint64_t, uint64_t, + int, + void (*) (void *, const void *, size_t, void *), + void (*) (void *, const void *, size_t, void *), + void *); /* Splay tree definitions. */ typedef struct splay_tree_node_s *splay_tree_node; diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index e76c126..ce6b719 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -628,3 +628,8 @@ GOMP_PLUGIN_1.3 { GOMP_PLUGIN_goacc_profiling_dispatch; GOMP_PLUGIN_goacc_thread; } GOMP_PLUGIN_1.2; + +GOMP_PLUGIN_1.4 { + global: + GOMP_PLUGIN_target_rev; +} GOMP_PLUGIN_1.3; diff --git a/libgomp/plugin/cuda-lib.def b/libgomp/plugin/cuda-lib.def index cd91b39..dff42d6 100644 --- a/libgomp/plugin/cuda-lib.def +++ b/libgomp/plugin/cuda-lib.def @@ -29,6 +29,7 @@ CUDA_ONE_CALL_MAYBE_NULL (cuLinkCreate_v2) CUDA_ONE_CALL (cuLinkDestroy) CUDA_ONE_CALL (cuMemAlloc) CUDA_ONE_CALL (cuMemAllocHost) +CUDA_ONE_CALL (cuMemHostAlloc) CUDA_ONE_CALL (cuMemcpy) CUDA_ONE_CALL (cuMemcpyDtoDAsync) CUDA_ONE_CALL (cuMemcpyDtoH) diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c index ba6b229..0768fca 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -40,6 +40,9 @@ #include "gomp-constants.h" #include "oacc-int.h" +/* For struct rev_offload + GOMP_REV_OFFLOAD_VAR. */ +#include "config/nvptx/libgomp-nvptx.h" + #include <pthread.h> #ifndef PLUGIN_NVPTX_INCLUDE_SYSTEM_CUDA_H # include "cuda/cuda.h" @@ -329,6 +332,7 @@ struct ptx_device pthread_mutex_t lock; } omp_stacks; + struct rev_offload *rev_data; struct ptx_device *next; }; @@ -423,7 +427,7 @@ nvptx_open_device (int n) struct ptx_device *ptx_dev; CUdevice dev, ctx_dev; CUresult r; - int async_engines, pi; + int pi; CUDA_CALL_ERET (NULL, cuDeviceGet, &dev, n); @@ -519,10 +523,12 @@ nvptx_open_device (int n) CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR, dev); ptx_dev->max_threads_per_multiprocessor = pi; - r = CUDA_CALL_NOCHECK (cuDeviceGetAttribute, &async_engines, - CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT, dev); - if (r != CUDA_SUCCESS) - async_engines = 1; + /* Required below for reverse offload as implemented, but with compute + capability >= 2.0 and 64bit device processes, this should be universally be + the case; hence, an assert. */ + r = CUDA_CALL_NOCHECK (cuDeviceGetAttribute, &pi, + CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, dev); + assert (r == CUDA_SUCCESS && pi); for (int i = 0; i != GOMP_DIM_MAX; i++) ptx_dev->default_dims[i] = 0; @@ -540,6 +546,8 @@ nvptx_open_device (int n) ptx_dev->omp_stacks.size = 0; pthread_mutex_init (&ptx_dev->omp_stacks.lock, NULL); + ptx_dev->rev_data = NULL; + return ptx_dev; } @@ -1179,8 +1187,10 @@ GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask) { int num_devices = nvptx_get_num_devices (); /* Return -1 if no omp_requires_mask cannot be fulfilled but - devices were present. */ - if (num_devices > 0 && omp_requires_mask != 0) + devices were present. Unified-shared address: see comment in + nvptx_open_device for CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING. */ + if (num_devices > 0 + && (omp_requires_mask & ~GOMP_REQUIRES_UNIFIED_ADDRESS) != 0) return -1; return num_devices; } @@ -1380,7 +1390,7 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data, else if (rev_fn_table) { CUdeviceptr var; - size_t bytes; + size_t bytes, i; r = CUDA_CALL_NOCHECK (cuModuleGetGlobal, &var, &bytes, module, "$offload_func_table"); if (r != CUDA_SUCCESS) @@ -1390,6 +1400,37 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data, r = CUDA_CALL_NOCHECK (cuMemcpyDtoH, *rev_fn_table, var, bytes); if (r != CUDA_SUCCESS) GOMP_PLUGIN_fatal ("cuMemcpyDtoH error: %s", cuda_error (r)); + /* Free if only NULL entries. */ + for (i = 0; i < fn_entries; ++i) + if ((*rev_fn_table)[i] != 0) + break; + if (i == fn_entries) + { + free (*rev_fn_table); + *rev_fn_table = NULL; + } + } + + if (rev_fn_table && *rev_fn_table && dev->rev_data == NULL) + { + /* cuMemHostAlloc memory is accessible on the device, if unified-shared + address is supported; this is assumed - see comment in + nvptx_open_device for CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING. */ + CUDA_CALL_ASSERT (cuMemHostAlloc, (void **) &dev->rev_data, + sizeof (*dev->rev_data), CU_MEMHOSTALLOC_DEVICEMAP); + CUdeviceptr dp = (CUdeviceptr) dev->rev_data; + CUdeviceptr device_rev_offload_var; + size_t device_rev_offload_size; + CUresult r = CUDA_CALL_NOCHECK (cuModuleGetGlobal, + &device_rev_offload_var, + &device_rev_offload_size, module, + XSTRING (GOMP_REV_OFFLOAD_VAR)); + if (r != CUDA_SUCCESS) + GOMP_PLUGIN_fatal ("cuModuleGetGlobal error - GOMP_REV_OFFLOAD_VAR: %s", cuda_error (r)); + r = CUDA_CALL_NOCHECK (cuMemcpyHtoD, device_rev_offload_var, &dp, + sizeof (dp)); + if (r != CUDA_SUCCESS) + GOMP_PLUGIN_fatal ("cuMemcpyHtoD error: %s", cuda_error (r)); } nvptx_set_clocktick (module, dev); @@ -2001,6 +2042,23 @@ nvptx_stacks_acquire (struct ptx_device *ptx_dev, size_t size, int num) return (void *) ptx_dev->omp_stacks.ptr; } + +void +rev_off_dev_to_host_cpy (void *dest, const void *src, size_t size, + CUstream stream) +{ + CUDA_CALL_ASSERT (cuMemcpyDtoHAsync, dest, (CUdeviceptr) src, size, stream); + CUDA_CALL_ASSERT (cuStreamSynchronize, stream); +} + +void +rev_off_host_to_dev_cpy (void *dest, const void *src, size_t size, + CUstream stream) +{ + CUDA_CALL_ASSERT (cuMemcpyHtoDAsync, (CUdeviceptr) dest, src, size, stream); + CUDA_CALL_ASSERT (cuStreamSynchronize, stream); +} + void GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args) { @@ -2035,6 +2093,8 @@ GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args) nvptx_adjust_launch_bounds (tgt_fn, ptx_dev, &teams, &threads); size_t stack_size = nvptx_stacks_size (); + bool reverse_offload = ptx_dev->rev_data != NULL; + CUstream copy_stream = NULL; pthread_mutex_lock (&ptx_dev->omp_stacks.lock); void *stacks = nvptx_stacks_acquire (ptx_dev, stack_size, teams * threads); @@ -2048,12 +2108,41 @@ GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args) GOMP_PLUGIN_debug (0, " %s: kernel %s: launch" " [(teams: %u), 1, 1] [(lanes: 32), (threads: %u), 1]\n", __FUNCTION__, fn_name, teams, threads); + if (reverse_offload) + CUDA_CALL_ASSERT (cuStreamCreate, ©_stream, CU_STREAM_NON_BLOCKING); r = CUDA_CALL_NOCHECK (cuLaunchKernel, function, teams, 1, 1, 32, threads, 1, 0, NULL, NULL, config); if (r != CUDA_SUCCESS) GOMP_PLUGIN_fatal ("cuLaunchKernel error: %s", cuda_error (r)); - - r = CUDA_CALL_NOCHECK (cuCtxSynchronize, ); + if (reverse_offload) + while (true) + { + r = CUDA_CALL_NOCHECK (cuStreamQuery, NULL); + if (r == CUDA_SUCCESS) + break; + if (r == CUDA_ERROR_LAUNCH_FAILED) + GOMP_PLUGIN_fatal ("cuStreamQuery error: %s %s\n", cuda_error (r), + maybe_abort_msg); + else if (r != CUDA_ERROR_NOT_READY) + GOMP_PLUGIN_fatal ("cuStreamQuery error: %s", cuda_error (r)); + + if (__atomic_load_n (&ptx_dev->rev_data->fn, __ATOMIC_ACQUIRE) != 0) + { + struct rev_offload *rev_data = ptx_dev->rev_data; + GOMP_PLUGIN_target_rev (rev_data->fn, rev_data->mapnum, + rev_data->addrs, rev_data->sizes, + rev_data->kinds, rev_data->dev_num, + rev_off_dev_to_host_cpy, + rev_off_host_to_dev_cpy, copy_stream); + CUDA_CALL_ASSERT (cuStreamSynchronize, copy_stream); + __atomic_store_n (&rev_data->fn, 0, __ATOMIC_RELEASE); + } + usleep (1); + } + else + r = CUDA_CALL_NOCHECK (cuCtxSynchronize, ); + if (reverse_offload) + CUDA_CALL_ASSERT (cuStreamDestroy, copy_stream); if (r == CUDA_ERROR_LAUNCH_FAILED) GOMP_PLUGIN_fatal ("cuCtxSynchronize error: %s %s\n", cuda_error (r), maybe_abort_msg); diff --git a/libgomp/target.c b/libgomp/target.c index 73c99c7c..189b201 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -2934,6 +2934,25 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum, htab_free (refcount_set); } +/* Handle reverse offload. This is called by the device plugins for a + reverse offload; it is not called if the outer target runs on the host. */ + +void +gomp_target_rev (uint64_t fn_ptr __attribute__ ((unused)), + uint64_t mapnum __attribute__ ((unused)), + uint64_t devaddrs_ptr __attribute__ ((unused)), + uint64_t sizes_ptr __attribute__ ((unused)), + uint64_t kinds_ptr __attribute__ ((unused)), + int dev_num __attribute__ ((unused)), + void (*dev_to_host_cpy) (void *, const void *, size_t, + void *) __attribute__ ((unused)), + void (*host_to_dev_cpy) (void *, const void *, size_t, + void *) __attribute__ ((unused)), + void *token __attribute__ ((unused))) +{ + __builtin_unreachable (); +} + /* Host fallback for GOMP_target_data{,_ext} routines. */ static void |