Fix mistranslation of OpAtomicCompareExchangeWeak

The semantics of OpAtomicCompareExchangeWeak are the same as those of
OpAtomicCompareExchange.  Each of the ops should be translated to
atomic_compare_exchange_strong_explicit() when using OpenCL 2.0 builtins,
and not atomic_compare_exchange_weak_explicit(), as use of the latter may
result in unreported spurious failures which are not detectable when
using OpAtomicCompareExchangeWeak, which does not give a direct
indication of its success.

Author: StuartDBrady

lib/SPIRV/SPIRVToOCL.h

@@ -187,7 +187,7 @@ class SPIRVToOCLBase : public InstVisitor<SPIRVToOCLBase> {
 
   /// Transform __spirv_OpAtomicCompareExchange and
   /// __spirv_OpAtomicCompareExchangeWeak
-  virtual Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI, Op OC) = 0;
+  virtual Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI) = 0;
 
   /// Transform __spirv_OpAtomicIIncrement/OpAtomicIDecrement to:
   /// - OCL2.0: atomic_fetch_add_explicit/atomic_fetch_sub_explicit
@@ -293,10 +293,10 @@ class SPIRVToOCL12Base : public SPIRVToOCLBase {
   /// (bool)atomic_xchg(*ptr, 1)
   Instruction *visitCallSPIRVAtomicFlagTestAndSet(CallInst *CI);
 
-  /// Transform __spirv_OpAtomicCompareExchange and
-  /// __spirv_OpAtomicCompareExchangeWeak into atomic_cmpxchg. There is no
-  /// weak version of function in OpenCL 1.2
-  Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI, Op OC) override;
+  /// Transform __spirv_OpAtomicCompareExchange/Weak into atomic_cmpxchg
+  /// OpAtomicCompareExchangeWeak is not "weak" at all, but instead has
+  /// the same semantics as OpAtomicCompareExchange.
+  Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI) override;
 
   /// Conduct generic mutations for all atomic builtins
   CallInst *mutateCommonAtomicArguments(CallInst *CI, Op OC) override;
@@ -371,8 +371,10 @@ class SPIRVToOCL20Base : public SPIRVToOCLBase {
   std::string mapFPAtomicName(Op OC) override;
 
   /// Transform __spirv_OpAtomicCompareExchange/Weak into
-  /// compare_exchange_strong/weak_explicit
-  Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI, Op OC) override;
+  /// atomic_compare_exchange_strong_explicit
+  /// OpAtomicCompareExchangeWeak is not "weak" at all, but instead has
+  /// the same semantics as OpAtomicCompareExchange.
+  Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI) override;
 };
 
 class SPIRVToOCL20Pass : public llvm::PassInfoMixin<SPIRVToOCL20Pass>,

lib/SPIRV/SPIRVToOCL12.cpp

@@ -205,8 +205,7 @@ SPIRVToOCL12Base::visitCallSPIRVAtomicFlagTestAndSet(CallInst *CI) {
       &Attrs);
 }
 
