blob: fbe030fe56d00b8105c72ee0ac70d95d83b3a1fc (
plain)
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
|
struct config{
int s0;
global short *s1;
};
global int i = 5;
global int bb = 4;
global int *global p;
/* array */
global int ba[12];
/* short/long data type */
global short s;
global short s2;
global long l;
/* pointer in constant AS to global */
global int * constant px =&i;
/* constant pointer relocation */
constant int x = 2;
constant int y =1;
constant int *constant z[2] = {&x, &y};
/* structure with pointer field */
global struct config c[2] = {{1, &s}, {2, &s2} };
global int a = 1;
global int b = 2;
global int * constant gArr[2]= {&a, &b};
global int a_var[1] = {0};
global int *p_var = a_var;
__kernel void compiler_program_global0(const global int *src, int dynamic) {
size_t gid = get_global_id(0);
/* global read/write */
p = &i;
*p += 1;
/* pointer in struct memory access */
*c[gid&1].s1 += 2;
s = 2;
l = 3;
/* constant AS pointer (points to global) memory access */
*px += *z[dynamic];
p = &bb;
/* array */
if (gid < 11)
ba[gid] = src[gid];
}
__kernel void compiler_program_global1(global int *dst, int dynamic) {
size_t gid = get_global_id(0);
// static global sg;
dst[11] = i;
dst[12] = *p;
dst[13] = s;
dst[14] = l;
if (p_var == a_var)
dst[15] = *gArr[dynamic];
if (gid < 11)
dst[gid] = ba[gid];
}
__kernel void nouse(int dynamic) {
c[0].s1 = &s2;
p_var = a+dynamic;
}
|
