offload/test/offloading/dyn_groupprivate.cpp

// RUN: %libomptarget-compilexx-generic -fopenmp-version=61
// RUN: %libomptarget-run-generic | %fcheck-generic
// RUN: %libomptarget-compileoptxx-generic -fopenmp-version=61
// RUN: %libomptarget-run-generic | %fcheck-generic
// REQUIRES: gpu
// XFAIL: intelgpu

#include <omp.h>
#include <stdio.h>

#define N 512

int main() {
  int Result[N], NumThreads;

// Verify the groupprivate buffer works as expected.
#pragma omp target teams num_teams(1) thread_limit(N)                          \
    dyn_groupprivate(fallback(abort) : N * sizeof(Result[0]))                  \
    map(from : Result, NumThreads)
  {
    int Buffer[N];
#pragma omp parallel
    {
      int *DynBuffer = (int *)omp_get_dyn_gprivate_nofb_ptr();
      int TId = omp_get_thread_num();
      if (TId == 0)
        NumThreads = omp_get_num_threads();
      Buffer[TId] = 7;
      DynBuffer[TId] = 3;
#pragma omp barrier
      int WrappedTId = (TId + 37) % NumThreads;
      Result[TId] = Buffer[WrappedTId] + DynBuffer[WrappedTId];
    }
  }

  if (NumThreads < N / 2 || NumThreads > N) {
    printf("Expected number of threads to be in [%i:%i], but got: %i", N / 2, N,
           NumThreads);
    return -1;
  }

  int Failed = 0;
  for (int i = 0; i < NumThreads; ++i) {
    if (Result[i] != 7 + 3) {
      printf("Result[%i] is %i, expected %i\n", i, Result[i], 7 + 3);
      ++Failed;
    }
  }

  // Verify that the routines in the host returns NULL and zero.
  if (omp_get_dyn_gprivate_ptr())
    ++Failed;
  if (omp_get_dyn_gprivate_nofb_ptr())
    ++Failed;
  if (omp_get_dyn_gprivate_size())
    ++Failed;

  size_t MaxSize = omp_get_gprivate_limit(0, omp_access_cgroup);
  size_t ExceededSize = MaxSize + 10;

// Verify that the fallback(default_mem) modifier works.
#pragma omp target dyn_groupprivate(fallback(default_mem) : ExceededSize)      \
    map(tofrom : Failed)
  {
    if (!omp_get_dyn_gprivate_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_ptr(0) == omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (!omp_get_dyn_gprivate_size())
      ++Failed;
    if (omp_get_dyn_gprivate_size() != ExceededSize)
      ++Failed;
    if (omp_get_dyn_gprivate_memspace() != omp_default_mem_space)
      ++Failed;
  }

// Verify that the fallback(null) modifier works.
#pragma omp target dyn_groupprivate(fallback(null) : ExceededSize)             \
    map(tofrom : Failed)
  {
    if (omp_get_dyn_gprivate_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_ptr(0) != omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_size())
      ++Failed;
    if (omp_get_dyn_gprivate_memspace() != omp_null_mem_space)
      ++Failed;
  }

// Verify that the default modifier is fallback(default_mem).
#pragma omp target dyn_groupprivate(ExceededSize)
  {
    if (!omp_get_dyn_gprivate_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_ptr(0) == omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (!omp_get_dyn_gprivate_size())
      ++Failed;
    if (omp_get_dyn_gprivate_size() != ExceededSize)
      ++Failed;
    if (omp_get_dyn_gprivate_memspace() != omp_default_mem_space)
      ++Failed;
  }

  // The dynamic size alone fits, but not together with the kernel's static
  // block memory, e.g., the device runtime's team state.
  size_t AlmostMaxSize = MaxSize - 1;

// Verify that the fallback(default_mem) also triggers in that case.
#pragma omp target dyn_groupprivate(fallback(default_mem) : AlmostMaxSize)     \
    map(tofrom : Failed)
  {
    if (!omp_get_dyn_gprivate_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_size() != AlmostMaxSize)
      ++Failed;
    if (omp_get_dyn_gprivate_memspace() != omp_default_mem_space)
      ++Failed;
  }

// Verify that the fallback(abort) modifier works.
#pragma omp target dyn_groupprivate(fallback(abort) : N) map(tofrom : Failed)
  {
    if (!omp_get_dyn_gprivate_ptr(0))
      ++Failed;
    if (!omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_ptr(0) != omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_ptr(5) != omp_get_dyn_gprivate_nofb_ptr(5))
      ++Failed;
    if (!omp_get_dyn_gprivate_size())
      ++Failed;
    if (omp_get_dyn_gprivate_size() != N)
      ++Failed;
    if (omp_get_dyn_gprivate_memspace() != omp_cgroup_mem_space)
      ++Failed;
  }

// Verify that the fallback(default_mem) does not trigger when not needed.
#pragma omp target dyn_groupprivate(fallback(default_mem) : N)                 \
    map(tofrom : Failed)
  {
    if (!omp_get_dyn_gprivate_ptr(0))
      ++Failed;
    if (!omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_ptr(0) != omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (!omp_get_dyn_gprivate_size())
      ++Failed;
    if (omp_get_dyn_gprivate_size() != N)
      ++Failed;
    if (omp_get_dyn_gprivate_memspace() != omp_cgroup_mem_space)
      ++Failed;
  }

// Verify that the clause works when passing a zero size.
#pragma omp target dyn_groupprivate(fallback(abort) : 0) map(tofrom : Failed)
  {
    if (omp_get_dyn_gprivate_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_ptr(0) != omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_size())
      ++Failed;
    if (omp_get_dyn_gprivate_memspace() != omp_null_mem_space)
      ++Failed;
  }

// Verify that the clause works when passing a zero size and
// fallback(default_mem).
#pragma omp target dyn_groupprivate(fallback(default_mem) : 0)                 \
    map(tofrom : Failed)
  {
    if (omp_get_dyn_gprivate_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_ptr(0) != omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_size())
      ++Failed;
    if (omp_get_dyn_gprivate_memspace() != omp_null_mem_space)
      ++Failed;
  }

// Verify that omitting the clause is the same as setting zero size.
#pragma omp target map(tofrom : Failed)
  {
    if (omp_get_dyn_gprivate_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_ptr(0) != omp_get_dyn_gprivate_nofb_ptr(0))
      ++Failed;
    if (omp_get_dyn_gprivate_size())
      ++Failed;
    if (omp_get_dyn_gprivate_memspace() != omp_null_mem_space)
      ++Failed;
  }

  // CHECK: PASS
  if (!Failed)
    printf("PASS\n");
}