composable_kernel icon indicating copy to clipboard operation
composable_kernel copied to clipboard

DTYPES : some newly created instances are not well protected

Open junliume opened this issue 2 years ago • 4 comments

[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

junliume avatar Sep 20 '23 04:09 junliume

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?

sgueko avatar Dec 22 '23 18:12 sgueko

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

illsilin avatar Jan 02 '24 16:01 illsilin