[Issue]: Compiling PyTorch tag v2.8.0 to support the ROCM backend resulted in an error: "composablekernel" compilation error

[WsqRichards]

ROCm/rocm-libraries#3153

[WsqRichards]

In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_256_256x256x32_32x32_4x4_16x16x1_16x16x1_1x16x1x16_4_Intrawave_v4.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:11:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_xdl_cshuffle_v3.hpp:9:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/common_header.hpp:33:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/thread_group.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:10:39: error: constexpr function never produces a constant expression [-Winvalid-constexpr]
10 | host device constexpr index_t get_warp_size()
| ^~~~~~~~~~~~~
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:13:12: note: non-constexpr function 'operator int' cannot be used in a constant expression
13 | return warpSize;
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_warp_functions.h:114:52: note: declared here
114 | device attribute((always_inline, const)) operator int() const noexcept {
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_256_256x256x32_32x32_4x4_4x64x1_4x64x1_1x16x1x16_4_Intrawave_v5.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:6:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/ck_types.h:21:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/data_type.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_ck_fp8.hpp:431:9: warning: '__assert_ocp_support' macro redefined [-Wmacro-redefined]
431 | #define __assert_ocp_support(interp)
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_hip_fp8.h:162:9: note: previous definition is here
162 | #define __assert_ocp_support(interp)
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_256_256x256x32_32x32_4x4_4x64x1_4x64x1_1x16x1x16_4_Intrawave_v5.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:6:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/ck_types.h:21:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/data_type.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_ck_fp8.hpp:439:9: warning: '__assert_fnuz_support' macro redefined [-Wmacro-redefined]
439 | #define __assert_fnuz_support(interp)
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_hip_fp8.h:168:9: note: previous definition is here
168 | #define __assert_fnuz_support(interp)
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_256_256x224x64_16x16_8x7_8x32x1_8x32x1_1x32x1x8_4_Intrawave_v3.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:11:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_xdl_cshuffle_v3.hpp:9:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/common_header.hpp:33:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/thread_group.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:10:39: error: constexpr function never produces a constant expression [-Winvalid-constexpr]
10 | host device constexpr index_t get_warp_size()
| ^~~~~~~~~~~~~
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:13:12: note: non-constexpr function 'operator int' cannot be used in a constant expression
13 | return warpSize;
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_warp_functions.h:114:52: note: declared here
114 | device attribute((always_inline, const)) operator int() const noexcept {
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_256_256x256x32_32x32_4x4_8x32x1_8x32x1_1x16x1x16_4_Intrawave_v4.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:6:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/ck_types.h:21:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/data_type.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_ck_fp8.hpp:431:9: warning: '__assert_ocp_support' macro redefined [-Wmacro-redefined]
431 | #define __assert_ocp_support(interp)
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_hip_fp8.h:162:9: note: previous definition is here
162 | #define __assert_ocp_support(interp)
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_256_256x256x32_32x32_4x4_8x32x1_8x32x1_1x16x1x16_4_Intrawave_v4.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:6:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/ck_types.h:21:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/data_type.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_ck_fp8.hpp:439:9: warning: '__assert_fnuz_support' macro redefined [-Wmacro-redefined]
439 | #define __assert_fnuz_support(interp)
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_hip_fp8.h:168:9: note: previous definition is here
168 | #define __assert_fnuz_support(interp)
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_64_16x16x64_16x16_1x1_8x8x1_8x8x1_1x16x1x4_4_Intrawave_v1.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:11:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_xdl_cshuffle_v3.hpp:9:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/common_header.hpp:33:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/thread_group.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:10:39: error: constexpr function never produces a constant expression [-Winvalid-constexpr]
10 | host device constexpr index_t get_warp_size()
| ^~~~~~~~~~~~~
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:13:12: note: non-constexpr function 'operator int' cannot be used in a constant expression
13 | return warpSize;
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_warp_functions.h:114:52: note: declared here
114 | device attribute((always_inline, const)) operator int() const noexcept {
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_64_16x16x64_16x16_1x1_8x8x1_8x8x1_1x16x1x4_4_Intrawave_v2.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:11:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_xdl_cshuffle_v3.hpp:9:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/common_header.hpp:33:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/thread_group.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:10:39: error: constexpr function never produces a constant expression [-Winvalid-constexpr]
10 | host device constexpr index_t get_warp_size()
| ^~~~~~~~~~~~~
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:13:12: note: non-constexpr function 'operator int' cannot be used in a constant expression
13 | return warpSize;
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_warp_functions.h:114:52: note: declared here
114 | device attribute((always_inline, const)) operator int() const noexcept {
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_256_256x256x32_32x32_4x4_4x64x1_4x64x1_1x16x1x16_4_Intrawave_v5.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:11:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_xdl_cshuffle_v3.hpp:9:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/common_header.hpp:33:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/thread_group.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:10:39: error: constexpr function never produces a constant expression [-Winvalid-constexpr]
10 | host device constexpr index_t get_warp_size()
| ^~~~~~~~~~~~~
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:13:12: note: non-constexpr function 'operator int' cannot be used in a constant expression
13 | return warpSize;
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_warp_functions.h:114:52: note: declared here
114 | device attribute((always_inline, const)) operator int() const noexcept {
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_256_256x16x64_16x16_4x1_8x32x1_8x16x1_1x32x1x8_2_Intrawave_v2.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:6:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/ck_types.h:21:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/data_type.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_ck_fp8.hpp:431:9: warning: '__assert_ocp_support' macro redefined [-Wmacro-redefined]
431 | #define __assert_ocp_support(interp)
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_hip_fp8.h:162:9: note: previous definition is here
162 | #define __assert_ocp_support(interp)
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_256_256x16x64_16x16_4x1_8x32x1_8x16x1_1x32x1x8_2_Intrawave_v2.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:6:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/ck_types.h:21:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/data_type.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_ck_fp8.hpp:439:9: warning: '__assert_fnuz_support' macro redefined [-Wmacro-redefined]
439 | #define __assert_fnuz_support(interp)
| ^
/home/richard/rocm-all-libs-build/rocm-install/include/hip/amd_detail/amd_hip_fp8.h:168:9: note: previous definition is here
168 | #define __assert_fnuz_support(interp)
| ^
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_256_256x256x32_32x32_4x4_8x32x1_8x32x1_1x16x1x16_4_Intrawave_v4.hip:3:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/native/hip/bgemm_kernels/bgemm_kernel_template.h:11:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multiple_d_xdl_cshuffle_v3.hpp:9:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/common_header.hpp:33:
In file included from /home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/thread_group.hpp:6:
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:10:39: error: constexpr function never produces a constant expression [-Winvalid-constexpr]
10 | host device constexpr index_t get_warp_size()
| ^~~~~~~~~~~~~
/home/richard/pytorch-build/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/get_id.hpp:13:12: note: non-constexpr function 'operator int' cannot be used in a constant expression
13 | return warpSize;

[WsqRichards]

rocm7.0.0