1 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda -fopenmp-target-new-runtime
2 // RUN: env LIBOMPTARGET_SHARED_MEMORY_SIZE=4 \
3 // RUN:   %libomptarget-run-nvptx64-nvidia-cuda | %fcheck-nvptx64-nvidia-cuda
4 // REQUIRES: nvptx64-nvidia-cuda
5 
6 #include <omp.h>
7 #include <stdio.h>
8 
9 void *get_dynamic_shared() { return NULL; }
10 #pragma omp begin declare variant match(device = {arch(nvptx64)})
11 extern void *__kmpc_get_dynamic_shared();
12 void *get_dynamic_shared() { return __kmpc_get_dynamic_shared(); }
13 #pragma omp end declare variant
14 
15 int main() {
16   int x;
17 #pragma omp target parallel map(from : x)
18   {
19     int *buf = get_dynamic_shared();
20 #pragma omp barrier
21     if (omp_get_thread_num() == 0)
22       *buf = 1;
23 #pragma omp barrier
24     if (omp_get_thread_num() == 1)
25       x = *buf;
26   }
27 
28   // CHECK: PASS
29   if (x == 1)
30     printf("PASS\n");
31 }
32