[gomp4] OpenACC async re-work

On 2017/6/29 6:31 AM, Cesar Philippidis wrote:
> On 06/27/2017 03:56 AM, Chung-Lin Tang wrote:
>> On 2017/6/27 6:45 AM, Cesar Philippidis wrote:
>>>> (1) Instead of essentially implementing the entire OpenACC async support
>>>> inside the plugin, we now use an opaque 'goacc_asyncqueue' implemented
>>>> by the plugin, along with core 'test', 'synchronize', 'serialize', etc.
>>>> plugin functions. Most of the OpenACC specific logic is pulled into
>>>> libgomp/oacc-async.c
>>> I'm not sure if plugins need to maintain backwards compatibility.
>>> However, I don't see any changes inside libgomp.map, so maybe it's not
>>> required.
>>
>> This patch is pretty large, but only inner workings (including libgomp vs. plugin interface) were modified.
>> No user API compatibility was altered.
>>
>>>> (3) For 'wait + async', we now add a local thread synchronize, instead
>>>> of just ordering the streams.
>>>>
>>>> (4) To work with the (3) change, some front end changes were added to
>>>> propagate argument-less wait clauses as 'wait(GOACC_ASYNC_NOVAL)' to
>>>> represent a 'wait all'.
>>> What's the significance of GOMP_ASYNC_NOVAL? Wouldn't it have been
>>> easier to make that change in the gimplifier?
>>
>> Actually, we were basically throwing away argument-less wait clauses in front-ends
>> before this patch; i.e. '#pragma acc parallel async' and '#pragma acc parallel wait async'
>> were internally the same.
>>
>> The use of GOMP_ASYNC_NOVAL (-1) was just following the current 'async' clause representation
>> convention.
> 
> So then then wait was implied before? Or maybe that's why 'wait async'
> didn't work.

Not really, 'wait async' didn't work because we were not synchronizing
the local thread.

>>> Overall, I like how you were able eliminate the externally managed map_*
>>> data structure which was used to pass in arguments to nvptx_exec.
>>> Although I wonder if we should just pass in those individual arguments
>>> directly to cuLaunchKernel. But that's a big change in itself.
>>
>> I didn't think of that when working on the current patch, maybe later.
> 
> Here's some more comments regarding the code below. One high-level
> comment regarding the usage of async-specific locks. Can't you get by
> with using the global device lock, instead of a special async queue or
> would that cause a deadlock?

I did not try to re-use the device lock here. I think using that might
cause more ordering/nesting problems (gomp mutexes are not pthread mutexes,
no 'try' mechanism here).

The intention of the async queue locks are just very short temporal
protection when looking up a struct goacc_asyncqueue. Also see below comments.

