Wrong generation for int8 when int8 is not supported
- Dominant language
- LLVM
- Stars
- 729
- Forks
- 110
- Avg merge
- 17h 51m
- Merged PRs (30d)
- 23
Description
This is following https://github.com/kpet/clvk/issues/716 and https://github.com/google/clspv/issues/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
```
Contributor guide
Assessment
This issue has not been assessed yet.