From 68fcec9e2c49fdfd1c9c61ed429e5c1721951ffa Mon Sep 17 00:00:00 2001 From: Syed Faaiz Hussain Date: Wed, 11 Sep 2024 21:14:29 -0700 Subject: [PATCH 1/6] change stride to 2 instead of 1 if int8 is not supported --- lib/SPIRVProducerPass.cpp | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/lib/SPIRVProducerPass.cpp b/lib/SPIRVProducerPass.cpp index 1216d68cf..365aa8651 100644 --- a/lib/SPIRVProducerPass.cpp +++ b/lib/SPIRVProducerPass.cpp @@ -2084,8 +2084,12 @@ SPIRVID SPIRVProducerPassImpl::getSPIRVType(Type *Ty, bool needs_layout) { // Ops[2] = Stride Number(Literal Number) Ops.clear(); - Ops << RID << spv::DecorationArrayStride - << static_cast(GetTypeAllocSize(EleTy, DL)); + // if stride is 1 automatically set it to 2 to avoid stride issues. + auto CurrStride = static_cast(GetTypeAllocSize(EleTy, DL)); + if (CurrStride == 1 && !Int8Support()) { + CurrStride = 2; + } + Ops << RID << spv::DecorationArrayStride << CurrStride; addSPIRVInst(spv::OpDecorate, Ops); } From ea8e346c080bebc1210ffb3c00aa7479ac36e9c0 Mon Sep 17 00:00:00 2001 From: Syed Faaiz Hussain Date: Mon, 16 Sep 2024 08:27:33 -0700 Subject: [PATCH 2/6] add unit test --- lib/SPIRVProducerPass.cpp | 5 ++++- test/correct_stride_int8_not_supported.ll | 16 ++++++++++++++++ 2 files changed, 20 insertions(+), 1 deletion(-) create mode 100644 test/correct_stride_int8_not_supported.ll diff --git a/lib/SPIRVProducerPass.cpp b/lib/SPIRVProducerPass.cpp index 365aa8651..bbbddad42 100644 --- a/lib/SPIRVProducerPass.cpp +++ b/lib/SPIRVProducerPass.cpp @@ -2087,7 +2087,7 @@ SPIRVID SPIRVProducerPassImpl::getSPIRVType(Type *Ty, bool needs_layout) { // if stride is 1 automatically set it to 2 to avoid stride issues. auto CurrStride = static_cast(GetTypeAllocSize(EleTy, DL)); if (CurrStride == 1 && !Int8Support()) { - CurrStride = 2; + CurrStride = 4; } Ops << RID << spv::DecorationArrayStride << CurrStride; @@ -6010,6 +6010,9 @@ void SPIRVProducerPassImpl::HandleDeferredDecorations() { // Ops[2] = Stride number (Literal Number) SPIRVOperandVec Ops; + if (stride == 1 && !Int8Support()) { + stride = 4; + } Ops << id << spv::DecorationArrayStride << 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..a2d4aa287 --- /dev/null +++ b/test/correct_stride_int8_not_supported.ll @@ -0,0 +1,16 @@ +// 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 \ No newline at end of file From 92caeb5bae1db0c59f31f66f8c324c9348edd046 Mon Sep 17 00:00:00 2001 From: Syed Faaiz Hussain Date: Mon, 16 Sep 2024 08:29:11 -0700 Subject: [PATCH 3/6] add EOF line --- test/correct_stride_int8_not_supported.ll | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/test/correct_stride_int8_not_supported.ll b/test/correct_stride_int8_not_supported.ll index a2d4aa287..1e398dc1e 100644 --- a/test/correct_stride_int8_not_supported.ll +++ b/test/correct_stride_int8_not_supported.ll @@ -13,4 +13,5 @@ void kernel __attribute__((reqd_work_group_size(4, 1, 1))) foo(global uchar* a) // CHECK-DAG: [[uint:%[^ ]+]] = OpTypeInt 32 0 // CHECK-DAG: OpTypePointer StorageBuffer [[struct:%[^ ]+]] // CHECK-DAG: [[struct]] = OpTypeStruct [[runtimearr:%[^ ]+]] -// CHECK-DAG: OpDecorate [[runtimearr]] ArrayStride 4 \ No newline at end of file +// CHECK-DAG: OpDecorate [[runtimearr]] ArrayStride 4 + From e2a9bcaa1420a1c824bfa99f9d5d7412fce50d88 Mon Sep 17 00:00:00 2001 From: Syed Faaiz Hussain Date: Mon, 16 Sep 2024 08:51:32 -0700 Subject: [PATCH 4/6] add second unit test from the issue --- ...rect_stride_int8_not_supported_char_arr.ll | 26 +++++++++++++++++++ 1 file changed, 26 insertions(+) create mode 100644 test/correct_stride_int8_not_supported_char_arr.ll 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 From ce8775adee3162e5c74cf245e15f22b7613cf2c6 Mon Sep 17 00:00:00 2001 From: Syed Faaiz Hussain Date: Mon, 23 Sep 2024 09:10:10 -0700 Subject: [PATCH 5/6] use func --- lib/SPIRVProducerPass.cpp | 19 ++++++++++--------- 1 file changed, 10 insertions(+), 9 deletions(-) diff --git a/lib/SPIRVProducerPass.cpp b/lib/SPIRVProducerPass.cpp index bbbddad42..83b059ab3 100644 --- a/lib/SPIRVProducerPass.cpp +++ b/lib/SPIRVProducerPass.cpp @@ -540,6 +540,13 @@ struct SPIRVProducerPassImpl { return false; } + uint32_t getCorrectedStride(uint32_t CurrStride) { + if (CurrStride == 1 && !Int8Support()) { + return 4; + } + return CurrStride; + } + // // Primary interface for adding SPIRVInstructions to a SPIRVSection. template @@ -2086,10 +2093,8 @@ SPIRVID SPIRVProducerPassImpl::getSPIRVType(Type *Ty, bool needs_layout) { // if stride is 1 automatically set it to 2 to avoid stride issues. auto CurrStride = static_cast(GetTypeAllocSize(EleTy, DL)); - if (CurrStride == 1 && !Int8Support()) { - CurrStride = 4; - } - Ops << RID << spv::DecorationArrayStride << CurrStride; + Ops << RID << spv::DecorationArrayStride + << getCorrectedStride(CurrStride); addSPIRVInst(spv::OpDecorate, Ops); } @@ -6009,11 +6014,7 @@ void SPIRVProducerPassImpl::HandleDeferredDecorations() { // Ops[1] = Decoration (ArrayStride) // Ops[2] = Stride number (Literal Number) SPIRVOperandVec Ops; - - if (stride == 1 && !Int8Support()) { - stride = 4; - } - Ops << id << spv::DecorationArrayStride << stride; + Ops << id << spv::DecorationArrayStride << getCorrectedStride(stride); addSPIRVInst(spv::OpDecorate, Ops); } From 0fe6a7a1d26e2659bcd74712d2d6fd13eda96ceb Mon Sep 17 00:00:00 2001 From: Syed Faaiz Hussain Date: Mon, 23 Sep 2024 09:19:08 -0700 Subject: [PATCH 6/6] fix comment --- lib/SPIRVProducerPass.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/lib/SPIRVProducerPass.cpp b/lib/SPIRVProducerPass.cpp index 83b059ab3..c470ddee5 100644 --- a/lib/SPIRVProducerPass.cpp +++ b/lib/SPIRVProducerPass.cpp @@ -2091,8 +2091,8 @@ SPIRVID SPIRVProducerPassImpl::getSPIRVType(Type *Ty, bool needs_layout) { // Ops[2] = Stride Number(Literal Number) Ops.clear(); - // if stride is 1 automatically set it to 2 to avoid stride issues. auto CurrStride = static_cast(GetTypeAllocSize(EleTy, DL)); + // if stride is 1 automatically set it to 4 to avoid stride issues. Ops << RID << spv::DecorationArrayStride << getCorrectedStride(CurrStride); @@ -6014,6 +6014,7 @@ void SPIRVProducerPassImpl::HandleDeferredDecorations() { // Ops[1] = Decoration (ArrayStride) // Ops[2] = Stride number (Literal Number) SPIRVOperandVec Ops; + // if stride is 1 automatically set it to 4 to avoid stride issues. Ops << id << spv::DecorationArrayStride << getCorrectedStride(stride); addSPIRVInst(spv::OpDecorate, Ops);