[Issue]: possible wrong vector width for fp8/int8 and fp16/bf16

[michael604work]

Problem Description

I think /aiter/csrc/kernels/quant_kernels.cu
vector widths must be updated to set the vector size to
16 when fp8/int8,
8 when fp16/bf16.

Operating System

CentOS Stream 9

CPU

AMD EPYC 9655

GPU

MI350X

ROCm Version

rocm7

ROCm Component

No response

Steps to Reproduce

No response

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

No response

Additional Information

No response

[valarLip]

any more details?

[michael604work]

Sorry for the lack of information.

I'm not sure if this is the issue with ck_tile actually.

When I use data_to_per_row_scale(), scaled_quant_impl(), or dynamic_per_token_scaled_quant() from /aiter/csrc/kernels/quant_kernels.cu,

my environment eventually uses ck_tile kernel of amd_buffer_load_impl() from /composable_kernel/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp

While amd_buffer_load_impl() does not support vector size of 16 or higher for bf16/fp16, and 32 or higher for fp8/int8, etc.,
quant_kernels.cu does not set the vector size appropriately with consideration of the datatype.

[valarLip]

aiter/aiter/jit/optCompilerConfig.json

Line 562 in decd64e

"'-DCK_TILE_USE_BUFFER_ADDRESSING_BUILTIN=0'"

this flag will help then, ck changed the default code path of amd_buffer_load_impl, so we need force -DCK_TILE_USE_BUFFER_ADDRESSING_BUILTIN=0

[michael604work]

Well that flag was already set at module_quant of optCompilerConfig.json :-(

I instead modified the aiter/csrc/kernels/quant_kernels.cu, to check for the datatypes whenever the vector width is 16 or higher, and limited the vector width when the datatype is fp8/int8 or fp16/bf16.

Do you think this is also valid?

[valarLip]

yes, it's valid and we are using it for a few month already...

[michael604work]

Oh I see :-)
Thanks!

BTW, what would you say the root cause of this is?
Should I be waiting for a new CK_tile drop in the future?

[valarLip]

ck has two buffer load path with same signature.... and they changed the default path for MI35x, so...