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

Wrong generation for int8 when int8 is not supported #1420

Open
rjodinchr opened this issue Nov 19, 2024 · 0 comments
Open

Wrong generation for int8 when int8 is not supported #1420

rjodinchr opened this issue Nov 19, 2024 · 0 comments

Comments

@rjodinchr
Copy link
Collaborator

rjodinchr commented Nov 19, 2024

This is following kpet/clvk#716 and #1396

When int8 is not supported, clspv translate:

  %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x i8] } zeroinitializer)
  %1 = getelementptr { [0 x i8] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0
  store i8 72, ptr addrspace(1) %1, align 1
  %2 = getelementptr { [0 x i8] }, ptr addrspace(1) %0, i32 0, i32 0, i32 1
  store i8 101, ptr addrspace(1) %2, align 1

into

         %19 = OpAccessChain %_ptr_StorageBuffer_uint %12 %uint_0 %uint_0
               OpStore %19 %uint_72 Aligned 1
         %22 = OpAccessChain %_ptr_StorageBuffer_uint %12 %uint_0 %uint_1
               OpStore %22 %uint_101 Aligned 1

Which is not correct as the bytes are now 4 bytes spaced instead of 1.

This is coming from compiling:

__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;
}

With the following options:

-cl-single-precision-constant -cl-kernel-arg-info -rounding-mode-rte=16,32,64 -int8=0 -std430-ubo-layout -decorate-nonuniform -hack-convert-to-float -arch=spir --use-native-builtins=ceil,copysign,fabs,fdim,floor,fmax,fmin,half_cos,half_exp,half_exp10,half_exp2,half_rsqrt,half_sin,half_sqrt,half_tan,isequal,isfinite,isgreater,isgreaterequal,isinf,isless,islessequal,islessgreater,isnan,isnormal,isnotequal,isordered,isunordered,mad,rint,round,rsqrt,signbit,sqrt,trunc, -spv-version=1.6 -max-pushconstant-size=256 -max-ubo-size=65536 -global-offset -long-vector -module-constants-in-storage-buffer -cl-arm-non-uniform-work-group-size -enable-printf -printf-buffer-size=1048576
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

No branches or pull requests

1 participant