Skip to content

Commit cb79377

Browse files
tonghainingtonghaining
committed
Add Supports for OpenCL Alignment Decoration
Signed-off-by: Haining Tong <[email protected]>
1 parent 97e583c commit cb79377

33 files changed

+1996
-49
lines changed

‎benchmarks/opencl/alignment/alignment1.cl

+21
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable -emit-llvm -c alignment1.cl -o a.bc
2+
// llvm-spirv a.bc -o a.spv
3+
// spirv-dis a.spv > alignment1.spv.dis
4+
5+
__kernel void test(global uint *x)
6+
{
7+
local uint3 data[2];
8+
9+
data[0].x = 0;
10+
data[0].y = 1;
11+
data[0].z = 2;
12+
13+
data[1].x = 3;
14+
data[1].y = 4;
15+
data[1].z = 5;
16+
17+
for (int i = 0; i < 8; i++)
18+
{
19+
x[i] = *(((uint*)data) + i);
20+
}
21+
}

‎benchmarks/opencl/alignment/alignment2.cl

+24
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable -emit-llvm -c alignment2.cl -o a.bc
2+
// llvm-spirv a.bc -o a.spv
3+
// spirv-dis a.spv > alignment2.spv.dis
4+
5+
struct aligned_struct {
6+
unsigned int e0 __attribute__ ((aligned (16)));
7+
unsigned int e1 __attribute__ ((aligned (16)));
8+
};
9+
10+
kernel void test(global int *x) {
11+
12+
local struct aligned_struct aligned;
13+
local unsigned int unaligned[2];
14+
15+
*(&aligned.e0) = 1;
16+
*(&aligned.e0 + 4) = 2;
17+
*(unaligned) = 3;
18+
*(unaligned + 1) = 4;
19+
20+
x[0] = aligned.e0;
21+
x[1] = aligned.e1;
22+
x[2] = unaligned[0];
23+
x[3] = unaligned[1];
24+
}

‎benchmarks/opencl/alignment/alignment3.cl

+42
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,42 @@
1+
// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable -emit-llvm -c alignment3.cl -o a.bc
2+
// llvm-spirv a.bc -o a.spv
3+
// spirv-dis a.spv > alignment3.spv.dis
4+
5+
struct aligned_struct {
6+
unsigned int e0[3] __attribute__ ((aligned (16)));
7+
unsigned int e1[3] __attribute__ ((aligned (16)));
8+
};
9+
10+
kernel void test(global int *x) {
11+
12+
local struct aligned_struct aligned;
13+
local unsigned int unaligned[2][3];
14+
15+
aligned.e0[0] = 0;
16+
aligned.e0[1] = 1;
17+
aligned.e0[2] = 2;
18+
aligned.e1[0] = 4;
19+
aligned.e1[1] = 5;
20+
aligned.e1[2] = 6;
21+
22+
unaligned[0][0] = 10;
23+
unaligned[0][1] = 11;
24+
unaligned[0][2] = 12;
25+
unaligned[1][0] = 13;
26+
unaligned[1][1] = 14;
27+
unaligned[1][2] = 15;
28+
29+
x[0] = aligned.e0[0];
30+
x[1] = aligned.e0[1];
31+
x[2] = aligned.e0[2];
32+
x[3] = aligned.e1[0];
33+
x[4] = aligned.e1[1];
34+
x[5] = aligned.e1[2];
35+
36+
x[6] = unaligned[0][0];
37+
x[7] = unaligned[0][1];
38+
x[8] = unaligned[0][2];
39+
x[9] = unaligned[1][0];
40+
x[10] = unaligned[1][1];
41+
x[11] = unaligned[1][2];
42+
}

‎benchmarks/opencl/alignment/alignment4.cl