-Instruction *SPIRVToOCL12Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI,
-                                                            Op OC) {
+Instruction *SPIRVToOCL12Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI) {
   AttributeList Attrs = CI->getCalledFunction()->getAttributes();
   return mutateCallInstOCL(
       M, CI,
@@ -247,7 +246,7 @@ Instruction *SPIRVToOCL12Base::visitCallSPIRVAtomicBuiltin(CallInst *CI,
     break;
   case OpAtomicCompareExchange:
   case OpAtomicCompareExchangeWeak:
-    NewCI = visitCallSPIRVAtomicCmpExchg(CI, OC);
+    NewCI = visitCallSPIRVAtomicCmpExchg(CI);
     break;
   default:
     NewCI = mutateCommonAtomicArguments(CI, OC);

lib/SPIRV/SPIRVToOCL20.cpp

@@ -159,7 +159,7 @@ Instruction *SPIRVToOCL20Base::visitCallSPIRVAtomicBuiltin(CallInst *CI,
     break;
   case OpAtomicCompareExchange:
   case OpAtomicCompareExchangeWeak:
-    NewCI = visitCallSPIRVAtomicCmpExchg(CIG, OC);
+    NewCI = visitCallSPIRVAtomicCmpExchg(CIG);
     break;
   default:
     NewCI = mutateAtomicName(CIG, OC);
@@ -232,8 +232,7 @@ CallInst *SPIRVToOCL20Base::mutateCommonAtomicArguments(CallInst *CI, Op OC) {
       &Attrs);
 }
 
-Instruction *SPIRVToOCL20Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI,
-                                                            Op OC) {
+Instruction *SPIRVToOCL20Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI) {
   assert(CI->getCalledFunction() && "Unexpected indirect call");
   AttributeList Attrs = CI->getCalledFunction()->getAttributes();
   Instruction *PInsertBefore = CI;
@@ -242,7 +241,7 @@ Instruction *SPIRVToOCL20Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI,
       M, CI,
       [=](CallInst *, std::vector<Value *> &Args, Type *&RetTy) {
         // OpAtomicCompareExchange[Weak] semantics is different from
-        // atomic_compare_exchange_[strong|weak] semantics as well as
+        // atomic_compare_exchange_strong semantics as well as
         // arguments order.
         // OCL built-ins returns boolean value and stores a new/original
         // value by pointer passed as 2nd argument (aka expected) while SPIR-V
@@ -263,7 +262,9 @@ Instruction *SPIRVToOCL20Base::visitCallSPIRVAtomicCmpExchg(CallInst *CI,
         std::swap(Args[3], Args[4]);
         std::swap(Args[2], Args[3]);
         RetTy = Type::getInt1Ty(*Ctx);
-        return OCLSPIRVBuiltinMap::rmap(OC);
+        // OpAtomicCompareExchangeWeak is not "weak" at all, but instead has
+        // the same semantics as OpAtomicCompareExchange.
+        return "atomic_compare_exchange_strong_explicit";
       },
       [=](CallInst *CI) -> Instruction * {
         // OCL built-ins atomic_compare_exchange_[strong|weak] return boolean

test/transcoding/AtomicCompareExchangeExplicit_cl20.cl

@@ -47,5 +47,5 @@ __kernel void testAtomicCompareExchangeExplicit_cl20(
 
 //CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected1.as, i32 %desired, i32 3, i32 0, i32 2)
 //CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected3.as, i32 %desired, i32 4, i32 0, i32 1)
-//CHECK-LLVM: call spir_func i1 @_Z37atomic_compare_exchange_weak_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected5.as, i32 %desired, i32 3, i32 0, i32 2)
-//CHECK-LLVM: call spir_func i1 @_Z37atomic_compare_exchange_weak_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected7.as, i32 %desired, i32 4, i32 0, i32 1)
+//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected5.as, i32 %desired, i32 3, i32 0, i32 2)
+//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected7.as, i32 %desired, i32 4, i32 0, i32 1)

test/transcoding/AtomicCompareExchange_cl20.ll

@@ -9,6 +9,13 @@ target triple = "spir-unknown-unknown"
 
 ; Check 'LLVM ==> SPIR-V ==> LLVM' conversion of atomic_compare_exchange_strong and atomic_compare_exchange_weak.
 
+; SPIR-V does not include an equivalent of atomic_compare_exchange_weak
+; (OpAtomicCompareExchangeWeak is identical to OpAtomicCompareExchange and
+; is deprecated, and removed in SPIR-V 1.4.)
+; This breaks the round trip for atomic_compare_exchange_weak, which must be
+; translated back to LLVM IR as atomic_compare_exchange_strong, regardless
+; of whether OpAtomicCompareExchange or OpAtomicCompareExchangeWeak is used.
+
 ; Function Attrs: nounwind
 
 ; CHECK-LABEL:   define spir_func void @test_strong
@@ -24,7 +31,7 @@ target triple = "spir-unknown-unknown"
 ; CHECK:         [[PTR_WEAK:%expected[0-9]*]] = alloca i32, align 4
 ; CHECK:         store i32 {{.*}}, i32* [[PTR_WEAK]]
 ; CHECK:         [[PTR_WEAK]].as = addrspacecast i32* [[PTR_WEAK]] to i32 addrspace(4)*
-; CHECK:         call spir_func i1 @_Z37atomic_compare_exchange_weak_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope{{.*}}(i32 {{.*}}* %object, i32 {{.*}}* [[PTR_WEAK]].as, i32 %desired, i32 5, i32 5, i32 2)
+; CHECK:         call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope{{.*}}(i32 {{.*}}* %object, i32 {{.*}}* [[PTR_WEAK]].as, i32 %desired, i32 5, i32 5, i32 2)
 ; CHECK:         load i32, i32 addrspace(4)* [[PTR_WEAK]].as
 
 ; Check that alloca for atomic_compare_exchange is being created in the entry block.

test/transcoding/atomics.spt

@@ -1,10 +1,10 @@
-119734787 65536 393230 32 0 
+119734787 65536 393230 33 0 
 2 Capability Addresses 
 2 Capability Kernel 
 5 ExtInstImport 1 "OpenCL.std"
 3 MemoryModel 2 2 
 8 EntryPoint 6 6 "test_atomic_global"
-13 String 31 "kernel_arg_type.test_atomic_global.int*,"
+13 String 32 "kernel_arg_type.test_atomic_global.int*,"
 3 Source 3 102000 
 3 Name 7 "dst"
 4 Name 8 "object"
@@ -18,7 +18,7 @@
 2 TypeVoid 2 
 4 TypePointer 4 5 3 
 7 TypeFunction 5 2 4 4 4 3 
-2 TypeBool 29 
+2 TypeBool 30 
 
 
 5 Function 2 6 0 5 
@@ -40,10 +40,11 @@
 7 AtomicXor 3 24 7 13 14 13 
 7 AtomicAnd 3 25 7 13 14 13 
 9 AtomicCompareExchange 3 26 7 13 14 14 13 17 
-7 AtomicExchange 3 27 7 13 14 13 
-6 AtomicLoad 3 28 8 13 14 
+9 AtomicCompareExchangeWeak 3 27 7 13 14 14 13 17 
+7 AtomicExchange 3 28 7 13 14 13 
+6 AtomicLoad 3 29 8 13 14 
 5 AtomicStore 8 13 14 10 
-6 AtomicFlagTestAndSet 29 30 8 13 14 
+6 AtomicFlagTestAndSet 30 31 8 13 14 
 4 AtomicFlagClear 8 13 14 
 1 Return 
 
@@ -66,7 +67,7 @@
 ; CHECK-LLVM-12: call spir_func i32 @_Z9atomic_orPU3AS1Vii(i32 addrspace(1)* %dst, i32 1) [[attr]]
 ; CHECK-LLVM-12: call spir_func i32 @_Z10atomic_xorPU3AS1Vii(i32 addrspace(1)* %dst, i32 1) [[attr]]
 ; CHECK-LLVM-12: call spir_func i32 @_Z10atomic_andPU3AS1Vii(i32 addrspace(1)* %dst, i32 1) [[attr]]
-; CHECK-LLVM-12: call spir_func i32 @_Z14atomic_cmpxchgPU3AS1Viii(i32 addrspace(1)* %dst, i32 0, i32 1) [[attr]]
+; CHECK-LLVM-12-COUNT-2: call spir_func i32 @_Z14atomic_cmpxchgPU3AS1Viii(i32 addrspace(1)* %dst, i32 0, i32 1) [[attr]]
 ; CHECK-LLVM-12: call spir_func i32 @_Z11atomic_xchgPU3AS1Vii(i32 addrspace(1)* %dst, i32 1) [[attr]]
 ; CHECK-LLVM-12: call spir_func i32 @_Z10atomic_addPU3AS1Vii(i32 addrspace(1)* %object, i32 0) [[attr]]
 ; CHECK-LLVM-12: call spir_func i32 @_Z11atomic_xchgPU3AS1Vii(i32 addrspace(1)* %object, i32 %desired) [[attr]]
@@ -89,11 +90,12 @@
 ; CHECK-LLVM-20: call spir_func i32 @_Z25atomic_fetch_xor_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %dst.as9, i32 1, i32 5, i32 2) [[attr]]
 ; CHECK-LLVM-20: call spir_func i32 @_Z25atomic_fetch_and_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %dst.as10, i32 1, i32 5, i32 2) [[attr]]
 ; CHECK-LLVM-20: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %dst.as11, i32 addrspace(4)* %expected12.as, i32 1, i32 5, i32 5, i32 2) [[attr]]
-; CHECK-LLVM-20: call spir_func i32 @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %dst.as13, i32 1, i32 5, i32 2) [[attr]]
+; CHECK-LLVM-20: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %dst.as13, i32 addrspace(4)* %expected14.as, i32 1, i32 5, i32 5, i32 2) [[attr]]
+; CHECK-LLVM-20: call spir_func i32 @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %dst.as15, i32 1, i32 5, i32 2) [[attr]]
 ; CHECK-LLVM-20: call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(i32 addrspace(4)* %object.as, i32 5, i32 2) [[attr]]
-; CHECK-LLVM-20: call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %object.as14, i32 %desired, i32 5, i32 2) [[attr]]
-; CHECK-LLVM-20: call spir_func i1 @_Z33atomic_flag_test_and_set_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(i32 addrspace(4)* %object.as15, i32 5, i32 2) [[attr]]
-; CHECK-LLVM-20: call spir_func void @_Z26atomic_flag_clear_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(i32 addrspace(4)* %object.as16, i32 5, i32 2) [[attr]]
+; CHECK-LLVM-20: call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(i32 addrspace(4)* %object.as16, i32 %desired, i32 5, i32 2) [[attr]]
+; CHECK-LLVM-20: call spir_func i1 @_Z33atomic_flag_test_and_set_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(i32 addrspace(4)* %object.as17, i32 5, i32 2) [[attr]]
+; CHECK-LLVM-20: call spir_func void @_Z26atomic_flag_clear_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(i32 addrspace(4)* %object.as18, i32 5, i32 2) [[attr]]
 
 ; RUN: llvm-spirv -r %t1.spv -o %t2.bc --spirv-target-env="SPV-IR"
 ; RUN: llvm-dis < %t2.bc | FileCheck %s --check-prefix=CHECK-LLVM-SPV-IR