files (via 'CRT_BEGIN' and 'CRT_END'): 'crtbegin.o' and 'crtend.o', but we
do so anyway, for symmetry with other configurations. */
+
+/* See 'crt0.c', 'mgomp.c'. */
+#if defined(__nvptx_softstack__) && defined(__nvptx_unisimt__)
+extern void *__nvptx_stacks[32] __attribute__((shared,nocommon));
+extern unsigned __nvptx_uni[32] __attribute__((shared,nocommon));
+#endif
+
+
#ifdef CRT_BEGIN
void
DO_GLOBAL_CTORS_BODY;
}
+/* Need '.entry' wrapper for offloading. */
+
+# if defined(__nvptx_softstack__) && defined(__nvptx_unisimt__)
+
+__attribute__((kernel)) void __do_global_ctors__entry__mgomp (void *);
+
+void
+__do_global_ctors__entry__mgomp (void *nvptx_stacks_0)
+{
+ __nvptx_stacks[0] = nvptx_stacks_0;
+ __nvptx_uni[0] = 0;
+
+ __do_global_ctors ();
+}
+
+# else
+
+__attribute__((kernel)) void __do_global_ctors__entry (void);
+
+void
+__do_global_ctors__entry (void)
+{
+ __do_global_ctors ();
+}
+
+# endif
+
#elif defined(CRT_END) /* ! CRT_BEGIN */
void
/* In this configuration here, there's no way that "this routine is run more
than once [...] when exit is called recursively": for nvptx target, the
call to '__do_global_dtors' is registered via 'atexit', which doesn't
- re-enter a function already run.
+ re-enter a function already run, and neither does nvptx offload target.
Therefore, we do *not* "arrange to remember where in the list we left off
processing". */
func_ptr *p;
(*p++) ();
}
+/* Need '.entry' wrapper for offloading. */
+
+# if defined(__nvptx_softstack__) && defined(__nvptx_unisimt__)
+
+__attribute__((kernel)) void __do_global_dtors__entry__mgomp (void *);
+
+void
+__do_global_dtors__entry__mgomp (void *nvptx_stacks_0)
+{
+ __nvptx_stacks[0] = nvptx_stacks_0;
+ __nvptx_uni[0] = 0;
+
+ __do_global_dtors ();
+}
+
+# else
+
+__attribute__((kernel)) void __do_global_dtors__entry (void);
+
+void
+__do_global_dtors__entry (void)
+{
+ __do_global_dtors ();
+}
+
+# endif
+
#else /* ! CRT_BEGIN && ! CRT_END */
#error "One of CRT_BEGIN or CRT_END must be defined."
#endif
default is set here. */
static unsigned lowlat_pool_size = 8*1024;
+static bool nvptx_do_global_cdtors (CUmodule, struct ptx_device *,
+ const char *);
+static size_t nvptx_stacks_size ();
+static void *nvptx_stacks_acquire (struct ptx_device *, size_t, int);
+
static inline struct nvptx_thread *
nvptx_thread (void)
{
if (!ptx_dev)
return true;
+ bool ret = true;
+
+ for (struct ptx_image_data *image = ptx_dev->images;
+ image != NULL;
+ image = image->next)
+ {
+ if (!nvptx_do_global_cdtors (image->module, ptx_dev,
+ "__do_global_dtors__entry"))
+ ret = false;
+ }
+
for (struct ptx_free_block *b = ptx_dev->free_blocks; b;)
{
struct ptx_free_block *b_next = b->next;
CUDA_CALL (cuCtxDestroy, ptx_dev->ctx);
free (ptx_dev);
- return true;
+
+ return ret;
}
static int
GOMP_PLUGIN_fatal ("cuMemcpyHtoD error: %s", cuda_error (r));
}
+/* Invoke MODULE's global constructors/destructors. */
+
+static bool
+nvptx_do_global_cdtors (CUmodule module, struct ptx_device *ptx_dev,
+ const char *funcname)
+{
+ bool ret = true;
+ char *funcname_mgomp = NULL;
+ CUresult r;
+ CUfunction funcptr;
+ r = CUDA_CALL_NOCHECK (cuModuleGetFunction,
+ &funcptr, module, funcname);
+ GOMP_PLUGIN_debug (0, "cuModuleGetFunction (%s): %s\n",
+ funcname, cuda_error (r));
+ if (r == CUDA_ERROR_NOT_FOUND)
+ {
+ /* Try '[funcname]__mgomp'. */
+
+ size_t funcname_len = strlen (funcname);
+ const char *mgomp_suffix = "__mgomp";
+ size_t mgomp_suffix_len = strlen (mgomp_suffix);
+ funcname_mgomp
+ = GOMP_PLUGIN_malloc (funcname_len + mgomp_suffix_len + 1);
+ memcpy (funcname_mgomp, funcname, funcname_len);
+ memcpy (funcname_mgomp + funcname_len,
+ mgomp_suffix, mgomp_suffix_len + 1);
+ funcname = funcname_mgomp;
+
+ r = CUDA_CALL_NOCHECK (cuModuleGetFunction,
+ &funcptr, module, funcname);
+ GOMP_PLUGIN_debug (0, "cuModuleGetFunction (%s): %s\n",
+ funcname, cuda_error (r));
+ }
+ if (r == CUDA_ERROR_NOT_FOUND)
+ ;
+ else if (r != CUDA_SUCCESS)
+ {
+ GOMP_PLUGIN_error ("cuModuleGetFunction (%s) error: %s",
+ funcname, cuda_error (r));
+ ret = false;
+ }
+ else
+ {
+ /* If necessary, set up soft stack. */
+ void *nvptx_stacks_0;
+ void *kargs[1];
+ if (funcname_mgomp)
+ {
+ size_t stack_size = nvptx_stacks_size ();
+ pthread_mutex_lock (&ptx_dev->omp_stacks.lock);
+ nvptx_stacks_0 = nvptx_stacks_acquire (ptx_dev, stack_size, 1);
+ nvptx_stacks_0 += stack_size;
+ kargs[0] = &nvptx_stacks_0;
+ }
+ r = CUDA_CALL_NOCHECK (cuLaunchKernel,
+ funcptr,
+ 1, 1, 1, 1, 1, 1,
+ /* sharedMemBytes */ 0,
+ /* hStream */ NULL,
+ /* kernelParams */ funcname_mgomp ? kargs : NULL,
+ /* extra */ NULL);
+ if (r != CUDA_SUCCESS)
+ {
+ GOMP_PLUGIN_error ("cuLaunchKernel (%s) error: %s",
+ funcname, cuda_error (r));
+ ret = false;
+ }
+
+ r = CUDA_CALL_NOCHECK (cuStreamSynchronize,
+ NULL);
+ if (r != CUDA_SUCCESS)
+ {
+ GOMP_PLUGIN_error ("cuStreamSynchronize (%s) error: %s",
+ funcname, cuda_error (r));
+ ret = false;
+ }
+
+ if (funcname_mgomp)
+ pthread_mutex_unlock (&ptx_dev->omp_stacks.lock);
+ }
+
+ if (funcname_mgomp)
+ free (funcname_mgomp);
+
+ return ret;
+}
+
/* Load the (partial) program described by TARGET_DATA to device
number ORD. Allocate and return TARGET_TABLE. If not NULL, REV_FN_TABLE
will contain the on-device addresses of the functions for reverse offload.
nvptx_set_clocktick (module, dev);
+ if (!nvptx_do_global_cdtors (module, dev, "__do_global_ctors__entry"))
+ return -1;
+
return fn_entries + var_entries + other_entries;
}
for (prev_p = &dev->images; (image = *prev_p) != 0; prev_p = &image->next)
if (image->target_data == target_data)
{
+ if (!nvptx_do_global_cdtors (image->module, dev,
+ "__do_global_dtors__entry"))
+ ret = false;
+
*prev_p = image->next;
if (CUDA_CALL_NOCHECK (cuModuleUnload, image->module) != CUDA_SUCCESS)
ret = false;
projects / thirdparty / gcc.git / commitdiff