+116
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,116 @@
1+
// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable -emit-llvm -c alignment4.cl -o a.bc
2+
// llvm-spirv a.bc -o a.spv
3+
// spirv-dis a.spv > alignment4.spv.dis
4+
5+
struct aligned_struct {
6+
unsigned int e0[3] __attribute__ ((aligned (16)));
7+
unsigned int e1[3] __attribute__ ((aligned (16)));
8+
};
9+
10+
struct nested_aligned_struct {
11+
struct aligned_struct as0[3] __attribute__ ((aligned (64)));
12+
struct aligned_struct as1[3] __attribute__ ((aligned (64)));
13+
};
14+
15+
kernel void test(global int *x) {
16+
local struct nested_aligned_struct nested_aligned;
17+
local unsigned int unaligned[2][2][3];
18+
19+
nested_aligned.as0[0].e0[0] = 0;
20+
nested_aligned.as0[0].e0[1] = 1;
21+
nested_aligned.as0[0].e0[2] = 2;
22+
nested_aligned.as0[0].e1[0] = 3;
23+
nested_aligned.as0[0].e1[1] = 4;
24+
nested_aligned.as0[0].e1[2] = 5;
25+
nested_aligned.as0[1].e0[0] = 6;
26+
nested_aligned.as0[1].e0[1] = 7;
27+
nested_aligned.as0[1].e0[2] = 8;
28+
nested_aligned.as0[1].e1[0] = 9;
29+
nested_aligned.as0[1].e1[1] = 10;
30+
nested_aligned.as0[1].e1[2] = 11;
31+
nested_aligned.as0[2].e0[0] = 12;
32+
nested_aligned.as0[2].e0[1] = 13;
33+
nested_aligned.as0[2].e0[2] = 14;
34+
nested_aligned.as0[2].e1[0] = 15;
35+
nested_aligned.as0[2].e1[1] = 16;
36+
nested_aligned.as0[2].e1[2] = 17;
37+
nested_aligned.as1[0].e0[0] = 18;
38+
nested_aligned.as1[0].e0[1] = 19;
39+
nested_aligned.as1[0].e0[2] = 20;
40+
nested_aligned.as1[0].e1[0] = 21;
41+
nested_aligned.as1[0].e1[1] = 22;
42+
nested_aligned.as1[0].e1[2] = 23;
43+
nested_aligned.as1[1].e0[0] = 24;
44+
nested_aligned.as1[1].e0[1] = 25;
45+
nested_aligned.as1[1].e0[2] = 26;
46+
nested_aligned.as1[1].e1[0] = 27;
47+
nested_aligned.as1[1].e1[1] = 28;
48+
nested_aligned.as1[1].e1[2] = 29;
49+
nested_aligned.as1[2].e0[0] = 30;
50+
nested_aligned.as1[2].e0[1] = 31;
51+
nested_aligned.as1[2].e0[2] = 32;
52+
nested_aligned.as1[2].e1[0] = 33;
53+
nested_aligned.as1[2].e1[1] = 34;
54+
nested_aligned.as1[2].e1[2] = 35;
55+
unaligned[0][0][0] = 36;
56+
unaligned[0][0][1] = 37;
57+
unaligned[0][0][2] = 38;
58+
unaligned[0][1][0] = 39;
59+
unaligned[0][1][1] = 40;
60+
unaligned[0][1][2] = 41;
61+
unaligned[1][0][0] = 42;
62+
unaligned[1][0][1] = 43;
63+
unaligned[1][0][2] = 44;
64+
unaligned[1][1][0] = 45;
65+
unaligned[1][1][1] = 46;
66+
unaligned[1][1][2] = 47;
67+
68+
x[0] = nested_aligned.as0[0].e0[0];
69+
x[1] = nested_aligned.as0[0].e0[1];
70+
x[2] = nested_aligned.as0[0].e0[2];
71+
x[3] = nested_aligned.as0[0].e1[0];
72+
x[4] = nested_aligned.as0[0].e1[1];
73+
x[5] = nested_aligned.as0[0].e1[2];
74+
x[6] = nested_aligned.as0[1].e0[0];
75+
x[7] = nested_aligned.as0[1].e0[1];
76+
x[8] = nested_aligned.as0[1].e0[2];
77+
x[9] = nested_aligned.as0[1].e1[0];
78+
x[10] = nested_aligned.as0[1].e1[1];
79+
x[11] = nested_aligned.as0[1].e1[2];
80+
x[12] = nested_aligned.as0[2].e0[0];
81+
x[13] = nested_aligned.as0[2].e0[1];
82+
x[14] = nested_aligned.as0[2].e0[2];
83+
x[15] = nested_aligned.as0[2].e1[0];
84+
x[16] = nested_aligned.as0[2].e1[1];
85+
x[17] = nested_aligned.as0[2].e1[2];
86+
x[18] = nested_aligned.as1[0].e0[0];
87+
x[19] = nested_aligned.as1[0].e0[1];
88+
x[20] = nested_aligned.as1[0].e0[2];
89+
x[21] = nested_aligned.as1[0].e1[0];
90+
x[22] = nested_aligned.as1[0].e1[1];
91+
x[23] = nested_aligned.as1[0].e1[2];
92+
x[24] = nested_aligned.as1[1].e0[0];
93+
x[25] = nested_aligned.as1[1].e0[1];
94+
x[26] = nested_aligned.as1[1].e0[2];
95+
x[27] = nested_aligned.as1[1].e1[0];
96+
x[28] = nested_aligned.as1[1].e1[1];
97+
x[29] = nested_aligned.as1[1].e1[2];
98+
x[30] = nested_aligned.as1[2].e0[0];
99+
x[31] = nested_aligned.as1[2].e0[1];
100+
x[32] = nested_aligned.as1[2].e0[2];
101+
x[33] = nested_aligned.as1[2].e1[0];
102+
x[34] = nested_aligned.as1[2].e1[1];
103+
x[35] = nested_aligned.as1[2].e1[2];
104+
x[36] = unaligned[0][0][0];
105+
x[37] = unaligned[0][0][1];
106+
x[38] = unaligned[0][0][2];
107+
x[39] = unaligned[0][1][0];
108+
x[40] = unaligned[0][1][1];
109+
x[41] = unaligned[0][1][2];
110+
x[42] = unaligned[1][0][0];
111+
x[43] = unaligned[1][0][1];
112+
x[44] = unaligned[1][0][2];
113+
x[45] = unaligned[1][1][0];
114+
x[46] = unaligned[1][1][1];
115+
x[47] = unaligned[1][1][2];
116+
}

