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
|
__kernel void compiler_atomic_functions_20(__global int *dst, __local int *tmp, __global int *src) {
int lid = get_local_id(0);
int i = lid % 12;
atomic_int* p = (atomic_int*)tmp;
if(lid == 0) {
for(int j=0; j<12; j=j+1) {
atomic_exchange(&p[j], 0);
}
atomic_exchange(&p[4], -1);
}
barrier(CLK_LOCAL_MEM_FENCE);
int compare = 0;
switch(i) {
case 0: atomic_inc(&tmp[i]); break;
case 1: atomic_dec(&tmp[i]); break;
case 2: atomic_fetch_add(&p[i], src[lid]); break;
case 3: atomic_fetch_sub(&p[i], src[lid]); break;
case 4: atomic_fetch_and(&p[i], ~(src[lid]<<(lid / 16))); break;
case 5: atomic_fetch_or (&p[i], src[lid]<<(lid / 16)); break;
case 6: atomic_fetch_xor(&p[i], src[lid]); break;
case 7: atomic_fetch_min(&p[i], -src[lid]); break;
case 8: atomic_fetch_max(&p[i], src[lid]); break;
case 9: atomic_fetch_min((atomic_uint*)&p[i], -src[lid]); break;
case 10: atomic_fetch_max((atomic_uint*)&p[i], src[lid]); break;
case 11: atomic_compare_exchange_strong(&p[i], &compare, src[10]); break;
default: break;
}
atomic_int* d = (atomic_int*)dst;
switch(i) {
case 0: atomic_inc(&dst[i]); break;
case 1: atomic_dec(&dst[i]); break;
case 2: atomic_fetch_add(&d[i], src[lid]); break;
case 3: atomic_fetch_sub(&d[i], src[lid]); break;
case 4: atomic_fetch_and(&d[i], ~(src[lid]<<(lid / 16))); break;
case 5: atomic_fetch_or (&d[i], src[lid]<<(lid / 16)); break;
case 6: atomic_fetch_xor(&d[i], src[lid]); break;
case 7: atomic_fetch_min(&d[i], -src[lid]); break;
case 8: atomic_fetch_max(&d[i], src[lid]); break;
case 9: atomic_fetch_min((atomic_uint*)&d[i], -src[lid]); break;
case 10: atomic_fetch_max((atomic_uint*)&d[i], src[lid]); break;
case 11: atomic_compare_exchange_strong(&d[i], &compare, src[10]); break;
default: break;
}
barrier(CLK_GLOBAL_MEM_FENCE);
if(get_global_id(0) == 0) {
for(i=0; i<12; i=i+1)
atomic_xchg(&dst[i+12], tmp[i]);
}
}
|
