1*dcc27ea4SJoseph Huber // RUN: %libomptarget-compile-amdgcn-amd-amdhsa -O2 -mllvm \ 2*dcc27ea4SJoseph Huber // RUN: -openmp-opt-inline-device 3330d8983SJohannes Doerfert // RUN: env LIBOMPTARGET_SHARED_MEMORY_SIZE=256 \ 4330d8983SJohannes Doerfert // RUN: %libomptarget-run-amdgcn-amd-amdhsa | %fcheck-amdgcn-amd-amdhsa 5330d8983SJohannes Doerfert // REQUIRES: amdgcn-amd-amdhsa 6330d8983SJohannes Doerfert 7330d8983SJohannes Doerfert #include <omp.h> 8330d8983SJohannes Doerfert #include <stdio.h> 9330d8983SJohannes Doerfert 10330d8983SJohannes Doerfert int main() { 11330d8983SJohannes Doerfert int x; 12330d8983SJohannes Doerfert #pragma omp target parallel map(from : x) 13330d8983SJohannes Doerfert { 14330d8983SJohannes Doerfert int *buf = llvm_omp_target_dynamic_shared_alloc() + 252; 15330d8983SJohannes Doerfert #pragma omp barrier 16330d8983SJohannes Doerfert if (omp_get_thread_num() == 0) 17330d8983SJohannes Doerfert *buf = 1; 18330d8983SJohannes Doerfert #pragma omp barrier 19330d8983SJohannes Doerfert if (omp_get_thread_num() == 1) 20330d8983SJohannes Doerfert x = *buf; 21330d8983SJohannes Doerfert } 22330d8983SJohannes Doerfert 23330d8983SJohannes Doerfert // CHECK: PASS 24330d8983SJohannes Doerfert if (x == 1 && llvm_omp_target_dynamic_shared_alloc() == NULL) 25330d8983SJohannes Doerfert printf("PASS\n"); 26330d8983SJohannes Doerfert } 27