[composable_kernel] gfx architectures need to be excluded due to missing support

[marbre]

Building composable kernel fails for several architectures that are not (yet) supported.

• gfx906

100.1	[187/694] Building CXX object library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/CMakeFiles/device_grouped_conv2d_bwd_data_instance.dir/xdl/device_grouped_conv2d_bwd_data_xdl_ngchw_gkcyx_ngkhw_bf16_instance.cpp.o
100.1	FAILED: library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/CMakeFiles/device_grouped_conv2d_bwd_data_instance.dir/xdl/device_grouped_conv2d_bwd_data_xdl_ngchw_gkcyx_ngkhw_bf16_instance.cpp.o 
100.1	ccache /__w/TheRock/TheRock/build/core/clr/dist/lib/llvm/bin/clang++ -DCK_ENABLE_BF16 -DCK_ENABLE_BF8 -DCK_ENABLE_FP16 -DCK_ENABLE_FP32 -DCK_ENABLE_FP64 -DCK_ENABLE_FP8 -DCK_ENABLE_INT8 -DCK_TIME_KERNEL=1 -DCK_USE_FNUZ_FP8 -DCK_USE_XDL -DDL_KERNELS -DDPP_KERNELS -DUSE_PROF_API=1 -D__HIP_PLATFORM_AMD__=1 -D__HIP_PLATFORM_HCC__=1 -I/__w/TheRock/TheRock/ml-libs/composable_kernel/library/include -I/__w/TheRock/TheRock/ml-libs/composable_kernel/include -I/__w/TheRock/TheRock/build/ml-libs/composable_kernel/build/include -I/__w/TheRock/TheRock/build/profiler/roctracer/stage/include -I/__w/TheRock/TheRock/build/base/half/stage/include -isystem /__w/TheRock/TheRock/build/core/clr/dist/include -Wno-documentation-unknown-command -Wno-documentation-pedantic -Wno-unused-command-line-argument -Wno-explicit-specialization-storage-class --hip-path=/__w/TheRock/TheRock/build/core/clr/dist --hip-device-lib-path=/__w/TheRock/TheRock/build/core/clr/dist/lib/llvm/amdgcn/bitcode -O3 -DNDEBUG -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-error=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-nvcc-compat -Wno-bit-int-extension -Wno-pass-failed -Wno-switch-default -Wno-unique-object-duplication -Wno-nrvo -fno-offload-uniform-block -mllvm --lsr-drop-solution=1 -mllvm -enable-post-misched=0 -mllvm -amdgpu-coerce-illegal-types=1 -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false -Werror -Weverything -fcolor-diagnostics --offload-compress -x hip --offload-arch=gfx906 --offload-arch=gfx908 --offload-arch=gfx90a -std=c++17 --offload-arch=gfx908 --offload-arch=gfx90a -MD -MT library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/CMakeFiles/device_grouped_conv2d_bwd_data_instance.dir/xdl/device_grouped_conv2d_bwd_data_xdl_ngchw_gkcyx_ngkhw_bf16_instance.cpp.o -MF library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/CMakeFiles/device_grouped_conv2d_bwd_data_instance.dir/xdl/device_grouped_conv2d_bwd_data_xdl_ngchw_gkcyx_ngkhw_bf16_instance.cpp.o.d -o library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/CMakeFiles/device_grouped_conv2d_bwd_data_instance.dir/xdl/device_grouped_conv2d_bwd_data_xdl_ngchw_gkcyx_ngkhw_bf16_instance.cpp.o -c /__w/TheRock/TheRock/ml-libs/composable_kernel/library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/xdl/device_grouped_conv2d_bwd_data_xdl_ngchw_gkcyx_ngkhw_bf16_instance.cpp
100.1	In file included from /__w/TheRock/TheRock/ml-libs/composable_kernel/library/src/tensor_operation_instance/gpu/grouped_conv2d_bwd_data/xdl/device_grouped_conv2d_bwd_data_xdl_ngchw_gkcyx_ngkhw_bf16_instance.cpp:5:
100.1	In file included from /__w/TheRock/TheRock/ml-libs/composable_kernel/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_bwd_data/device_grouped_conv_bwd_data_xdl_instance.hpp:8:
100.1	/__w/TheRock/TheRock/ml-libs/composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp:97:27: error: unused parameter 'KBatch' [-Werror,-Wunused-parameter]
100.1	   97 |             const index_t KBatch)
100.1	      |                           ^
100.1	1 error generated when compiling for gfx906.

• gfx1010, gfx1011, gfx1102:

9.6	FAILED: library/src/utility/CMakeFiles/utility.dir/convolution_parameter.cpp.o 
9.6	ccache /__w/TheRock/TheRock/build/core/clr/dist/lib/llvm/bin/clang++ -DCK_ENABLE_BF16 -DCK_ENABLE_BF8 -DCK_ENABLE_FP16 -DCK_ENABLE_FP32 -DCK_ENABLE_FP64 -DCK_ENABLE_FP8 -DCK_ENABLE_INT8 -DCK_TIME_KERNEL=1 -DDL_KERNELS -DDPP_KERNELS -DUSE_PROF_API=1 -D__HIP_PLATFORM_AMD__=1 -D__HIP_PLATFORM_HCC__=1 -I/__w/TheRock/TheRock/ml-libs/composable_kernel/library/include -I/__w/TheRock/TheRock/ml-libs/composable_kernel/include -I/__w/TheRock/TheRock/build/ml-libs/composable_kernel/build/include -I/__w/TheRock/TheRock/build/profiler/roctracer/stage/include -I/__w/TheRock/TheRock/build/base/half/stage/include -isystem /__w/TheRock/TheRock/build/core/clr/dist/include -Wno-documentation-unknown-command -Wno-documentation-pedantic -Wno-unused-command-line-argument -Wno-explicit-specialization-storage-class --hip-path=/__w/TheRock/TheRock/build/core/clr/dist --hip-device-lib-path=/__w/TheRock/TheRock/build/core/clr/dist/lib/llvm/amdgcn/bitcode -O3 -DNDEBUG -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-error=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-nvcc-compat -Wno-bit-int-extension -Wno-pass-failed -Wno-switch-default -Wno-unique-object-duplication -Wno-nrvo -fno-offload-uniform-block -mllvm --lsr-drop-solution=1 -mllvm -enable-post-misched=0 -mllvm -amdgpu-coerce-illegal-types=1 -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false -Werror -Weverything -fcolor-diagnostics -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-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-nvcc-compat -x hip --offload-arch=gfx1010 --offload-arch=gfx1011 --offload-arch=gfx1012 -std=c++17 -MD -MT library/src/utility/CMakeFiles/utility.dir/convolution_parameter.cpp.o -MF library/src/utility/CMakeFiles/utility.dir/convolution_parameter.cpp.o.d -o library/src/utility/CMakeFiles/utility.dir/convolution_parameter.cpp.o -c /__w/TheRock/TheRock/ml-libs/composable_kernel/library/src/utility/convolution_parameter.cpp
9.6	In file included from /__w/TheRock/TheRock/ml-libs/composable_kernel/library/src/utility/convolution_parameter.cpp:4:
9.6	In file included from /__w/TheRock/TheRock/ml-libs/composable_kernel/include/ck/host_utility/io.hpp:11:
9.6	In file included from /__w/TheRock/TheRock/ml-libs/composable_kernel/include/ck/tensor_description/tensor_descriptor.hpp:6:
9.6	In file included from /__w/TheRock/TheRock/ml-libs/composable_kernel/include/ck/utility/common_header.hpp:37:
9.6	/__w/TheRock/TheRock/ml-libs/composable_kernel/include/ck/utility/amd_buffer_addressing_builtins.hpp:32:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
9.6	   32 |     wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD;
9.6	      |                                                ^
9.6	/__w/TheRock/TheRock/ml-libs/composable_kernel/include/ck/utility/amd_buffer_addressing_builtins.hpp:47:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
9.6	   47 |     wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD;
9.6	      |                                                ^
9.6	/__w/TheRock/TheRock/ml-libs/composable_kernel/include/ck/utility/amd_buffer_addressing_builtins.hpp:60:22: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
9.6	   60 |     auto flags     = CK_BUFFER_RESOURCE_3RD_DWORD;
9.6	      |                      ^
9.6	/__w/TheRock/TheRock/ml-libs/composable_kernel/include/ck/utility/amd_buffer_addressing_builtins.hpp:72:22: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
9.6	   72 |     auto flags     = CK_BUFFER_RESOURCE_3RD_DWORD;
9.6	      |                      ^
9.6	4 errors generated when compiling for gfx1010.

Building MIOpen with CK support was already disabled with 42b760e.

[ScottTodd]

For gfx906, does it build without -Werror or -Wunused-parameter?

[marbre]

For gfx906, does it build without -Werror or -Wunused-parameter?

I haven't tested. If building for gfx906 it is at least not enabled to be used in MIOpen.

[TheTrustedComputer]

Regarding the gfx906, try removing all instances of -Werror from the CMake files (CMakeLists.txt, EnableCompilerWarnings.cmake, etc) and recompile.

As for gfx1010, this discussion enables experimental build support for RDNA1 targets.

[zackfuchtel]

Regarding the gfx906, try removing all instances of -Werror from the CMake files (CMakeLists.txt, EnableCompilerWarnings.cmake, etc) and recompile.

Like this?

grep -rl "\-Werror\ " . | xargs sed -i 's/\-Werror\ //g'

[GrecAndrei]

Note: gfx103x rows in therock_amdgpu_targets.cmake previously linked composable_kernel excludes to #4836 (MIOPEN mixed-arch), which is a different issue and already fixed.

TheRock#6444 retargets those excludes to this issue (#1245), which is the correct “CK not supported on these arches” tracker.