Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

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

Open
wants to merge 4 commits into
base: main
Choose a base branch
from

Conversation

Rekt3421
Copy link
Collaborator

@Rekt3421 Rekt3421 commented Sep 12, 2024

please see issue #1396 for more details

@Rekt3421
Copy link
Collaborator Author

Is a unit test needed for this? Seems a little trivial @rjodinchr ?

Copy link
Collaborator

@rjodinchr rjodinchr left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is not the only place where we create spv::DecorationArrayStride.
I would like to have every place fixed.

Also, why 2? when int8 is not supported, it is converted to int32 not int16.

We would definitely need tests here. At least the one from the issue.

@rjodinchr
Copy link
Collaborator

This one feels like a good one as well:

// 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

You will note that even if you change ArrayStride 4 to ArrayStride 2 in the test, it is still not passing with 2. It would be failing with the following message:

Structure id 15 decorated as Block for variable in StorageBuffer storage class must follow standard storage buffer layout rules: member 0 contains an array with stride 2 not satisfying alignment to 4

@Rekt3421 Rekt3421 changed the title change stride to 2 instead of 1 if int8 is not supported change stride to 4 instead of 1 if int8 is not supported Sep 16, 2024
@Rekt3421 Rekt3421 marked this pull request as draft September 16, 2024 15:31
@Rekt3421
Copy link
Collaborator Author

Will add the test from the issue as well.

@Rekt3421 Rekt3421 marked this pull request as ready for review September 16, 2024 15:52
Comment on lines +2089 to +2091
if (CurrStride == 1 && !Int8Support()) {
CurrStride = 4;
}
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please avoid code duplication. Make a function for it.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants