clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp

   1 // REQUIRES: amdgpu-registered-target
   2 // RUN: %clang_cc1 %s -emit-llvm -O0 -o - \
   3 // RUN:   -triple=amdgcn-amd-amdhsa  | opt -S | FileCheck %s
   4
   5 void test_memory_fence_success() {
   6   // CHECK-LABEL: test_memory_fence_success
   7
   8   // CHECK: fence syncscope("workgroup") seq_cst
   9   __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
  10
  11   // CHECK: fence syncscope("agent") acquire
  12   __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent");
  13
  14   // CHECK: fence seq_cst
  15   __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
  16
  17   // CHECK: fence syncscope("agent") acq_rel
  18   __builtin_amdgcn_fence(4, "agent");
  19
  20   // CHECK: fence syncscope("workgroup") release
  21   __builtin_amdgcn_fence(3, "workgroup");
  22 }