void
case_map_pop(struct bench_ctx *ctx)
{
struct map *m = map();
/* suite */
int i;
char keys[ctx->n][4];
for (i = 0; i < ctx->n; i++) {
sprintf(keys[i], "%d", i & 999);
map_set(m, keys[i], "val");
}
/* bench */
bench_ctx_reset_start_at(ctx);
for (i = 0; i < ctx->n; i++) {
map_pop(m, keys[i]);
}
bench_ctx_reset_end_at(ctx);
map_free(m);
}
void
nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
int async, unsigned *dims, void *targ_mem_desc)
{
struct targ_fn_descriptor *targ_fn = (struct targ_fn_descriptor *) fn;
CUfunction function;
CUresult r;
int i;
struct ptx_stream *dev_str;
void *kargs[1];
void *hp, *dp;
struct nvptx_thread *nvthd = nvptx_thread ();
const char *maybe_abort_msg = "(perhaps abort was called)";
function = targ_fn->fn;
dev_str = select_stream_for_async (async, pthread_self (), false, NULL);
assert (dev_str == nvthd->current_stream);
/* Initialize the launch dimensions. Typically this is constant,
provided by the device compiler, but we must permit runtime
values. */
for (i = 0; i != 3; i++)
if (targ_fn->launch->dim[i])
dims[i] = targ_fn->launch->dim[i];
/* This reserves a chunk of a pre-allocated page of memory mapped on both
the host and the device. HP is a host pointer to the new chunk, and DP is
the corresponding device pointer. */
map_push (dev_str, async, mapnum * sizeof (void *), &hp, &dp);
GOMP_PLUGIN_debug (0, " %s: prepare mappings\n", __FUNCTION__);
/* Copy the array of arguments to the mapped page. */
for (i = 0; i < mapnum; i++)
((void **) hp)[i] = devaddrs[i];
/* Copy the (device) pointers to arguments to the device (dp and hp might in
fact have the same value on a unified-memory system). */
r = cuMemcpy ((CUdeviceptr)dp, (CUdeviceptr)hp, mapnum * sizeof (void *));
if (r != CUDA_SUCCESS)
GOMP_PLUGIN_fatal ("cuMemcpy failed: %s", cuda_error (r));
GOMP_PLUGIN_debug (0, " %s: kernel %s: launch"
" gangs=%u, workers=%u, vectors=%u\n",
__FUNCTION__, targ_fn->launch->fn,
dims[0], dims[1], dims[2]);
// OpenACC CUDA
//
// num_gangs nctaid.x
// num_workers ntid.y
// vector length ntid.x
kargs[0] = &dp;
r = cuLaunchKernel (function,
dims[GOMP_DIM_GANG], 1, 1,
dims[GOMP_DIM_VECTOR], dims[GOMP_DIM_WORKER], 1,
0, dev_str->stream, kargs, 0);
if (r != CUDA_SUCCESS)
GOMP_PLUGIN_fatal ("cuLaunchKernel error: %s", cuda_error (r));
#ifndef DISABLE_ASYNC
if (async < acc_async_noval)
{
r = cuStreamSynchronize (dev_str->stream);
if (r == CUDA_ERROR_LAUNCH_FAILED)
GOMP_PLUGIN_fatal ("cuStreamSynchronize error: %s %s\n", cuda_error (r),
maybe_abort_msg);
else if (r != CUDA_SUCCESS)
GOMP_PLUGIN_fatal ("cuStreamSynchronize error: %s", cuda_error (r));
}
else
{
CUevent *e;
e = (CUevent *)GOMP_PLUGIN_malloc (sizeof (CUevent));
r = cuEventCreate (e, CU_EVENT_DISABLE_TIMING);
if (r == CUDA_ERROR_LAUNCH_FAILED)
GOMP_PLUGIN_fatal ("cuEventCreate error: %s %s\n", cuda_error (r),
maybe_abort_msg);
else if (r != CUDA_SUCCESS)
GOMP_PLUGIN_fatal ("cuEventCreate error: %s", cuda_error (r));
event_gc (true);
r = cuEventRecord (*e, dev_str->stream);
if (r != CUDA_SUCCESS)
GOMP_PLUGIN_fatal ("cuEventRecord error: %s", cuda_error (r));
event_add (PTX_EVT_KNL, e, (void *)dev_str);
}
#else
r = cuCtxSynchronize ();
if (r == CUDA_ERROR_LAUNCH_FAILED)
GOMP_PLUGIN_fatal ("cuCtxSynchronize error: %s %s\n", cuda_error (r),
maybe_abort_msg);
else if (r != CUDA_SUCCESS)
GOMP_PLUGIN_fatal ("cuCtxSynchronize error: %s", cuda_error (r));
#endif
GOMP_PLUGIN_debug (0, " %s: kernel %s: finished\n", __FUNCTION__,
targ_fn->launch->fn);
#ifndef DISABLE_ASYNC
if (async < acc_async_noval)
#endif
map_pop (dev_str);
}
static void
event_gc (bool memmap_lockable)
{
struct ptx_event *ptx_event = ptx_events;
struct nvptx_thread *nvthd = nvptx_thread ();
pthread_mutex_lock (&ptx_event_lock);
while (ptx_event != NULL)
{
CUresult r;
struct ptx_event *e = ptx_event;
ptx_event = ptx_event->next;
if (e->ord != nvthd->ptx_dev->ord)
continue;
r = cuEventQuery (*e->evt);
if (r == CUDA_SUCCESS)
{
CUevent *te;
te = e->evt;
switch (e->type)
{
case PTX_EVT_MEM:
case PTX_EVT_SYNC:
break;
case PTX_EVT_KNL:
map_pop (e->addr);
break;
case PTX_EVT_ASYNC_CLEANUP:
{
/* The function gomp_plugin_async_unmap_vars needs to claim the
memory-map splay tree lock for the current device, so we
can't call it when one of our callers has already claimed
the lock. In that case, just delay the GC for this event
until later. */
if (!memmap_lockable)
continue;
GOMP_PLUGIN_async_unmap_vars (e->addr);
}
break;
}
cuEventDestroy (*te);
free ((void *)te);
if (ptx_events == e)
ptx_events = ptx_events->next;
else
{
struct ptx_event *e_ = ptx_events;
while (e_->next != e)
e_ = e_->next;
e_->next = e_->next->next;
}
free (e);
}
}
pthread_mutex_unlock (&ptx_event_lock);
}
