SPIR-V Alignment and ArrayStride

Signed-off-by: Natalia Gavrilenko <natalia.gavrilenko@huawei.com>

Author: natgavrilenko

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/alignment1.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);
+    }
+}

benchmarks/opencl/alignment/alignment10.cl

@@ -0,0 +1,57 @@
+// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -fno-discard-value-names -cl-opt-disable -emit-llvm -c alignment2-struct-local.cl -o a.bc
+// llvm-spirv a.bc -o a.spv
+// spirv-dis a.spv > alignment2-struct-local.spvasm
+
+// clspv -O0 alignment2-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 > alignment2-struct-local.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;
+
+__kernel void test(global int *r_aligned, global int* r_unaligned) {
+    local aligned_t aligned[3];
+    local unaligned_t unaligned[3];
+
+    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);
+    }
+}