Add Support for OpenCL Alignment Decoration

benchmarks/lkmm/C-PaulEMcKenney-MP+o-r+ai-mb-o.c

@@ -10,7 +10,7 @@ int r0,r1;
 
 void *thread_1(void *arg)
 {
-	WRITE_ONCE(x, 1);
+	atomic_set(&x, 1);
 	r0 = atomic_xchg_release(&y, 5);
 	return NULL;
 }
@@ -19,7 +19,7 @@ void *thread_2(void *arg)
 {
 	atomic_inc(&y);
 	smp_mb();
-	r1 = READ_ONCE(x);
+	r1 = atomic_read(&x);
 	return NULL;
 }
 

benchmarks/lkmm/C-WWC+o-branch-o+o-branch-o.c

@@ -14,27 +14,27 @@ int r4_1;
 
 void *thread_1(void *unused)
 {
-	r1_0 = READ_ONCE(x);
+	r1_0 = atomic_read(&x);
 	r3_0 = (r1_0 != 0);
 	if (r3_0) {
-		WRITE_ONCE(y, 1);
+		atomic_set(&y, 1);
 	}
     return NULL;
 }
 
 void *thread_2(void *unused)
 {
-	r2_1 = READ_ONCE(y);
+	r2_1 = atomic_read(&y);
 	r4_1 = (r2_1 != 0);
 	if (r4_1) {
-		WRITE_ONCE(x, 1);
+		atomic_set(&x, 1);
 	}
     return NULL;
 }
 
 void *thread_3(void *unused)
 {
-	WRITE_ONCE(x, 2);
+	atomic_set(&x, 2);
     return NULL;
 }
 
@@ -50,7 +50,7 @@ int main()
 	pthread_join(t2, NULL);
 	pthread_join(t3, NULL);
 
-	assert(!(r1_0 == 2 && r2_1 == 1 && READ_ONCE(x) == 2));
+	assert(!(r1_0 == 2 && r2_1 == 1 && atomic_read(&x) == 2));
 
 	return 0;
 }

benchmarks/lkmm/C-atomic-op-return-simple-02-2.c

@@ -36,7 +36,7 @@ int main()
 	pthread_join(t1, NULL);
 	pthread_join(t2, NULL);
 
-	assert(!(r0_0 == 1 && r1_0 == 0 && r0_1 == 1 && r1_1 == 0 && READ_ONCE(x) == 1 && READ_ONCE(y) == 1));
+	assert(!(r0_0 == 1 && r1_0 == 0 && r0_1 == 1 && r1_1 == 0 && atomic_read(&x) == 1 && atomic_read(&y) == 1));
 
 	return 0;
 }

benchmarks/opencl/alignment/alignment1-array-global.cl

@@ -0,0 +1,32 @@
+// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -fno-discard-value-names -cl-opt-disable -emit-llvm -c alignment1-array-global.cl -o a.bc
+// llvm-spirv a.bc -o a.spv
+// spirv-dis a.spv > alignment1-array-global.spvasm
+
+typedef int aligned_t __attribute__((vector_size(3 * sizeof(int))));
+typedef int unaligned_t[3];
+
+global static aligned_t aligned[3];
+global static unaligned_t unaligned[3];
+
+__kernel void test(global int *r_aligned, global int* r_unaligned) {
+    aligned[0][0] = 0;
+    aligned[0][1] = 1;
+    aligned[0][2] = 2;
+
+    aligned[1][0] = 3;
+    aligned[1][1] = 4;
+    aligned[1][2] = 5;
+
+    unaligned[0][0] = 6;
+    unaligned[0][1] = 7;
+    unaligned[0][2] = 8;
+
+    unaligned[1][0] = 9;
+    unaligned[1][1] = 10;
+    unaligned[1][2] = 11;
+
+    for (int i = 0; i < 8; i++) {
+        r_aligned[i] = *(((int*) aligned) + i);
+        r_unaligned[i] = *(((int*) unaligned) + i);
+    }
+}

benchmarks/opencl/alignment/alignment1-array-local.cl

