change stride to 4 instead of 1 if int8 is not supported (#1397)

* change stride to 2 instead of 1 if int8 is not supported

* add unit test

* add EOF line

* add second unit test from the issue

* use func

* fix comment

lib/SPIRVProducerPass.cpp

@@ -555,6 +555,13 @@ struct SPIRVProducerPassImpl {
     }
   }
 
+  uint32_t getCorrectedStride(uint32_t CurrStride) {
+    if (CurrStride == 1 && !Int8Support()) {
+      return 4;
+    }
+    return CurrStride;
+  }
+
   //
   // Primary interface for adding SPIRVInstructions to a SPIRVSection.
   template <enum SPIRVSection TSection = kFunctions>
@@ -2101,8 +2108,10 @@ SPIRVID SPIRVProducerPassImpl::getSPIRVType(Type *Ty, bool needs_layout) {
         // Ops[2] = Stride Number(Literal Number)
         Ops.clear();
 
+        auto CurrStride = static_cast<uint32_t>(GetTypeAllocSize(EleTy, DL));
+        // if stride is 1 automatically set it to 4 to avoid stride issues.
         Ops << RID << spv::DecorationArrayStride
-            << static_cast<uint32_t>(GetTypeAllocSize(EleTy, DL));
+            << getCorrectedStride(CurrStride);
 
         addSPIRVInst<kAnnotations>(spv::OpDecorate, Ops);
       }
@@ -6242,8 +6251,8 @@ void SPIRVProducerPassImpl::HandleDeferredDecorations() {
     // Ops[1] = Decoration (ArrayStride)
     // Ops[2] = Stride number (Literal Number)
     SPIRVOperandVec Ops;
-
-    Ops << id << spv::DecorationArrayStride << stride;
+    // if stride is 1 automatically set it to 4 to avoid stride issues.
+    Ops << id << spv::DecorationArrayStride << getCorrectedStride(stride);
 
     addSPIRVInst<kAnnotations>(spv::OpDecorate, Ops);
   }

test/correct_stride_int8_not_supported.ll

@@ -0,0 +1,17 @@
+// RUN: clspv -int8=0 %target %s -o %t.spv
+// RUN: spirv-dis -o %t2.spvasm %t.spv
+// RUN: FileCheck %s < %t2.spvasm
+// RUN: spirv-val --target-env vulkan1.0 %t.spv
+
+constant uchar b[4] = {[0]=42, [1]=13, [2]=0, [3]=5};
+
+void kernel __attribute__((reqd_work_group_size(4, 1, 1))) foo(global uchar* a)
+{
+  *a = b[get_local_id(0)];
+}
+
+// CHECK-DAG:  [[uint:%[^ ]+]] = OpTypeInt 32 0
+// CHECK-DAG:  OpTypePointer StorageBuffer [[struct:%[^ ]+]]
+// CHECK-DAG:  [[struct]] = OpTypeStruct [[runtimearr:%[^ ]+]]
+// CHECK-DAG:  OpDecorate [[runtimearr]] ArrayStride 4
+

test/correct_stride_int8_not_supported_char_arr.ll

@@ -0,0 +1,26 @@
+// RUN: clspv -int8=0 %target %s -o %t.spv
+// RUN: spirv-dis -o %t2.spvasm %t.spv
+// RUN: FileCheck %s < %t2.spvasm
+// RUN: spirv-val --target-env vulkan1.0 %t.spv
+
+__kernel void helloWorld(__global char* data){
+    data[0] = 'H';
+    data[1] = 'e';
+    data[2] = 'l';
+    data[3] = 'l';
+    data[4] = 'o';
+    data[5] = ' ';
+    data[6] = 'W';
+    data[7] = 'o';
+    data[8] = 'r';
+    data[9] = 'l';
+    data[10] = 'd';
+    data[11] = '!';
+    data[12] = '\n';
+    data[13] = 0;
+}
+
+// CHECK-DAG:  [[uint:%[^ ]+]] = OpTypeInt 32 0
+// CHECK-DAG:  OpTypePointer StorageBuffer [[struct:%[^ ]+]]
+// CHECK-DAG:  [[struct]] = OpTypeStruct [[runtimearr:%[^ ]+]]
+// CHECK-DAG:  OpDecorate [[runtimearr]] ArrayStride 4