From b1ecd4a3959baa869b4136db53cb32692c05ad98 Mon Sep 17 00:00:00 2001 From: Sven van Haastregt Date: Thu, 16 May 2019 15:09:32 +0100 Subject: [PATCH] Set atomic_compare_exchange argument address space to generic 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". --- lib/SPIRV/SPIRVToOCL20.cpp | 6 +++++- .../AtomicCompareExchangeExplicit_cl20.cl | 8 ++++---- test/transcoding/AtomicCompareExchange_cl12.ll | 3 ++- test/transcoding/AtomicCompareExchange_cl20.ll | 12 +++++++----- 4 files changed, 18 insertions(+), 11 deletions(-) diff --git a/lib/SPIRV/SPIRVToOCL20.cpp b/lib/SPIRV/SPIRVToOCL20.cpp index 298850ced4..de934e5ea5 100644 --- a/lib/SPIRV/SPIRVToOCL20.cpp +++ b/lib/SPIRV/SPIRVToOCL20.cpp @@ -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); diff --git a/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl b/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl index 72ef810be8..4509631cae 100644 --- a/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl +++ b/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl @@ -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) diff --git a/test/transcoding/AtomicCompareExchange_cl12.ll b/test/transcoding/AtomicCompareExchange_cl12.ll index bda7e03f21..64e5df1626 100644 --- a/test/transcoding/AtomicCompareExchange_cl12.ll +++ b/test/transcoding/AtomicCompareExchange_cl12.ll @@ -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: diff --git a/test/transcoding/AtomicCompareExchange_cl20.ll b/test/transcoding/AtomicCompareExchange_cl20.ll index b0f29d9e10..0331310fad 100644 --- a/test/transcoding/AtomicCompareExchange_cl20.ll +++ b/test/transcoding/AtomicCompareExchange_cl20.ll @@ -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. @@ -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 {