1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
|
// RUN: %clang_cc1 -triple spir-unknown-unknown -O1 -cl-std=CL2.0 -fdeclare-opencl-builtins -finclude-default-header -emit-llvm-bc %s -o %t.bc
// RUN: llvm-spirv %t.bc -spirv-text -o %t.txt
// RUN: FileCheck < %t.txt %s --check-prefix=CHECK-SPIRV
// RUN: llvm-spirv %t.bc -o %t.spv
// RUN: spirv-val %t.spv
// RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.rev.bc
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
#define DEFINE_KERNEL(TYPE) \
__kernel void testAtomicCompareExchangeExplicit_cl20_##TYPE( \
volatile global atomic_##TYPE* object, \
global TYPE* expected, \
TYPE desired) \
{ \
/* Values of memory order and memory scope arguments correspond to SPIR-2.0 spec. */ \
atomic_compare_exchange_strong_explicit(object, expected, desired, \
memory_order_release, /* 3 */ \
memory_order_relaxed /* 0 */ \
); /* by default, assume device scope = 2 */ \
atomic_compare_exchange_strong_explicit(object, expected, desired, \
memory_order_acq_rel, /* 4 */ \
memory_order_relaxed, /* 0 */ \
memory_scope_work_group /* 1 */ \
); \
atomic_compare_exchange_weak_explicit(object, expected, desired, \
memory_order_release, /* 3 */ \
memory_order_relaxed /* 0 */ \
); /* by default, assume device scope = 2 */ \
atomic_compare_exchange_weak_explicit(object, expected, desired, \
memory_order_acq_rel, /* 4 */ \
memory_order_relaxed, /* 0 */ \
memory_scope_work_group /* 1 */ \
); \
}
DEFINE_KERNEL(int)
DEFINE_KERNEL(float)
DEFINE_KERNEL(double)
//CHECK-SPIRV: TypeInt [[int32:[0-9]+]] 32 0
//CHECK-SPIRV: TypeInt [[int64:[0-9]+]] 64 0
//; Constants below correspond to the SPIR-V spec
//CHECK-SPIRV-DAG: Constant [[int32]] [[DeviceScope:[0-9]+]] 1
//CHECK-SPIRV-DAG: Constant [[int32]] [[WorkgroupScope:[0-9]+]] 2
//CHECK-SPIRV-DAG: Constant [[int32]] [[ReleaseMemSem:[0-9]+]] 4
//CHECK-SPIRV-DAG: Constant [[int32]] [[RelaxedMemSem:[0-9]+]] 0
//CHECK-SPIRV-DAG: Constant [[int32]] [[AcqRelMemSem:[0-9]+]] 8
//CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int32]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int32]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchange [[int64]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchange [[int64]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int64]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
//CHECK-SPIRV: AtomicCompareExchangeWeak [[int64]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
//CHECK-LLVM-LABEL: define spir_kernel void @testAtomicCompareExchangeExplicit_cl20_int(
//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)
//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)
//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)
//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)
//CHECK-LLVM-LABEL: define spir_kernel void @testAtomicCompareExchangeExplicit_cl20_float(
//CHECK-LLVM: [[OBJECT:%[0-9]+]] = addrspacecast float addrspace(1)* %object to float addrspace(4)*
//CHECK-LLVM: [[EXPECTED:%[0-9]+]] = addrspacecast float addrspace(1)* %expected to float addrspace(4)*
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast float addrspace(4)* [[OBJECT]] to i32 addrspace(4)*
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast float addrspace(4)* [[EXPECTED]] to i32 addrspace(4)*
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast float %desired to i32
//CHECK-LLVM: %exp = load i32, i32 addrspace(4)* [[EXPECTED_BC]], align 4
//CHECK-LLVM: store i32 %exp, i32* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 4
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i32* [[EXPECTED_ALLOCA]] to i32 addrspace(4)*
//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)
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast float addrspace(4)* [[OBJECT]] to i32 addrspace(4)*
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast float addrspace(4)* [[EXPECTED]] to i32 addrspace(4)*
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast float %desired to i32
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i32, i32 addrspace(4)* [[EXPECTED_BC]], align 4
//CHECK-LLVM: store i32 [[LOAD]], i32* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 4
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i32* [[EXPECTED_ALLOCA]] to i32 addrspace(4)*
//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)
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast float addrspace(4)* [[OBJECT]] to i32 addrspace(4)*
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast float addrspace(4)* [[EXPECTED]] to i32 addrspace(4)*
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast float %desired to i32
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i32, i32 addrspace(4)* [[EXPECTED_BC]], align 4
//CHECK-LLVM: store i32 [[LOAD]], i32* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 4
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i32* [[EXPECTED_ALLOCA]] to i32 addrspace(4)*
//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)
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast float addrspace(4)* [[OBJECT]] to i32 addrspace(4)*
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast float addrspace(4)* [[EXPECTED]] to i32 addrspace(4)*
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast float %desired to i32
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i32, i32 addrspace(4)* [[EXPECTED_BC]], align 4
//CHECK-LLVM: store i32 [[LOAD]], i32* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 4
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i32* [[EXPECTED_ALLOCA]] to i32 addrspace(4)*
//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)
//CHECK-LLVM-LABEL: define spir_kernel void @testAtomicCompareExchangeExplicit_cl20_double(
//CHECK-LLVM: [[OBJECT:%[0-9]+]] = addrspacecast double addrspace(1)* %object to double addrspace(4)*
//CHECK-LLVM: [[EXPECTED:%[0-9]+]] = addrspacecast double addrspace(1)* %expected to double addrspace(4)*
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast double addrspace(4)* [[OBJECT]] to i64 addrspace(4)*
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast double addrspace(4)* [[EXPECTED]] to i64 addrspace(4)*
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast double %desired to i64
//CHECK-LLVM: %exp = load i64, i64 addrspace(4)* [[EXPECTED_BC]], align 8
//CHECK-LLVM: store i64 %exp, i64* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 8
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i64* [[EXPECTED_ALLOCA]] to i64 addrspace(4)*
//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)
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast double addrspace(4)* [[OBJECT]] to i64 addrspace(4)*
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast double addrspace(4)* [[EXPECTED]] to i64 addrspace(4)*
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast double %desired to i64
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i64, i64 addrspace(4)* [[EXPECTED_BC]], align 8
//CHECK-LLVM: store i64 [[LOAD]], i64* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 8
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i64* [[EXPECTED_ALLOCA]] to i64 addrspace(4)*
//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)
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast double addrspace(4)* [[OBJECT]] to i64 addrspace(4)*
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast double addrspace(4)* [[EXPECTED]] to i64 addrspace(4)*
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast double %desired to i64
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i64, i64 addrspace(4)* [[EXPECTED_BC]], align 8
//CHECK-LLVM: store i64 [[LOAD]], i64* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 8
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i64* [[EXPECTED_ALLOCA]] to i64 addrspace(4)*
//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)
//CHECK-LLVM: [[OBJECT_BC:%[0-9]+]] = bitcast double addrspace(4)* [[OBJECT]] to i64 addrspace(4)*
//CHECK-LLVM: [[EXPECTED_BC:%[0-9]+]] = bitcast double addrspace(4)* [[EXPECTED]] to i64 addrspace(4)*
//CHECK-LLVM: [[DESIRED_CAST:%[0-9]+]] = bitcast double %desired to i64
//CHECK-LLVM: [[LOAD:%exp[0-9]+]] = load i64, i64 addrspace(4)* [[EXPECTED_BC]], align 8
//CHECK-LLVM: store i64 [[LOAD]], i64* [[EXPECTED_ALLOCA:%expected[0-9]+]], align 8
//CHECK-LLVM: [[EXPECTED_AS:%expected.*]] = addrspacecast i64* [[EXPECTED_ALLOCA]] to i64 addrspace(4)*
//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)
|
