OMP: zero-length array sections

[Thomas Schwinge]

Hi Jakub!

Andrew and I are currently trying to sort out zero-length array sections
behavior in the OpenACC context.  From that fell out a testcase that I
created to learn/verify (some of?) the OpenMP behavior.  OK to push?


Oh, and any thought on this:

    int dev = omp_get_default_device();
    bool shared_mem = false;
    #pragma omp target map(alloc:shared_mem)
    shared_mem = true;
    #if 1 //TODO
    /* 'omp_target_is_present' (and a few others) have behavior different from
       most others, where 'devicep == NULL' is handled the same as 'device_num 
==
       GOMP_DEVICE_HOST_FALLBACK'.  Is that difference intentional?  */
    /* Given that GCC doesn't support shared-memory offload devices, we fake the
       expected thing as follows.  */
    if (shared_mem)
      dev = omp_get_initial_device(); // GOMP_DEVICE_HOST_FALLBACK
    #endif


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander 
Walter

#include <assert.h>
#include <stdbool.h>
#include <stdlib.h>
#include <omp.h>

static int dev;
static bool shared_mem = false;

static void test(bool mapped_outside, int *a, int n_map, int n_iterate)
{
#pragma omp target enter data map(alloc:a[0:n_map])
  if (!mapped_outside && n_map == 0)
    assert(!!omp_target_is_present(a, dev) == shared_mem);
  else
    assert(omp_target_is_present(a, dev));

  {
    bool a_NULL;
#pragma omp target map(from:a[0:n_map]) map(from:a_NULL)
    {
      a_NULL = (a == NULL);

      for (int i = 0; i < n_iterate; ++i)
	a[i] = n_iterate + i;
    }
    if (!shared_mem && !mapped_outside && n_map == 0)
      assert(a_NULL);
    else
      assert(!a_NULL);
  }

#pragma omp target exit data map(from:a[0:n_map])
  if (mapped_outside)
    assert(omp_target_is_present(a, dev));
  else
    assert(!!omp_target_is_present(a, dev) == shared_mem);
}

int main(int argc, char *argv[])
{
  dev = omp_get_default_device();
#pragma omp target map(alloc:shared_mem)
  shared_mem = true;
#if 1 //TODO
  /* 'omp_target_is_present' (and a few others) have behavior different from
     most others, where 'devicep == NULL' is handled the same as 'device_num ==
     GOMP_DEVICE_HOST_FALLBACK'.  Is that difference intentional?  */
  /* Given that GCC doesn't support shared-memory offload devices, we fake the
     expected thing as follows.  */
  if (shared_mem)
    dev = omp_get_initial_device(); // GOMP_DEVICE_HOST_FALLBACK
#endif

  const int n = 38;

  int *a = malloc(n * sizeof *a);
  assert(a);
  assert(!!omp_target_is_present(a, dev) == shared_mem);

  for (int i = 0; i < n; ++i)
    a[i] = -i;
  test(false, a, n, n);
  for (int i = 0; i < n; ++i)
    assert(a[i] == n + i);

  for (int i = 0; i < n; ++i)
    a[i] = -i;
  test(false, a, 0, 0);
  for (int i = 0; i < n; ++i)
    assert(a[i] == -i);

  for (int i = 0; i < n; ++i)
    a[i] = -i;
#pragma omp target enter data map(alloc:a[0:n])
  assert(omp_target_is_present(a, dev));
  test(true, a, n, n);
  for (int i = 0; i < n; ++i)
    if (shared_mem)
      assert(a[i] == n + i);
    else
      assert(a[i] == -i);
#pragma omp target exit data map(from:a[0:n])
  assert(!!omp_target_is_present(a, dev) == shared_mem);
  for (int i = 0; i < n; ++i)
    assert(a[i] == n + i);

  for (int i = 0; i < n; ++i)
    a[i] = -i;
#pragma omp target enter data map(alloc:a[0:n])
  assert(omp_target_is_present(a, dev));
  test(true, a, 0, n);
  for (int i = 0; i < n; ++i)
    if (shared_mem)
      assert(a[i] == n + i);
    else
      assert(a[i] == -i);
#pragma omp target exit data map(from:a[0:n])
  assert(!!omp_target_is_present(a, dev) == shared_mem);
  for (int i = 0; i < n; ++i)
    assert(a[i] == n + i);

  free(a);

  return 0;
}