Re: [gomp4] Redesign oacc_parallel launch API

[Cesar Philippidis <cesar_philippidis@mentor.com>]

[-- Attachment #1: Type: text/plain, Size: 1891 bytes --]

On 07/28/2015 09:52 AM, Nathan Sidwell wrote:
> I've committed this patch to the gomp4 branch to redo the launch API. 
> I'll post a version for trunk once the versioning patch gets approved &
> committed.
> 
> This changes the API in a number of ways, allowing device-specific
> knowledge to be moved into the device compiler and out of the host
> compiler.
> 
> Firstly, we attach a tuple of launch dimensions as an attribute to the
> offloaded function's 'oacc function' attribute.  These are the constant
> launch dimensions.  Dynamic dimensions get a zero for their slot in this
> list.  Further this list can be extended in the future to an alist keyed
> by device_type.
> 
> Dynamic dimensions are computed on the host.  however they are passed
> via varadic args to the GOACC_parallel function (which is renamed).  The
> varadic args are passed using key/value representation, and 3 keys are
> currently defined:
> END -- end of the varadic list
> DIM - set of runtime-computed dimensions.  Only the dynamic ones are
> passed.
> ASYNC_WAIT - an async and a set of waits (possibly zero).
> 
> I have arranged for the key to have a slot that can later be filled by
> device_type, and hence support multiple device types.
> 
> The constant dimensions can be used in expansion of the GOACC_nid
> function in the device compiler.  The device compiler could also process
> that list to select the device_type slot that is appropriate.
> 
> For PTX the backend is augmented to emit the launch dimensions into the
> target data, from whence the ptx plugin can pick them up and overwrite
> with any dynamic ones passed in from the launch function.

Looking at set_oacc_fn_attrib, it appears that const values are also
considered dynamic. See the attached test case more more info. Is that
the expected behavior? If not, I could take a look at this after I
finished my reduction patch.

Cesar

[-- Attachment #2: vlength.c --]
[-- Type: text/x-csrc, Size: 250 bytes --]

#include <stdio.h>

const int vl = 32;

int
main ()
{
  unsigned int red = 0;

#pragma acc parallel loop vector_length (vl) vector reduction (+:red) copy (red)
  for (int i = 0; i < 100; i++)
    red ++;

  printf ("red = %d\n", red);

  return 0;
}