Re: [PATCH, OpenMP] Implement local clause for declare target

[Tobias Burnus]

Tobias Burnus:

And it has another effect: with 'requires self_maps', the device has
its own copy of the variable - which might permit for much faster
access than accessing the host variable.

I tried the following code - and it fails (when a suitable offload
device is available):

Re-reading the OpenMP specification and what Deepak wrote in omp-lang,
I created a new testcase.

I am not sure whether it clashes with a previous testcase or not, but
it currently fails with both 'requires self_maps' and without.

The problem is that the local variable has predefined mapping such that:


(A) "A variable that is a groupprivate variable or a
     device-local variable must not appear as a list item in a 'map' clause."

Can you add an error diagnostic and a a test case for this?
Currently, it is accepted ...


Note for Fortran, this is handled correctly:

    4 | !$omp target map(x)
      |                  1
Error: ‘x’ argument to MAP clause at (1) must not be a device-local variable, 
including GROUPPRIVATE



(B) For

    #pragma omp target
      addr = &device_local_var[0];

the address is always the local address of the global variable and not
of some mapped variable.

See attached testcase for the second item.


Tobias

#if USE_SELF_MAPS
  #pragma omp self_maps
#endif

int var[2] = {5,6};
#pragma omp declare target local(var)

int get_var_inc () {
  var[0]++; return var[0];
}

__UINTPTR_TYPE__ get_addr () {
  return (__UINTPTR_TYPE__) &var[0];
}

int main()
{
  __UINTPTR_TYPE__ addr_tgt, addr_fn;
  bool is_initial_dev;
  int val;
  var[0] = 7;

  #pragma omp target map(from: val, is_initial_dev, addr_tgt, addr_fn)
  {
    is_initial_dev = __builtin_omp_is_initial_device ();
    addr_tgt = (__UINTPTR_TYPE__) &var[0];
    addr_fn = get_addr ();

    var[0] += 5;
    val = get_var_inc ();
    // Those are the device-local variable, i.e
    // - on the host:   7 + 5 + 1 -> val == var[0] = 13
    // - on the device: 5 + 5 + 1 -> val == var[0] == 11
    // This does not depend on the self-map clause
  }

  __builtin_printf ("var = %d, val = %d, is_initial_dev = %d, %p, %p, %p\n",
		    var[0], val, is_initial_dev, &var[0],
		    (void*) addr_tgt, (void*) addr_fn);

  if ((is_initial_dev && var[0] != 13) || (!is_initial_dev && var[0] != 11))
    __builtin_abort ();

  if (val != var[0])
    __builtin_abort ();

  if (addr_tgt != addr_fn)
    __builtin_abort ();


  if ((is_initial_dev && &var[0] != (void*) addr_fn)
      || (!is_initial_dev && &var[0] == (void*) addr_fn))
    __builtin_abort ();
}