6.2.4 compilation issue - error: Illegal instruction detected

[cmyster]

I'm trying to compile and install ROCm 6.2.4 and I'm building it one step at a time, now stuck on composable_kernel with this error (full build log is here)

[90/1272] /usr/bin/hipcc -DCK_ENABLE_BF16 -DCK_ENABLE_BF8 -DCK_ENABLE_FP16 -DCK_ENABLE_FP32 -DCK_ENABLE_FP64 -DCK_ENABLE_FP8 -DCK_ENABLE_INT8 -DCK_USE_WMMA -DCK_USE_XDL -DUSE_PROF_API=1 -D__HIP_PLATFORM_AMD__=1 -D__HIP_PLATFORM_HCC__=1 -I/tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4/library/include -I/tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4/include -I/tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4_build/include -march=native -O2 -pipe -DNDEBUG -std=c++17 -fPIC -Wall -Wextra -Wcomment -Wendif-labels -Wformat -Winit-self -Wreturn-type -Wsequence-point -Wswitch -Wtrigraphs -Wundef -Wuninitialized -Wunreachable-code -Wunused -Wno-reserved-identifier -Wno-option-ignored -Wsign-compare -Wno-extra-semi-stmt -Wno-unused-template -Wno-missing-field-initializers -Wno-deprecated-declarations -Wall -Wextra -Wcomment -Wendif-labels -Wformat -Winit-self -Wreturn-type -Wsequence-point -Wswitch -Wtrigraphs -Wundef -Wuninitialized -Wunreachable-code -Wunused -Wno-reserved-identifier -Wno-option-ignored -Wsign-compare -Wno-extra-semi-stmt -Wno-unused-template -Weverything -Wno-c++98-compat -Wno-c++98-compat-pedantic-Wno-conversion -Wno-double-promotion -Wno-exit-time-destructors -Wno-extra-semi -Wno-float-conversion -Wno-gnu-anonymous-struct -Wno-gnu-zero-variadic-macro-arguments -Wno-missing-prototypes-Wno-nested-anon-types -Wno-padded -Wno-return-std-move-in-c++11 -Wno-shorten-64-to-32 -Wno-sign-conversion -Wno-unknown-warning-option -Wno-unused-command-line-argument -Wno-weak-vtables -Wno-covered-switch-default -Wno-unsafe-buffer-usage -Wno-unused-lambda-capture -Wno-bit-int-extension -Wno-pass-failed -Wno-switch-default -fno-offload-uniform-block -mllvm -enable-post-misched=0 -x hip --offload-arch=gfx1030 --offload-arch=gfx1100 --offload-arch=gfx906 --offload-arch=gfx908 --offload-arch=gfx90a --offload-arch=gfx942 --offload-arch=gfx908 --offload-arch=gfx90a --offload-arch=gfx942 -MD -MT library/src/tensor_operation_instance/gpu/contraction_bilinear/CMakeFiles/device_contraction_bilinear_instance.dir/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp.o -MF library/src/tensor_operation_instance/gpu/contraction_bilinear/CMakeFiles/device_contraction_bilinear_instance.dir/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp.o.d -o library/src/tensor_operation_instance/gpu/contraction_bilinear/CMakeFiles/device_contraction_bilinear_instance.dir/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp.o -c /tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4/library/src/tensor_operation_instance/gpu/contraction_bilinear/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp
FAILED: library/src/tensor_operation_instance/gpu/contraction_bilinear/CMakeFiles/device_contraction_bilinear_instance.dir/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp.o
/usr/bin/hipcc -DCK_ENABLE_BF16 -DCK_ENABLE_BF8 -DCK_ENABLE_FP16 -DCK_ENABLE_FP32 -DCK_ENABLE_FP64 -DCK_ENABLE_FP8 -DCK_ENABLE_INT8 -DCK_USE_WMMA -DCK_USE_XDL -DUSE_PROF_API=1 -D__HIP_PLATFORM_AMD__=1 -D__HIP_PLATFORM_HCC__=1 -I/tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4/library/include -I/tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4/include -I/tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4_build/include -march=native -O2 -pipe -DNDEBUG -std=c++17 -fPIC -Wall -Wextra -Wcomment -Wendif-labels -Wformat -Winit-self -Wreturn-type -Wsequence-point -Wswitch -Wtrigraphs -Wundef -Wuninitialized -Wunreachable-code -Wunused -Wno-reserved-identifier -Wno-option-ignored -Wsign-compare -Wno-extra-semi-stmt -Wno-unused-template -Wno-missing-field-initializers -Wno-deprecated-declarations -Wall -Wextra -Wcomment -Wendif-labels -Wformat -Winit-self -Wreturn-type -Wsequence-point -Wswitch -Wtrigraphs -Wundef -Wuninitialized -Wunreachable-code -Wunused -Wno-reserved-identifier -Wno-option-ignored -Wsign-compare -Wno-extra-semi-stmt -Wno-unused-template -Weverything -Wno-c++98-compat -Wno-c++98-compat-pedantic -Wno-conversion -Wno-double-promotion -Wno-exit-time-destructors -Wno-extra-semi -Wno-float-conversion -Wno-gnu-anonymous-struct -Wno-gnu-zero-variadic-macro-arguments -Wno-missing-prototypes -Wno-nested-anon-types -Wno-padded -Wno-return-std-move-in-c++11 -Wno-shorten-64-to-32 -Wno-sign-conversion -Wno-unknown-warning-option -Wno-unused-command-line-argument -Wno-weak-vtables -Wno-covered-switch-default -Wno-unsafe-buffer-usage -Wno-unused-lambda-capture -Wno-bit-int-extension -Wno-pass-failed -Wno-switch-default -fno-offload-uniform-block -mllvm -enable-post-misched=0 -x hip --offload-arch=gfx1030 --offload-arch=gfx1100 --offload-arch=gfx906 --offload-arch=gfx908 --offload-arch=gfx90a --offload-arch=gfx942 --offload-arch=gfx908 --offload-arch=gfx90a --offload-arch=gfx942 -MD -MT library/src/tensor_operation_instance/gpu/contraction_bilinear/CMakeFiles/device_contraction_bilinear_instance.dir/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp.o -MF library/src/tensor_operation_instance/gpu/contraction_bilinear/CMakeFiles/device_contraction_bilinear_instance.dir/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp.o.d -o library/src/tensor_operation_instance/gpu/contraction_bilinear/CMakeFiles/device_contraction_bilinear_instance.dir/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp.o -c /tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4/library/src/tensor_operation_instance/gpu/contraction_bilinear/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp
error: Illegal instruction detected: Operand has incorrect register class.
renamable $vgpr1_vgpr2_vgpr3_vgpr4 = V_ACCVGPR_READ_B32_e64 killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
error: Illegal instruction detected: Operand has incorrect register class.
renamable $vgpr1_vgpr2_vgpr3_vgpr4 = V_ACCVGPR_READ_B32_e64 killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
error: Illegal instruction detected: Operand has incorrect register class.
renamable $vgpr12_vgpr13_vgpr14_vgpr15 = V_ACCVGPR_READ_B32_e64 killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
error: Illegal instruction detected: Operand has incorrect register class.
renamable $vgpr1_vgpr2_vgpr3_vgpr4 = V_ACCVGPR_READ_B32_e64 killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
error: Illegal instruction detected: Operand has incorrect register class.
renamable $vgpr1_vgpr2_vgpr3_vgpr4 = V_ACCVGPR_READ_B32_e64 killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
error: Illegal instruction detected: Operand has incorrect register class.
renamable $vgpr11_vgpr12_vgpr13_vgpr14 = V_ACCVGPR_READ_B32_e64 killed $agpr0_agpr1_agpr2_agpr3, implicit $exec
6 errors generated when compiling for gfx908.
failed to execute:/usr/lib/llvm/bin/clang++ --offload-arch=gfx1030 --offload-arch=gfx1100 --offload-arch=gfx906 --offload-arch=gfx908 --offload-arch=gfx90a --offload-arch=gfx942 --offload-arch=gfx908 --offload-arch=gfx90a --offload-arch=gfx942 -DCK_ENABLE_BF16 -DCK_ENABLE_BF8 -DCK_ENABLE_FP16 -DCK_ENABLE_FP32 -DCK_ENABLE_FP64 -DCK_ENABLE_FP8 -DCK_ENABLE_INT8 -DCK_USE_WMMA -DCK_USE_XDL -DUSE_PROF_API=1 -D__HIP_PLATFORM_AMD__=1 -D__HIP_PLATFORM_HCC__=1 -I/tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4/library/include -I/tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4/include -I/tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4_build/include -march=native -O2 -pipe -DNDEBUG -std=c++17 -fPIC -Wall -Wextra -Wcomment -Wendif-labels -Wformat -Winit-self -Wreturn-type -Wsequence-point -Wswitch -Wtrigraphs -Wundef -Wuninitialized -Wunreachable-code -Wunused -Wno-reserved-identifier -Wno-option-ignored -Wsign-compare -Wno-extra-semi-stmt -Wno-unused-template -Wno-missing-field-initializers -Wno-deprecated-declarations -Wall -Wextra -Wcomment -Wendif-labels -Wformat -Winit-self -Wreturn-type -Wsequence-point -Wswitch -Wtrigraphs -Wundef -Wuninitialized -Wunreachable-code -Wunused -Wno-reserved-identifier -Wno-option-ignored -Wsign-compare -Wno-extra-semi-stmt -Wno-unused-template -Weverything -Wno-c++98-compat -Wno-c++98-compat-pedantic -Wno-conversion -Wno-double-promotion -Wno-exit-time-destructors -Wno-extra-semi -Wno-float-conversion -Wno-gnu-anonymous-struct -Wno-gnu-zero-variadic-macro-arguments -Wno-missing-prototypes -Wno-nested-anon-types -Wno-padded -Wno-return-std-move-in-c++11 -Wno-shorten-64-to-32 -Wno-sign-conversion -Wno-unknown-warning-option -Wno-unused-command-line-argument -Wno-weak-vtables -Wno-covered-switch-default -Wno-unsafe-buffer-usage -Wno-unused-lambda-capture -Wno-bit-int-extension -Wno-pass-failed -Wno-switch-default -fno-offload-uniform-block -mllvm -enable-post-misched=0 -x hip -MD -MT library/src/tensor_operation_instance/gpu/contraction_bilinear/CMakeFiles/device_contraction_bilinear_instance.dir/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp.o -MF library/src/tensor_operation_instance/gpu/contraction_bilinear/CMakeFiles/device_contraction_bilinear_instance.dir/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp.o.d -o "library/src/tensor_operation_instance/gpu/contraction_bilinear/CMakeFiles/device_contraction_bilinear_instance.dir/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp.o" -c /tmp/var/tmp/portage/sci-libs/composable-kernel-6.2.4/work/composable_kernel-rocm-6.2.4/library/src/tensor_operation_instance/gpu/contraction_bilinear/6D/device_contraction_bilinear_m6_n6_k6_xdl_c_shuffle_f64_f64_f64_f64_compute_f32_knnn_instance.cpp

[ppanchad-amd]

Hi @cmyster. Internal ticket has been created to investigate your issue. Thanks!

[zichguan-amd]

Hi @cmyster, based on your log it seems that you are using hipcc but hipcc is not using the ROCm packaged llvm

-- The CXX compiler identification is Clang 18.1.8 
-- The HIP compiler identification is Clang 18.1.8

I can observe the error with upstream llvm 18.1.8, using the ROCm packaged llvm should resolve the issue.
You can make sure that hipcc uses the right clang by setting HIP_CLANG_PATH, or specify the compiler for cmake with -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/amdclang++.