offload/test/api/omp_dynamic_shared_memory_amdgpu.c

   1 // RUN: %libomptarget-compile-amdgcn-amd-amdhsa -O1 -mllvm -openmp-opt-inline-device
   2 // RUN: env LIBOMPTARGET_SHARED_MEMORY_SIZE=256 \
   3 // RUN:   %libomptarget-run-amdgcn-amd-amdhsa | %fcheck-amdgcn-amd-amdhsa
   4 // REQUIRES: amdgcn-amd-amdhsa
   5 // XFAIL: amdgcn-amd-amdhsa
   6
   7 #include <omp.h>
   8 #include <stdio.h>
   9
  10 int main() {
  11   int x;
  12 #pragma omp target parallel map(from : x)
  13   {
  14     int *buf = llvm_omp_target_dynamic_shared_alloc() + 252;
  15 #pragma omp barrier
  16     if (omp_get_thread_num() == 0)
  17       *buf = 1;
  18 #pragma omp barrier
  19     if (omp_get_thread_num() == 1)
  20       x = *buf;
  21   }
  22
  23   // CHECK: PASS
  24   if (x == 1 && llvm_omp_target_dynamic_shared_alloc() == NULL)
  25     printf("PASS\n");
  26 }