This repository has been archived by the owner on Oct 3, 2018. It is now read-only.
/
atomics-base.cl
86 lines (75 loc) · 3.61 KB
/
atomics-base.cl
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
// RUN: %webcl-validator "%s" 2>&1 | grep -v CHECK | %FileCheck "%s"
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
// check only one return argument per argument length, per argument type and per address space
// long global
// CHECK-DAG: return atomic_add(_
// CHECK-DAG: return atomic_inc(_
// CHECK-DAG: return atomic_cmpxchg(_
// ulong global
// CHECK-DAG: return atomic_add(_
// CHECK-DAG: return atomic_inc(_
// CHECK-DAG: return atomic_cmpxchg(_
// long local
// CHECK-DAG: return atomic_add(_
// CHECK-DAG: return atomic_inc(_
// CHECK-DAG: return atomic_cmpxchg(_
// ulong local
// CHECK-DAG: return atomic_add(_
// CHECK-DAG: return atomic_inc(_
// CHECK-DAG: return atomic_cmpxchg(_
__kernel void atoms(volatile __global long* global_long_data, volatile __local long* local_long_data,
volatile __global ulong* global_ulong_data, volatile __local ulong* local_ulong_data,
volatile __global float* global_float_data, volatile __local float* local_float_data)
{
long val_long = 1;
ulong val_ulong = 2;
ulong cmp = 1;
// CHECK: (_wcl_allocs, global_long_data, val_long);
val_long = atomic_add (global_long_data, val_long);
// CHECK: (_wcl_allocs, global_long_data, val_long);
val_long = atomic_sub (global_long_data, val_long);
// CHECK: (_wcl_allocs, global_long_data, val_long);
val_long = atomic_xchg (global_long_data, val_long);
// CHECK: (_wcl_allocs, global_long_data);
val_long = atomic_inc (global_long_data);
// CHECK: (_wcl_allocs, global_long_data);
val_long = atomic_dec (global_long_data);
// CHECK: (_wcl_allocs, global_long_data, cmp, val_long);
val_long = atomic_cmpxchg(global_long_data, cmp, val_long);
// CHECK: (_wcl_allocs, global_ulong_data, val_ulong);
val_ulong = atomic_add (global_ulong_data, val_ulong);
// CHECK: (_wcl_allocs, global_ulong_data, val_ulong);
val_ulong = atomic_sub (global_ulong_data, val_ulong);
// CHECK: (_wcl_allocs, global_ulong_data, val_ulong);
val_ulong = atomic_xchg (global_ulong_data, val_ulong);
// CHECK: (_wcl_allocs, global_ulong_data);
val_ulong = atomic_inc (global_ulong_data);
// CHECK: (_wcl_allocs, global_ulong_data);
val_ulong = atomic_dec (global_ulong_data);
// CHECK: (_wcl_allocs, global_ulong_data, cmp, val_ulong);
val_ulong = atomic_cmpxchg(global_ulong_data, cmp, val_ulong);
// CHECK: (_wcl_allocs, local_long_data, val_long);
val_long = atomic_add (local_long_data, val_long);
// CHECK: (_wcl_allocs, local_long_data, val_long);
val_long = atomic_sub (local_long_data, val_long);
// CHECK: (_wcl_allocs, local_long_data, val_long);
val_long = atomic_xchg (local_long_data, val_long);
// CHECK: (_wcl_allocs, local_long_data);
val_long = atomic_inc (local_long_data);
// CHECK: (_wcl_allocs, local_long_data);
val_long = atomic_dec (local_long_data);
// CHECK: (_wcl_allocs, local_long_data, cmp, val_long);
val_long = atomic_cmpxchg(local_long_data, cmp, val_long);
// CHECK: (_wcl_allocs, local_ulong_data, val_ulong);
val_ulong = atomic_add (local_ulong_data, val_ulong);
// CHECK: (_wcl_allocs, local_ulong_data, val_ulong);
val_ulong = atomic_sub (local_ulong_data, val_ulong);
// CHECK: (_wcl_allocs, local_ulong_data, val_ulong);
val_ulong = atomic_xchg (local_ulong_data, val_ulong);
// CHECK: (_wcl_allocs, local_ulong_data);
val_ulong = atomic_inc (local_ulong_data);
// CHECK: (_wcl_allocs, local_ulong_data);
val_ulong = atomic_dec (local_ulong_data);
// CHECK: (_wcl_allocs, local_ulong_data, cmp, val_ulong);
val_ulong = atomic_cmpxchg(local_ulong_data, cmp, val_ulong);
}