‎benchmarks/opencl/alignment/alignment5.cl

+32
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable -emit-llvm -c alignment5.cl -o a.bc
2+
// llvm-spirv a.bc -o a.spv
3+
// spirv-dis a.spv > alignment5.spv.dis
4+
5+
struct struct1 {
6+
int a; // 4 bytes, aligned to 4
7+
char b; // 1 byte
8+
}__attribute__ ((aligned (16)));
9+
10+
struct struct2 {
11+
char b; // 1 byte
12+
int a; // 4 bytes, aligned to 4
13+
};
14+
15+
16+
__kernel void manual_vs_struct(__global uchar *out)
17+
{
18+
local struct struct1 s1;
19+
local struct struct2 s2;
20+
21+
s1.a = 0;
22+
s1.b = 1;
23+
s2.a = 2;
24+
s2.b = 3;
25+
*((int *)(&s1.a + 1)) = 22;
26+
*((int *)(&s2.a + 1)) = 33;
27+
28+
out[0] = s1.a;
29+
out[1] = s1.b;
30+
out[2] = s2.a;
31+
out[3] = s2.b;
32+
}

‎benchmarks/opencl/alignment/alignment6.cl

+43
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,43 @@
1+
// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable -emit-llvm -c alignment6.cl -o a.bc
2+
// llvm-spirv a.bc -o a.spv
3+
// spirv-dis a.spv > alignment6.spv.dis
4+
5+
struct aligned_struct {
6+
unsigned int e0[3] __attribute__ ((aligned (16)));
7+
unsigned int e1[3] __attribute__ ((aligned (16)));
8+
};
9+
10+
global struct aligned_struct aligned;
11+
12+
kernel void test(global int *x, global struct aligned_struct *aligned) {
13+
14+
local unsigned int unaligned[2][3];
15+
16+
aligned->e0[0] = 0;
17+
aligned->e0[1] = 1;
18+
aligned->e0[2] = 2;
19+
aligned->e1[0] = 4;
20+
aligned->e1[1] = 5;
21+
aligned->e1[2] = 6;
22+
23+
unaligned[0][0] = 10;
24+
unaligned[0][1] = 11;
25+
unaligned[0][2] = 12;
26+
unaligned[1][0] = 13;
27+
unaligned[1][1] = 14;
28+
unaligned[1][2] = 15;
29+
30+
x[0] = aligned->e0[0];
31+
x[1] = aligned->e0[1];
32+
x[2] = aligned->e0[2];
33+
x[3] = aligned->e1[0];
34+
x[4] = aligned->e1[1];
35+
x[5] = aligned->e1[2];
36+
37+
x[6] = unaligned[0][0];
38+
x[7] = unaligned[0][1];
39+
x[8] = unaligned[0][2];
40+
x[9] = unaligned[1][0];
41+
x[10] = unaligned[1][1];
42+
x[11] = unaligned[1][2];
43+
}