@@ -0,0 +1,36 @@
+// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -fno-discard-value-names -cl-opt-disable -emit-llvm -c alignment1-array-local.cl -o a.bc
+// llvm-spirv a.bc -o a.spv
+// spirv-dis a.spv > alignment1-array-local.spvasm
+
+// clspv -O0 alignment1-array-local.cl --cl-std=CL2.0 --inline-entry-points --spv-version=1.6
+// spirv-opt --upgrade-memory-model -o a.opt.spv a.spv
+// spirv-dis a.opt.spv > alignment1-array-local.spvasm
+
+typedef int aligned_t __attribute__((vector_size(3 * sizeof(int))));
+typedef int unaligned_t[3];
+
+__kernel void test(global int *r_aligned, global int* r_unaligned) {
+    local aligned_t aligned[3];
+    local unaligned_t unaligned[3];
+
+    aligned[0][0] = 0;
+    aligned[0][1] = 1;
+    aligned[0][2] = 2;
+
+    aligned[1][0] = 3;
+    aligned[1][1] = 4;
+    aligned[1][2] = 5;
+
+    unaligned[0][0] = 6;
+    unaligned[0][1] = 7;
+    unaligned[0][2] = 8;
+
+    unaligned[1][0] = 9;
+    unaligned[1][1] = 10;
+    unaligned[1][2] = 11;
+
+    for (int i = 0; i < 8; i++) {
+        r_aligned[i] = *(((int*) aligned) + i);
+        r_unaligned[i] = *(((int*) unaligned) + i);
+    }
+}

benchmarks/opencl/alignment/alignment1-array-pointer.cl

@@ -0,0 +1,33 @@
+// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -fno-discard-value-names -cl-opt-disable -emit-llvm -c alignment1-array-pointer.cl -o a.bc
+// llvm-spirv a.bc -o a.spv
+// spirv-dis a.spv > alignment1-array-pointer.spvasm
+
+// clspv -O0 alignment1-array-pointer.cl --cl-std=CL2.0 --inline-entry-points --spv-version=1.6
+// spirv-opt --upgrade-memory-model -o a.opt.spv a.spv
+// spirv-dis a.opt.spv > alignment1-array-pointer.spvasm
+
+typedef int aligned_t __attribute__((vector_size(3 * sizeof(int))));
+typedef int unaligned_t[3] __attribute__((aligned (4)));
+
+__kernel void test(global aligned_t* aligned, global unaligned_t* unaligned, global int *r_aligned, global int* r_unaligned) {
+    aligned[0][0] = 0;
+    aligned[0][1] = 1;
+    aligned[0][2] = 2;
+
+    aligned[1][0] = 3;
+    aligned[1][1] = 4;
+    aligned[1][2] = 5;
+
+    unaligned[0][0] = 6;
+    unaligned[0][1] = 7;
+    unaligned[0][2] = 8;
+
+    unaligned[1][0] = 9;
+    unaligned[1][1] = 10;
+    unaligned[1][2] = 11;
+
+    for (int i = 0; i < 8; i++) {
+        r_aligned[i] = *(((int*) aligned) + i);
+        r_unaligned[i] = *(((int*) unaligned) + i);
+    }
+}

benchmarks/opencl/alignment/alignment1-struct-global.cl

@@ -0,0 +1,41 @@
+// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -fno-discard-value-names -cl-opt-disable -emit-llvm -c alignment1-struct-global.cl -o a.bc
+// llvm-spirv a.bc -o a.spv
+// spirv-dis a.spv > alignment1-struct-global.spvasm
+
+typedef struct __attribute__ ((aligned (16))) {
+    int x;
+    int y;
+    int z;
+} aligned_t;
+
+typedef struct {
+    int x;
+    int y;
+    int z;
+} unaligned_t;
+
+global static aligned_t aligned[3];
+global static unaligned_t unaligned[3];
+
+__kernel void test(global int *r_aligned, global int* r_unaligned) {
+    aligned[0].x = 0;
+    aligned[0].y = 1;
+    aligned[0].z = 2;
+
+    aligned[1].x = 3;
+    aligned[1].y = 4;
+    aligned[1].z = 5;
+
+    unaligned[0].x = 6;
+    unaligned[0].y = 7;
+    unaligned[0].z = 8;
+
+    unaligned[1].x = 9;
+    unaligned[1].y = 10;
+    unaligned[1].z = 11;
+
+    for (int i = 0; i < 8; i++) {
+        r_aligned[i] = *(((int*) aligned) + i);
+        r_unaligned[i] = *(((int*) unaligned) + i);
+    }
+}

benchmarks/opencl/alignment/alignment1-struct-local.cl

