Skip to content

Commit d838f77

Browse files
wenju-hewenju-he
authored
[Backport to 14][OCLToSPIRV] Translate floating-point atomic_compare_exchange as integer (#2668) (#2671)
OpenCL spec supports atomic_float/atomic_double type for atomic_compare_exchange* functions. However, value and return type in OpAtomicCompareExchange in SPIR-V spec must be integer type. Therefore, in OCLToSPIRV translation we need to translate floating-point type to corresponding integer variant that has the same type size. Floating-point value is bitcasted so that bits remain the same. (cherry picked from commit e554401)
1 parent 144fa52 commit d838f77

File tree

3 files changed

+210
-43
lines changed

3 files changed

+210
-43
lines changed

‎lib/SPIRV/OCLToSPIRV.cpp

Lines changed: 17 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -472,10 +472,24 @@ CallInst *OCLToSPIRVBase::visitCallAtomicCmpXchg(CallInst *CI) {
472472
mutateCallInstOCL(
473473
M, CI,
474474
[&](CallInst *CI, std::vector<Value *> &Args, Type *&RetTy) {
475-
Expected = Args[1]; // temporary save second argument.
476-
Args[1] = new LoadInst(Args[1]->getType()->getPointerElementType(),
477-
Args[1], "exp", false, CI);
478475
RetTy = Args[2]->getType();
476+
if (RetTy->isFloatTy() || RetTy->isDoubleTy()) {
477+
RetTy = RetTy->isFloatTy() ? Type::getInt32Ty(*Ctx)
478+
: Type::getInt64Ty(*Ctx);
479+
Args[0] = new BitCastInst(
480+
Args[0],
481+
PointerType::get(RetTy,
482+
Args[0]->getType()->getPointerAddressSpace()),
483+
"", CI);
484+
Args[1] = new BitCastInst(
485+
Args[1],
486+
PointerType::get(RetTy,
487+
Args[1]->getType()->getPointerAddressSpace()),
488+
"", CI);
489+
Args[2] = new BitCastInst(Args[2], RetTy, "", CI);
490+
}
491+
Expected = Args[1]; // temporary save second argument.
492+
Args[1] = new LoadInst(RetTy, Args[1], "exp", false, CI);
479493
assert(Args[0]->getType()->getPointerElementType()->isIntegerTy() &&
480494
Args[1]->getType()->isIntegerTy() &&
481495
Args[2]->getType()->isIntegerTy() &&

‎test/transcoding/AtomicCompareExchangeExplicit_cl20.cl

Lines changed: 121 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -6,46 +6,127 @@
66
// RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.rev.bc
77
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
88

9-
__kernel void testAtomicCompareExchangeExplicit_cl20(
10-
volatile global atomic_int* object,
11-
global int* expected,
12-
int desired)
13-
{
14-
// Values of memory order and memory scope arguments correspond to SPIR-2.0 spec.
15-
atomic_compare_exchange_strong_explicit(object, expected, desired,
16-
memory_order_release, // 3
17-
memory_order_relaxed // 0
18-
); // by default, assume device scope = 2
19-
atomic_compare_exchange_strong_explicit(object, expected, desired,
20-
memory_order_acq_rel, // 4
21-
memory_order_relaxed, // 0
22-
memory_scope_work_group // 1
23-
);
24-
atomic_compare_exchange_weak_explicit(object, expected, desired,
25-
memory_order_release, // 3
26-
memory_order_relaxed // 0
27-
); // by default, assume device scope = 2
28-
atomic_compare_exchange_weak_explicit(object, expected, desired,
29-
memory_order_acq_rel, // 4
30-
memory_order_relaxed, // 0
31-
memory_scope_work_group // 1
32-
);
9+
#define DEFINE_KERNEL(TYPE) \
10+
__kernel void testAtomicCompareExchangeExplicit_cl20_##TYPE( \
11+
volatile global atomic_##TYPE* object, \
12+
global TYPE* expected, \
13+
TYPE desired) \
14+
{ \
15+
/* Values of memory order and memory scope arguments correspond to SPIR-2.0 spec. */ \
16+
atomic_compare_exchange_strong_explicit(object, expected, desired, \
17+
memory_order_release, /* 3 */ \
18+
memory_order_relaxed /* 0 */ \
19+
); /* by default, assume device scope = 2 */ \
20+
atomic_compare_exchange_strong_explicit(object, expected, desired, \
21+
memory_order_acq_rel, /* 4 */ \
22+
memory_order_relaxed, /* 0 */ \
23+
memory_scope_work_group /* 1 */ \
24+
); \
25+
atomic_compare_exchange_weak_explicit(object, expected, desired, \
26+
memory_order_release, /* 3 */ \
27+
memory_order_relaxed /* 0 */ \
28+
); /* by default, assume device scope = 2 */ \
29+
atomic_compare_exchange_weak_explicit(object, expected, desired, \
30+
memory_order_acq_rel, /* 4 */ \
31+
memory_order_relaxed, /* 0 */ \
32+
memory_scope_work_group /* 1 */ \
33+
); \
3334
}
3435

35-
//CHECK-SPIRV: TypeInt [[int:[0-9]+]] 32 0
36+
DEFINE_KERNEL(int)
37+
DEFINE_KERNEL(float)
38+
DEFINE_KERNEL(double)
39+
40+
//CHECK-SPIRV: TypeInt [[int32:[0-9]+]] 32 0
41+
//CHECK-SPIRV: TypeInt [[int64:[0-9]+]] 64 0
3642
//; Constants below correspond to the SPIR-V spec
37-
//CHECK-SPIRV-DAG: Constant [[int]] [[DeviceScope:[0-9]+]] 1
38-
//CHECK-SPIRV-DAG: Constant [[int]] [[WorkgroupScope:[0-9]+]] 2
39-
//CHECK-SPIRV-DAG: Constant [[int]] [[ReleaseMemSem:[0-9]+]] 4
40-
//CHECK-SPIRV-DAG: Constant [[int]] [[RelaxedMemSem:[0-9]+]] 0
41-
//CHECK-SPIRV-DAG: Constant [[int]] [[AcqRelMemSem:[0-9]+]] 8
42-
43-
//CHECK-SPIRV: AtomicCompareExchange {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
44-
//CHECK-SPIRV: AtomicCompareExchange {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
45-
//CHECK-SPIRV: AtomicCompareExchangeWeak {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
46-
//CHECK-SPIRV: AtomicCompareExchangeWeak {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
47-
48-
//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)
49-
//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)
50-
//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)
51-
//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)
43+
//CHECK-SPIRV-DAG: Constant [[int32]] [[DeviceScope:[0-9]+]] 1
44+
//CHECK-SPIRV-DAG: Constant [[int32]] [[WorkgroupScope:[0-9]+]] 2
45+
//CHECK-SPIRV-DAG: Constant [[int32]] [[ReleaseMemSem:[0-9]+]] 4
46+
//CHECK-SPIRV-DAG: Constant [[int32]] [[RelaxedMemSem:[0-9]+]] 0
47+
//CHECK-SPIRV-DAG: Constant [[int32]] [[AcqRelMemSem:[0-9]+]] 8
48+
49+
//CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
50+
//CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
51+
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
52+
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int32]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
53+
54+
//CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
55+
//CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
56+
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
57+
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int32]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
58+
59+
//CHECK-SPIRV: AtomicCompareExchange [[int64]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
60+
//CHECK-SPIRV: AtomicCompareExchange [[int64]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
61+
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int64]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
62+
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int64]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
63+
64+
//CHECK-LLVM-LABEL: define spir_kernel void @testAtomicCompareExchangeExplicit_cl20_int(
65+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected{{.*}}, i32 %desired, i32 3, i32 0, i32 2)
66+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected{{.*}}, i32 %desired, i32 4, i32 0, i32 1)
67+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected{{.*}}, i32 %desired, i32 3, i32 0, i32 2)
68+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* %0, i32 addrspace(4)* %expected{{.*}}, i32 %desired, i32 4, i32 0, i32 1)
69+
70+
//CHECK-LLVM-LABEL: define spir_kernel void @testAtomicCompareExchangeExplicit_cl20_float(
71+
//CHECK-LLVM: [[OBJECT:%[0-9]+]] = addrspacecast float addrspace(1)* %object to float addrspace(4)*
72+
//CHECK-LLVM: [[EXPECTED:%[0-9]+]] = addrspacecast float addrspace(1)* %expected to float addrspace(4)*
73+
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast float addrspace(4)* [[OBJECT]] to i32 addrspace(4)*
74+
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast float addrspace(4)* [[EXPECTED]] to i32 addrspace(4)*
75+
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast float %desired to i32
76+
//CHECK-LLVM: %exp = load i32, i32 addrspace(4)* [[EXPECTED_BC]], align 4
77+
//CHECK-LLVM: store i32 %exp, i32* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 4
78+
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i32* [[EXPECTED_ALLOCA]] to i32 addrspace(4)*
79+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* [[OBJECT_BC]], i32 addrspace(4)* [[EXPECTED_AS]], i32 [[DESIRED_CAST]], i32 3, i32 0, i32 2)
80+
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast float addrspace(4)* [[OBJECT]] to i32 addrspace(4)*
81+
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast float addrspace(4)* [[EXPECTED]] to i32 addrspace(4)*
82+
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast float %desired to i32
83+
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i32, i32 addrspace(4)* [[EXPECTED_BC]], align 4
84+
//CHECK-LLVM: store i32 [[LOAD]], i32* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 4
85+
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i32* [[EXPECTED_ALLOCA]] to i32 addrspace(4)*
86+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* [[OBJECT_BC]], i32 addrspace(4)* [[EXPECTED_AS]], i32 [[DESIRED_CAST]], i32 4, i32 0, i32 1)
87+
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast float addrspace(4)* [[OBJECT]] to i32 addrspace(4)*
88+
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast float addrspace(4)* [[EXPECTED]] to i32 addrspace(4)*
89+
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast float %desired to i32
90+
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i32, i32 addrspace(4)* [[EXPECTED_BC]], align 4
91+
//CHECK-LLVM: store i32 [[LOAD]], i32* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 4
92+
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i32* [[EXPECTED_ALLOCA]] to i32 addrspace(4)*
93+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* [[OBJECT_BC]], i32 addrspace(4)* [[EXPECTED_AS]], i32 [[DESIRED_CAST]], i32 3, i32 0, i32 2)
94+
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast float addrspace(4)* [[OBJECT]] to i32 addrspace(4)*
95+
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast float addrspace(4)* [[EXPECTED]] to i32 addrspace(4)*
96+
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast float %desired to i32
97+
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i32, i32 addrspace(4)* [[EXPECTED_BC]], align 4
98+
//CHECK-LLVM: store i32 [[LOAD]], i32* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 4
99+
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i32* [[EXPECTED_ALLOCA]] to i32 addrspace(4)*
100+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(i32 addrspace(4)* [[OBJECT_BC]], i32 addrspace(4)* [[EXPECTED_AS]], i32 [[DESIRED_CAST]], i32 4, i32 0, i32 1)
101+
102+
//CHECK-LLVM-LABEL: define spir_kernel void @testAtomicCompareExchangeExplicit_cl20_double(
103+
//CHECK-LLVM: [[OBJECT:%[0-9]+]] = addrspacecast double addrspace(1)* %object to double addrspace(4)*
104+
//CHECK-LLVM: [[EXPECTED:%[0-9]+]] = addrspacecast double addrspace(1)* %expected to double addrspace(4)*
105+
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast double addrspace(4)* [[OBJECT]] to i64 addrspace(4)*
106+
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast double addrspace(4)* [[EXPECTED]] to i64 addrspace(4)*
107+
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast double %desired to i64
108+
//CHECK-LLVM: %exp = load i64, i64 addrspace(4)* [[EXPECTED_BC]], align 8
109+
//CHECK-LLVM: store i64 %exp, i64* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 8
110+
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i64* [[EXPECTED_ALLOCA]] to i64 addrspace(4)*
111+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiclPU3AS4ll12memory_orderS4_12memory_scope(i64 addrspace(4)* [[OBJECT_BC]], i64 addrspace(4)* [[EXPECTED_AS]], i64 [[DESIRED_CAST]], i32 3, i32 0, i32 2)
112+
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast double addrspace(4)* [[OBJECT]] to i64 addrspace(4)*
113+
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast double addrspace(4)* [[EXPECTED]] to i64 addrspace(4)*
114+
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast double %desired to i64
115+
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i64, i64 addrspace(4)* [[EXPECTED_BC]], align 8
116+
//CHECK-LLVM: store i64 [[LOAD]], i64* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 8
117+
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i64* [[EXPECTED_ALLOCA]] to i64 addrspace(4)*
118+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiclPU3AS4ll12memory_orderS4_12memory_scope(i64 addrspace(4)* [[OBJECT_BC]], i64 addrspace(4)* [[EXPECTED_AS]], i64 [[DESIRED_CAST]], i32 4, i32 0, i32 1)
119+
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast double addrspace(4)* [[OBJECT]] to i64 addrspace(4)*
120+
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast double addrspace(4)* [[EXPECTED]] to i64 addrspace(4)*
121+
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast double %desired to i64
122+
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i64, i64 addrspace(4)* [[EXPECTED_BC]], align 8
123+
//CHECK-LLVM: store i64 [[LOAD]], i64* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 8
124+
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i64* [[EXPECTED_ALLOCA]] to i64 addrspace(4)*
125+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiclPU3AS4ll12memory_orderS4_12memory_scope(i64 addrspace(4)* [[OBJECT_BC]], i64 addrspace(4)* [[EXPECTED_AS]], i64 [[DESIRED_CAST]], i32 3, i32 0, i32 2)
126+
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast double addrspace(4)* [[OBJECT]] to i64 addrspace(4)*
127+
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast double addrspace(4)* [[EXPECTED]] to i64 addrspace(4)*
128+
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast double %desired to i64
129+
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i64, i64 addrspace(4)* [[EXPECTED_BC]], align 8
130+
//CHECK-LLVM: store i64 [[LOAD]], i64* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 8
131+
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i64* [[EXPECTED_ALLOCA]] to i64 addrspace(4)*
132+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiclPU3AS4ll12memory_orderS4_12memory_scope(i64 addrspace(4)* [[OBJECT_BC]], i64 addrspace(4)* [[EXPECTED_AS]], i64 [[DESIRED_CAST]], i32 4, i32 0, i32 1)

0 commit comments

Comments
 (0)