‎benchmarks/opencl/alignment/alignment7.cl

+31
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable -emit-llvm -c alignment7.cl -o a.bc
2+
// llvm-spirv a.bc -o a.spv
3+
// spirv-dis a.spv > alignment7.spv.dis
4+
5+
struct struct1 {
6+
int a; // 4 bytes, aligned to 4
7+
char b; // 1 byte
8+
}__attribute__ ((aligned (16)));
9+
10+
struct struct2 {
11+
char b; // 1 byte
12+
int a; // 4 bytes, aligned to 4
13+
};
14+
15+
global struct struct1 s1;
16+
global struct struct2 s2;
17+
18+
__kernel void manual_vs_struct(__global uchar *out, __global struct struct1 *s1, __global struct struct2 *s2)
19+
{
20+
s1->a = 0;
21+
s1->b = 1;
22+
s2->a = 2;
23+
s2->b = 3;
24+
*((int *)(&s1->a + 1)) = 22;
25+
*((int *)(&s2->a + 1)) = 33;
26+
27+
out[0] = s1->a;
28+
out[1] = s1->b;
29+
out[2] = s2->a;
30+
out[3] = s2->b;
31+
}

‎benchmarks/opencl/alignment/alignment8.cl

+25
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable -emit-llvm -c alignment8.cl -o a.bc
2+
// llvm-spirv a.bc -o a.spv
3+
// spirv-dis a.spv > alignment8.spv.dis
4+
5+
struct aligned_struct {
6+
unsigned int e0 __attribute__ ((aligned (16)));
7+
unsigned int e1 __attribute__ ((aligned (16)));
8+
};
9+
10+
global struct aligned_struct aligned;
11+
12+
kernel void test(global int *x, global struct aligned_struct *aligned) {
13+
14+
local unsigned int unaligned[2];
15+
16+
*(&aligned->e0) = 1;
17+
*(&aligned->e0 + 4) = 2;
18+
*(unaligned) = 3;
19+
*(unaligned + 1) = 4;
20+
21+
x[0] = aligned->e0;
22+
x[1] = aligned->e1;
23+
x[2] = unaligned[0];
24+
x[3] = unaligned[1];
25+
}

‎benchmarks/opencl/alignment/alignment9.cl

+25
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable -emit-llvm -c alignment9.cl -o a.bc
2+
// llvm-spirv a.bc -o a.spv
3+
// spirv-dis a.spv > alignment9.spv.dis
4+
5+
struct aligned_struct {
6+
uint3 data[2];
7+
};
8+
9+
global struct aligned_struct aligned;
10+
11+
__kernel void test(global uint *x, global struct aligned_struct *aligned)
12+
{
13+
aligned->data[0].x = 0;
14+
aligned->data[0].y = 1;
15+
aligned->data[0].z = 2;
16+
17+
aligned->data[1].x = 3;
18+
aligned->data[1].y = 4;
19+
aligned->data[1].z = 5;
20+
21+
for (int i = 0; i < 8; i++)
22+
{
23+
x[i] = *(((uint*)aligned) + i);
24+
}
25+
}

‎dartagnan/src/main/java/com/dat3m/dartagnan/expression/ExpressionFactory.java

+5
Original file line numberDiff line numberDiff line change
@@ -27,13 +27,18 @@ public final class ExpressionFactory {
2727
private final BooleanType booleanType = types.getBooleanType();
2828
private final BoolLiteral falseConstant = new BoolLiteral(booleanType, false);
2929
private final BoolLiteral trueConstant = new BoolLiteral(booleanType, true);
30+
private final Expression defaultAlignment = makeValue(8, types.getArchType());
3031

3132
private ExpressionFactory() {}
3233

3334
public static ExpressionFactory getInstance() {
3435
return instance;
3536
}
3637

38+
public Expression getDefaultAlignment() {
39+
return defaultAlignment;
40+
}
41+
3742
// -----------------------------------------------------------------------------------------------------------------
3843
// Boolean
3944

0 commit comments

Comments
 (0)