Skip to content

Commit

Permalink
Set atomic_compare_exchange argument address space to generic
Browse files Browse the repository at this point in the history 
Ensure that the `expected` argument of atomic_compare_exchange_* lives
in the generic address space in accordance with the OpenCL 2.0
specification Section 6.13.11, "The atomic_compare_exchange
functions".
  • Loading branch information
@svenvh @AlexeySotkin
svenvh authored and AlexeySotkin committed May 24, 2019
1 parent 54b49fa commit b1ecd4a
Show file tree
Hide file tree
Showing 4 changed files with 18 additions and 11 deletions.
6 changes: 5 additions & 1 deletion lib/SPIRV/SPIRVToOCL20.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -382,7 +382,11 @@ void SPIRVToOCL20::visitCallSPIRVAtomicBuiltin(CallInst *CI, Op OC) {
.getFirstInsertionPt()));
PExpected->setAlignment(CI->getType()->getScalarSizeInBits() / 8);
new StoreInst(Args[1], PExpected, PInsertBefore);
Args[1] = PExpected;
unsigned AddrSpc = SPIRAS_Generic;
Type *PtrTyAS =
PExpected->getType()->getElementType()->getPointerTo(AddrSpc);
Args[1] = CastInst::CreatePointerBitCastOrAddrSpaceCast(
PExpected, PtrTyAS, PExpected->getName() + ".as", PInsertBefore);
std::swap(Args[3], Args[4]);
std::swap(Args[2], Args[3]);
RetTy = Type::getInt1Ty(*Ctx);
Expand Down
8 changes: 4 additions & 4 deletions test/transcoding/AtomicCompareExchangeExplicit_cl20.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -45,7 +45,7 @@ __kernel void testAtomicCompareExchangeExplicit_cl20(
//CHECK-SPIRV: AtomicCompareExchangeWeak {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchangeWeak {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]

//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32* %expected1, i32 %desired, i32 3, i32 0, i32 2)
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32* %expected2, i32 %desired, i32 4, i32 0, i32 1)
//CHECK-LLVM: call spir_func i1 @_Z37atomic_compare_exchange_weak_explicitPU3AS4VU7_AtomiciPii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32* %expected3, i32 %desired, i32 3, i32 0, i32 2)
//CHECK-LLVM: call spir_func i1 @_Z37atomic_compare_exchange_weak_explicitPU3AS4VU7_AtomiciPii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32* %expected4, 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)* %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)* %expected2.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)* %expected3.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)* %expected4.as, i32 %desired, i32 4, i32 0, i32 1)
3 changes: 2 additions & 1 deletion test/transcoding/AtomicCompareExchange_cl12.ll
View file
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,8 @@ target triple = "spir-unknown-unknown"
; CHECK-LABEL: entry
; CHECK: [[PTR:%expected[0-9]*]] = alloca i32, align 4
; CHECK: store i32 {{.*}}, i32* [[PTR]]
; CHECK: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicit{{.*}}%object, i32* [[PTR]], i32 %desired, i32 5, i32 5, i32 2)
; CHECK: [[PTR]].as = addrspacecast i32* [[PTR]] to i32 addrspace(4)*
; CHECK: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicit{{.*}}%object, i32 addrspace(4)* [[PTR]].as, i32 %desired, i32 5, i32 5, i32 2)
; CHECK-NEXT; load i32* [[PTR]]
define spir_func i32 @test(i32 addrspace(1)* %object, i32 %expected, i32 %desired) #0 {
entry:
Expand Down
12 changes: 7 additions & 5 deletions test/transcoding/AtomicCompareExchange_cl20.ll
View file
Original file line number Diff line number Diff line change
Expand Up @@ -15,15 +15,17 @@ target triple = "spir-unknown-unknown"
; CHECK-NEXT: entry:
; CHECK: [[PTR_STRONG:%expected[0-9]*]] = alloca i32, align 4
; CHECK: store i32 {{.*}}, i32* [[PTR_STRONG]]
; CHECK: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicit{{.*}}(i32 {{.*}}* %object, i32* [[PTR_STRONG]], i32 %desired, i32 5, i32 5, i32 2)
; CHECK: load i32, i32* [[PTR_STRONG]]
; CHECK: [[PTR_STRONG]].as = addrspacecast i32* [[PTR_STRONG]] to i32 addrspace(4)*
; CHECK: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicit{{.*}}(i32 {{.*}}* %object, i32 {{.*}}* [[PTR_STRONG]].as, i32 %desired, i32 5, i32 5, i32 2)
; CHECK: load i32, i32 addrspace(4)* [[PTR_STRONG]].as

; CHECK-LABEL: define spir_func void @test_weak
; CHECK-NEXT: entry:
; CHECK: [[PTR_WEAK:%expected[0-9]*]] = alloca i32, align 4
; CHECK: store i32 {{.*}}, i32* [[PTR_WEAK]]
; CHECK: call spir_func i1 @_Z37atomic_compare_exchange_weak_explicitPU3AS4VU7_AtomiciPii12memory_orderS4_12memory_scope{{.*}}(i32 {{.*}}* %object, i32* [[PTR_WEAK]], i32 %desired, i32 5, i32 5, i32 2)
; CHECK: load 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: load i32, i32 addrspace(4)* [[PTR_WEAK]].as

; Check that alloca for atomic_compare_exchange is being created in the entry block.

Expand All @@ -32,7 +34,7 @@ target triple = "spir-unknown-unknown"
; CHECK: %expected{{[0-9]*}} = alloca i32
; CHECK-LABEL: for.body:
; CHECK-NOT: %expected{{[0-9]*}} = alloca i32
; CHECK: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicit{{.*}}(i32 {{.*}}* {{.*}}, i32* {{.*}}, i32 {{.*}}, i32 5, i32 5, i32 2)
; CHECK: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicit{{.*}}(i32 {{.*}}* {{.*}}, i32 addrspace(4)* {{.*}}, i32 {{.*}}, i32 5, i32 5, i32 2)

; Function Attrs: nounwind
define spir_func void @test_strong(i32 addrspace(4)* %object, i32 addrspace(4)* %expected, i32 %desired) #0 {
Expand Down

0 comments on commit b1ecd4a

Please sign in to comment.