>> +attribute_hidden struct goacc_asyncqueue *
>> +lookup_goacc_asyncqueue (struct goacc_thread *thr, bool create, int
> async)
>> +{
>> +  /* The special value acc_async_noval (-1) maps to the thread-specific
>> +     default async stream.  */
>> +  if (async == acc_async_noval)
>> +    async = thr->default_async;
> 
> Is the default async queue device independent? I thought the default
> async queue is defined in the acc_async_t enum. Maybe set
> 
>  async = acc_async_default

It's now thread dependent and user settable in OpenACC 2.5
(i.e. the acc_set/get_default_async user API routines)

>> +  if (async == acc_async_sync)
>> +    return NULL;
>> +
>> +  if (async < 0)
>> +    gomp_fatal ("bad async %d", async);
>> +
>> +  struct gomp_device_descr *dev = thr->dev;
>> +
>> +  if (!create
>> +      && (async >= dev->openacc.async.nasyncqueue
>> +	  || !dev->openacc.async.asyncqueue[async]))
>> +    return NULL;
>> +
>> +  gomp_mutex_lock (&dev->openacc.async.lock);
> Is this lock sufficient? What happens if the device is released?
> 
>> +  if (async >= dev->openacc.async.nasyncqueue)
>> +    {
> Not your fault, but I wonder if we would be better off just hard-capping
> the number of async queues. What happens if the user does something like
> wait (1<<30)? That can be addressed later.

Yeah, looks like we should use some sort of map or hash (above some threshold)
instead.

>>  int
>>  acc_async_test (int async)
>>  {
>> @@ -54,15 +129,14 @@ acc_async_test (int async)
>>    if (!thr || !thr->dev)
>>      gomp_fatal ("no device active");
>>
>> -  int res = thr->dev->openacc.async_test_func (async);
>> -
>>    if (profiling_setup_p)
>>      {
>>        thr->prof_info = NULL;
>>        thr->api_info = NULL;
>>      }
>> -
>> -  return res;
>> +
>> +  goacc_aq aq = lookup_goacc_asyncqueue (thr, true, async);
>> +  return thr->dev->openacc.async.test_func (aq);
> I'm not sure how the profling stuff works. Should the profling state be
> state be set after calling thr->dev->openacc.async.test_func?

I've updated the patch to place those profiling code (which seems to be
end of routine cleanups) to the very end.

>> -  void *ret = NULL;
>>    if (thr && thr->dev && thr->dev->openacc.cuda.get_stream_func)
>> -    ret = thr->dev->openacc.cuda.get_stream_func (async);
>> -
>> -  if (profiling_setup_p)
>>      {
>> -      thr->prof_info = NULL;
>> -      thr->api_info = NULL;
>> +      goacc_aq aq = lookup_goacc_asyncqueue (thr, false, async);
>> +      return aq ? thr->dev->openacc.cuda.get_stream_func (aq) : NULL;
> Again, strange ordering fo profiling_setup_p.
> 
> By the way, why not use get_goacc_thread here and other places in this
> function? Again, that's a problem for another day.

get_goacc_thread immediately gomp_fatal()s if the 'thr' or device is NULL,
while the original behavior here fails more elegantly. The supposed behavior
can be discussed later, but here I'm just trying to preserve things
as much as possible.

>> -  if (async > acc_async_sync)
>> -    acc_dev->openacc.async_set_async_func (async);
>> +  goacc_aq aq = get_goacc_asyncqueue (async);
>>
>>    if (is_dev)
>> -    acc_dev->host2dev_func (acc_dev->target_id, d, h, s);
>> +    gomp_copy_host2dev (acc_dev, aq, d, h, s);
>>    else
>> -    acc_dev->dev2host_func (acc_dev->target_id, h, d, s);
>> +    gomp_copy_dev2host (acc_dev, aq, h, d, s);
>>
>> -  if (async > acc_async_sync)
>> -    acc_dev->openacc.async_set_async_func (acc_async_sync);
>> -
> Why did you remove this, but not add a clal to set_goacc_asyncqueue?
> Maybe it's redundant.

get_goacc_asyncqueue acquires the async queue lock internally. This is
related to your above comment on sharing the device lock. IMHO, this
may cause more complex lock ordering issues, but might be revisited later
to see if that simplifies things.

>> Index: libgomp/oacc-parallel.c
>> ===================================================================
>> --- libgomp/oacc-parallel.c	(revision 249620)
>> +++ libgomp/oacc-parallel.c	(working copy)
>> @@ -215,7 +215,9 @@ GOACC_parallel_keyed (int device, void (*fn) (void
>>        fn (hostaddrs);
>>        goto out;
>>      }
>> -
>> +  else if (profiling_dispatch_p)
>> +    api_info.device_api = acc_device_api_cuda;
>> +
> That seems target specific. Does that belong in the generic code path?

I don't really understand the profiling code. This change appeared to
fix regressions. Thomas might know better.

Attached is the updated version of the patch, re-tested.

Thomas, do you need some more time to look over it? Or should I commit it first?

Thanks,
Chung-Lin

[gomp4] OpenACC async re-work

Commit Message

Comments

Patch