ROCm / composable_kernel

Composable Kernel: Performance Portable Programming Model for Machine Learning Tensor Operators
https://rocm.docs.amd.com/projects/composable_kernel/en/latest/
Other
289 stars 110 forks source link

DTYPES : some newly created instances are not well protected #928

Open junliume opened 11 months ago

junliume commented 11 months ago

[Reproduce]

CXX=/opt/rocm/llvm/bin/clang++ cmake -DCMAKE_CXX_COMPILER_LAUNCHER="${COMPILER_LAUNCHER}" -DCMAKE_PREFIX_PATH=/opt/rocm -DDTYPES="fp16;fp32;bf16" -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON -DGPU_TARGETS="gfx1100" ..

hence -DDTYPES="fp16;fp32;bf16", then

make -j$(nproc)

[Observation]

/home/junliu/composable_kernel/library/src/tensor_operation_instance/gpu/grouped_gemm_fixed_nk/device_grouped_gemm_xdl_fixed_nk_f16_f8_f16_mk_kn_mn_instance.cpp:18:17: error: no type named 'f8_t' in namespace 'ck'
using F8  = ck::f8_t;
            ~~~~^

Since fp8;bf8 was not included in the cmake command, resulting in the instances partially not protected. Hence

#if defined CK_ENABLE_FP8
using f8_t = _BitInt(8);
#endif
#if defined CK_ENABLE_BF8
using bf8_t = unsigned _BitInt(8);
#endif
sgueko commented 8 months ago

i have this same error

[ 18%] Building CXX object library/src/tensor_operation_instance/gpu/gemm/CMakeFiles/device_gemm_instance.dir/device_gemm_dpp_f16_f16_f16_km_kn_mn_instance.cpp.o In file included from /tmp/ROCm-SBo/composable_kernel/library/src/tensor_operation_instance/gpu/gemm/device_gemm_dpp_f16_f16_f16_km_kn_mn_instance.cpp:9: In file included from /tmp/ROCm-SBo/composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_gemm_dpp.hpp:14: In file included from /tmp/ROCm-SBo/composable_kernel/include/ck/tensor_operation/gpu/grid/gridwise_gemm_dpp.hpp:13: In file included from /tmp/ROCm-SBo/composable_kernel/include/ck/tensor_operation/gpu/block/blockwise_gemm_dpp.hpp:9: /tmp/ROCm-SBo/composable_kernel/include/ck/tensor_operation/gpu/warp/dpp_gemm.hpp:453:81: error: use of undeclared identifier 'f8_t' 453 | is_same<BaseType, int8_t>::value || is_same<BaseType, f8_t>::value, | ^ 1 error generated when compiling for gfx900.

cmake ... -DDTYPES="fp32;fp16;bf16" ...

resolve this by adding fp8;bf8 in my dtypes?

illsilin commented 8 months ago

Yes, since we've added a few mixed-type kernels, we cannot decouple the fp16 and fp8 types.