Skip to content

Commit e554401

Browse files
wenju-hewenju-he
authored
[OCLToSPIRV] Translate floating-point atomic_compare_exchange as integer (#2668)
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.
1 parent 62ea823 commit e554401

File tree

3 files changed

+186
-40
lines changed

3 files changed

+186
-40
lines changed

‎lib/SPIRV/OCLToSPIRV.cpp

+12
Original file line numberDiff line numberDiff line change
@@ -485,6 +485,18 @@ CallInst *OCLToSPIRVBase::visitCallAtomicCmpXchg(CallInst *CI) {
485485
auto Mutator = mutateCallInst(CI, kOCLBuiltinName::AtomicCmpXchgStrong);
486486
Value *Expected = Mutator.getArg(1);
487487
Type *MemTy = Mutator.getArg(2)->getType();
488+
if (MemTy->isFloatTy() || MemTy->isDoubleTy()) {
489+
MemTy =
490+
MemTy->isFloatTy() ? Type::getInt32Ty(*Ctx) : Type::getInt64Ty(*Ctx);
491+
Mutator.replaceArg(
492+
0,
493+
{Mutator.getArg(0),
494+
TypedPointerType::get(
495+
MemTy, Mutator.getArg(0)->getType()->getPointerAddressSpace())});
496+
Mutator.mapArg(2, [=](IRBuilder<> &Builder, Value *V) {
497+
return Builder.CreateBitCast(V, MemTy);
498+
});
499+
}
488500
assert(MemTy->isIntegerTy() &&
489501
"In SPIR-V 1.0 arguments of OpAtomicCompareExchange must be "
490502
"an integer type scalars");

‎test/transcoding/AtomicCompareExchangeExplicit_cl20.cl

+106-40
Original file line numberDiff line numberDiff line change
@@ -6,46 +6,112 @@
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: AtomicCompareExchange {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
46-
//CHECK-SPIRV: AtomicCompareExchange {{[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(ptr addrspace(4) %0, ptr addrspace(4) %expected5.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(ptr addrspace(4) %0, ptr addrspace(4) %expected8.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(ptr addrspace(4) %0, ptr addrspace(4) %expected11.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(ptr addrspace(4) %0, ptr addrspace(4) %expected14.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: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
52+
//CHECK-SPIRV: AtomicCompareExchange [[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: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
57+
//CHECK-SPIRV: AtomicCompareExchange [[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: AtomicCompareExchange [[int64]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
62+
//CHECK-SPIRV: AtomicCompareExchange [[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(ptr addrspace(4) %0, ptr 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(ptr addrspace(4) %0, ptr 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(ptr addrspace(4) %0, ptr 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(ptr addrspace(4) %0, ptr 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 ptr addrspace(1) %object to ptr addrspace(4)
72+
//CHECK-LLVM: [[EXPECTED:%[0-9]+]] = addrspacecast ptr addrspace(1) %expected to ptr addrspace(4)
73+
//CHECK-LLVM: [[CAST1:%[0-9]+]] = bitcast float %desired to i32
74+
//CHECK-LLVM: %exp = load i32, ptr addrspace(4) [[EXPECTED]], align 4
75+
//CHECK-LLVM: store i32 %exp, ptr [[EXPECTED_ALLOCA:%expected[0-9]+]], align 4
76+
//CHECK-LLVM: [[EXPECTED_AS1:%expected.*]] = addrspacecast ptr [[EXPECTED_ALLOCA]] to ptr addrspace(4)
77+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(ptr addrspace(4) [[OBJECT]], ptr addrspace(4) [[EXPECTED_AS1]], i32 [[CAST1]], i32 3, i32 0, i32 2)
78+
//CHECK-LLVM: [[CAST2:%[0-9]+]] = bitcast float %desired to i32
79+
//CHECK-LLVM: [[LOAD2:%exp[0-9]+]] = load i32, ptr addrspace(4) [[EXPECTED]], align 4
80+
//CHECK-LLVM: store i32 [[LOAD2]], ptr [[EXPECTED_ALLOCA2:%expected[0-9]+]], align 4
81+
//CHECK-LLVM: [[EXPECTED_AS2:%expected.*]] = addrspacecast ptr [[EXPECTED_ALLOCA2]] to ptr addrspace(4)
82+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(ptr addrspace(4) [[OBJECT]], ptr addrspace(4) [[EXPECTED_AS2]], i32 [[CAST2]], i32 4, i32 0, i32 1)
83+
//CHECK-LLVM: [[CAST3:%[0-9]+]] = bitcast float %desired to i32
84+
//CHECK-LLVM: [[LOAD3:%exp[0-9]+]] = load i32, ptr addrspace(4) [[EXPECTED]], align 4
85+
//CHECK-LLVM: store i32 [[LOAD3]], ptr [[EXPECTED_ALLOCA3:%expected[0-9]+]], align 4
86+
//CHECK-LLVM: [[EXPECTED_AS3:%expected.*]] = addrspacecast ptr [[EXPECTED_ALLOCA3]] to ptr addrspace(4)
87+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(ptr addrspace(4) [[OBJECT]], ptr addrspace(4) [[EXPECTED_AS3]], i32 [[CAST3]], i32 3, i32 0, i32 2)
88+
//CHECK-LLVM: [[CAST4:%[0-9]+]] = bitcast float %desired to i32
89+
//CHECK-LLVM: [[LOAD4:%exp[0-9]+]] = load i32, ptr addrspace(4) [[EXPECTED]], align 4
90+
//CHECK-LLVM: store i32 [[LOAD4]], ptr [[EXPECTED_ALLOCA4:%expected[0-9]+]], align 4
91+
//CHECK-LLVM: [[EXPECTED_AS4:%expected.*]] = addrspacecast ptr [[EXPECTED_ALLOCA4]] to ptr addrspace(4)
92+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(ptr addrspace(4) [[OBJECT]], ptr addrspace(4) [[EXPECTED_AS4]], i32 [[CAST4]], i32 4, i32 0, i32 1)
93+
94+
95+
//CHECK-LLVM-LABEL: define spir_kernel void @testAtomicCompareExchangeExplicit_cl20_double(
96+
//CHECK-LLVM: [[OBJECT:%[0-9]+]] = addrspacecast ptr addrspace(1) %object to ptr addrspace(4)
97+
//CHECK-LLVM: [[EXPECTED:%[0-9]+]] = addrspacecast ptr addrspace(1) %expected to ptr addrspace(4)
98+
//CHECK-LLVM: [[CAST1:%[0-9]+]] = bitcast double %desired to i64
99+
//CHECK-LLVM: %exp = load i64, ptr addrspace(4) [[EXPECTED]], align 8
100+
//CHECK-LLVM: store i64 %exp, ptr [[EXPECTED_ALLOCA:%expected[0-9]+]], align 8
101+
//CHECK-LLVM: [[EXPECTED_AS1:%expected.*]] = addrspacecast ptr [[EXPECTED_ALLOCA]] to ptr addrspace(4)
102+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiclPU3AS4ll12memory_orderS4_12memory_scope(ptr addrspace(4) [[OBJECT]], ptr addrspace(4) [[EXPECTED_AS1]], i64 [[CAST1]], i32 3, i32 0, i32 2)
103+
//CHECK-LLVM: [[CAST2:%[0-9]+]] = bitcast double %desired to i64
104+
//CHECK-LLVM: [[LOAD2:%exp[0-9]+]] = load i64, ptr addrspace(4) [[EXPECTED]], align 8
105+
//CHECK-LLVM: store i64 [[LOAD2]], ptr [[EXPECTED_ALLOCA2:%expected[0-9]+]], align 8
106+
//CHECK-LLVM: [[EXPECTED_AS2:%expected.*]] = addrspacecast ptr [[EXPECTED_ALLOCA2]] to ptr addrspace(4)
107+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiclPU3AS4ll12memory_orderS4_12memory_scope(ptr addrspace(4) [[OBJECT]], ptr addrspace(4) [[EXPECTED_AS2]], i64 [[CAST2]], i32 4, i32 0, i32 1)
108+
//CHECK-LLVM: [[CAST3:%[0-9]+]] = bitcast double %desired to i64
109+
//CHECK-LLVM: [[LOAD3:%exp[0-9]+]] = load i64, ptr addrspace(4) [[EXPECTED]], align 8
110+
//CHECK-LLVM: store i64 [[LOAD3]], ptr [[EXPECTED_ALLOCA3:%expected[0-9]+]], align 8
111+
//CHECK-LLVM: [[EXPECTED_AS3:%expected.*]] = addrspacecast ptr [[EXPECTED_ALLOCA3]] to ptr addrspace(4)
112+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiclPU3AS4ll12memory_orderS4_12memory_scope(ptr addrspace(4) [[OBJECT]], ptr addrspace(4) [[EXPECTED_AS3]], i64 [[CAST3]], i32 3, i32 0, i32 2)
113+
//CHECK-LLVM: [[CAST4:%[0-9]+]] = bitcast double %desired to i64
114+
//CHECK-LLVM: [[LOAD4:%exp[0-9]+]] = load i64, ptr addrspace(4) [[EXPECTED]], align 8
115+
//CHECK-LLVM: store i64 [[LOAD4]], ptr [[EXPECTED_ALLOCA4:%expected[0-9]+]], align 8
116+
//CHECK-LLVM: [[EXPECTED_AS4:%expected.*]] = addrspacecast ptr [[EXPECTED_ALLOCA4]] to ptr addrspace(4)
117+
//CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiclPU3AS4ll12memory_orderS4_12memory_scope(ptr addrspace(4) [[OBJECT]], ptr addrspace(4) [[EXPECTED_AS4]], i64 [[CAST4]], i32 4, i32 0, i32 1)

0 commit comments

Comments
 (0)