@@ -0,0 +1,45 @@
+// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -fno-discard-value-names -cl-opt-disable -emit-llvm -c alignment1-struct-local.cl -o a.bc
+// llvm-spirv a.bc -o a.spv
+// spirv-dis a.spv > alignment1-struct-local.spvasm
+
+// clspv -O0 alignment1-struct-local.cl --cl-std=CL2.0 --inline-entry-points --spv-version=1.6
+// spirv-opt --upgrade-memory-model -o a.opt.spv a.spv
+// spirv-dis a.opt.spv > alignment1-struct-local.spvasm
+
+typedef struct __attribute__ ((aligned (16))) {
+    int x;
+    int y;
+    int z;
+} aligned_t;
+
+typedef struct {
+    int x;
+    int y;
+    int z;
+} unaligned_t;
+
+__kernel void test(global int *r_aligned, global int* r_unaligned) {
+    local aligned_t aligned[3];
+    local unaligned_t unaligned[3];
+
+    aligned[0].x = 0;
+    aligned[0].y = 1;
+    aligned[0].z = 2;
+
+    aligned[1].x = 3;
+    aligned[1].y = 4;
+    aligned[1].z = 5;
+
+    unaligned[0].x = 6;
+    unaligned[0].y = 7;
+    unaligned[0].z = 8;
+
+    unaligned[1].x = 9;
+    unaligned[1].y = 10;
+    unaligned[1].z = 11;
+
+    for (int i = 0; i < 8; i++) {
+        r_aligned[i] = *(((int*) aligned) + i);
+        r_unaligned[i] = *(((int*) unaligned) + i);
+    }
+}

benchmarks/opencl/alignment/alignment1-struct-pointer.cl

@@ -0,0 +1,42 @@
+// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -fno-discard-value-names -cl-opt-disable -emit-llvm -c alignment1-struct-pointer.cl -o a.bc
+// llvm-spirv a.bc -o a.spv
+// spirv-dis a.spv > alignment1-struct-pointer.spvasm
+
+// clspv -O0 alignment1-struct-pointer.cl --cl-std=CL2.0 --inline-entry-points --spv-version=1.6
+// spirv-opt --upgrade-memory-model -o a.opt.spv a.spv
+// spirv-dis a.opt.spv > alignment1-struct-pointer.spvasm
+
+typedef struct __attribute__ ((aligned (16))) {
+    int x;
+    int y;
+    int z;
+} aligned_t;
+
+typedef struct {
+    int x;
+    int y;
+    int z;
+} unaligned_t;
+
+__kernel void test(global aligned_t* aligned, global unaligned_t* unaligned, global int *r_aligned, global int* r_unaligned) {
+    aligned[0].x = 0;
+    aligned[0].y = 1;
+    aligned[0].z = 2;
+
+    aligned[1].x = 3;
+    aligned[1].y = 4;
+    aligned[1].z = 5;
+
+    unaligned[0].x = 6;
+    unaligned[0].y = 7;
+    unaligned[0].z = 8;
+
+    unaligned[1].x = 9;
+    unaligned[1].y = 10;
+    unaligned[1].z = 11;
+
+    for (int i = 0; i < 8; i++) {
+        r_aligned[i] = *(((int*) aligned) + i);
+        r_unaligned[i] = *(((int*) unaligned) + i);
+    }
+}

benchmarks/opencl/alignment/alignment2-struct-global.cl

@@ -0,0 +1,53 @@
+// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -fno-discard-value-names -cl-opt-disable -emit-llvm -c alignment2-struct-global.cl -o a.bc
+// llvm-spirv a.bc -o a.spv
+// spirv-dis a.spv > alignment2-struct-global.spvasm
+
+typedef struct __attribute__ ((aligned (32))) {
+    int a;
+    int b;
+    int c;
+    int d;
+    int e;
+} aligned_t;
+
+typedef struct {
+    int a;
+    int b;
+    int c;
+    int d;
+    int e;
+} unaligned_t;
+
+global static aligned_t aligned[3];
+global static unaligned_t unaligned[3];
+
+__kernel void test(global int *r_aligned, global int* r_unaligned) {
+    aligned[0].a = 0;
+    aligned[0].b = 1;
+    aligned[0].c = 2;
+    aligned[0].d = 3;
+    aligned[0].e = 4;
+
+    aligned[1].a = 5;
+    aligned[1].b = 6;
+    aligned[1].c = 7;
+    aligned[1].d = 8;
+    aligned[1].e = 9;
+
+    unaligned[0].a = 10;
+    unaligned[0].b = 11;
+    unaligned[0].c = 12;
+    unaligned[0].d = 13;
+    unaligned[0].e = 14;
+
+    unaligned[1].a = 15;
+    unaligned[1].b = 16;
+    unaligned[1].c = 17;
+    unaligned[1].d = 18;
+    unaligned[1].e = 19;
+
+    for (int i = 0; i < 16; i++) {
+        r_aligned[i] = *(((int*) aligned) + i);
+        r_unaligned[i] = *(((int*) unaligned) + i);
+    }
+}