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 @@ -2101,8 +2108,10 @@ SPIRVID SPIRVProducerPassImpl::getSPIRVType(Type *Ty, bool needs_layout) { // Ops[2] = Stride Number(Literal Number) Ops.clear(); + auto CurrStride = static_cast(GetTypeAllocSize(EleTy, DL)); + // if stride is 1 automatically set it to 4 to avoid stride issues. Ops << RID << spv::DecorationArrayStride - << static_cast(GetTypeAllocSize(EleTy, DL)); + << getCorrectedStride(CurrStride); addSPIRVInst(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(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