From efc952843c0b1ae91625d78808338c16378d57f3 Mon Sep 17 00:00:00 2001
From: Syed Faaiz <syedfaaiz@google.com>
Date: Tue, 24 Sep 2024 00:45:03 -0700
Subject: [PATCH] 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                     | 15 ++++++++---
 test/correct_stride_int8_not_supported.ll     | 17 ++++++++++++
 ...rect_stride_int8_not_supported_char_arr.ll | 26 +++++++++++++++++++
 3 files changed, 55 insertions(+), 3 deletions(-)
 create mode 100644 test/correct_stride_int8_not_supported.ll
 create mode 100644 test/correct_stride_int8_not_supported_char_arr.ll

diff --git a/lib/SPIRVProducerPass.cpp b/lib/SPIRVProducerPass.cpp
index c02155a8a..5356979a9 100644
--- a/lib/SPIRVProducerPass.cpp
+++ b/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);
   }
diff --git a/test/correct_stride_int8_not_supported.ll b/test/correct_stride_int8_not_supported.ll
new file mode 100644
index 000000000..1e398dc1e
--- /dev/null
+++ b/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
+
diff --git a/test/correct_stride_int8_not_supported_char_arr.ll b/test/correct_stride_int8_not_supported_char_arr.ll
new file mode 100644
index 000000000..481b1dc76
--- /dev/null
+++ b/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