openmp/libomptarget/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
   6 #include <omp.h>
   7 #include <stdio.h>
   8
   9 int main() {
  10   int x;
  11 #pragma omp target parallel map(from : x)
  12   {
  13     int *buf = llvm_omp_target_dynamic_shared_alloc() + 252;
  14 #pragma omp barrier
  15     if (omp_get_thread_num() == 0)
  16       *buf = 1;
  17 #pragma omp barrier
  18     if (omp_get_thread_num() == 1)
  19       x = *buf;
  20   }
  21
  22   // CHECK: PASS
  23   if (x == 1 && llvm_omp_target_dynamic_shared_alloc() == NULL)
  24     printf("PASS\n");
  25 }