diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp index 042950d0c1..4fb73cf672 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp @@ -25,7 +25,7 @@ struct ConvTraits builder::ElementwiseOperation weight_element_op; builder::ElementwiseOperation output_element_op; - builder::GemmPadding gemm_padding; + builder::GemmPadding gemm_padding = builder::GemmPadding::DEFAULT; std::variant diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_bwd_weight_wmma_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_bwd_weight_wmma_cshuffle.hpp index 2538c48372..a483da7113 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_bwd_weight_wmma_cshuffle.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_bwd_weight_wmma_cshuffle.hpp @@ -27,7 +27,7 @@ instance_to_conv_traits() .spatial_dim = InstTraits::kSpatialDim, .direction = conv_direction(), .layout = conv_layout(), - .data_type = conv_data_type(), + .data_type = conv_data_type(), .input_element_op = elementwise_op(), .weight_element_op = elementwise_op(), .output_element_op = elementwise_op(), diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp index 0b9d8f5ea1..0aa19c64a3 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp @@ -25,11 +25,10 @@ constexpr ConvTraits instance_to_conv_traits() .spatial_dim = InstTraits::kSpatialDim, .direction = conv_direction(), .layout = bwd_wei_conv_layout(), - .data_type = conv_data_type(), + .data_type = conv_data_type(), .input_element_op = elementwise_op(), .weight_element_op = elementwise_op(), .output_element_op = elementwise_op(), - .gemm_padding = gemm_spec(), .conv_specialization = conv_spec(), .thread_block_size = InstTraits::kBlockSize, .tile_dims = conv_traits_data_tile(InstTraits::kK0PerBlock), diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp index 4e3e5a5e9a..37cd09f1e6 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp @@ -25,7 +25,7 @@ constexpr ConvTraits instance_to_conv_traits() .spatial_dim = InstTraits::kSpatialDim, .direction = conv_direction(), .layout = fwd_conv_layout(), - .data_type = conv_data_type(), + .data_type = conv_data_type(), .input_element_op = elementwise_op(), .weight_element_op = elementwise_op(), .output_element_op = elementwise_op(), diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp index 32cf532ccf..37cc9a2c5e 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp @@ -25,7 +25,7 @@ constexpr ConvTraits instance_to_conv_traits() .spatial_dim = InstTraits::kSpatialDim, .direction = conv_direction(), .layout = fwd_conv_layout(), - .data_type = conv_data_type(), + .data_type = conv_data_type(), .input_element_op = elementwise_op(), .weight_element_op = elementwise_op(), .output_element_op = elementwise_op(), diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp index 9c907ae875..5e88ee1783 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp @@ -25,7 +25,7 @@ constexpr ConvTraits instance_to_conv_traits() .spatial_dim = InstTraits::kSpatialDim, .direction = conv_direction(), .layout = fwd_conv_layout(), - .data_type = conv_data_type(), + .data_type = conv_data_type(), .input_element_op = elementwise_op(), .weight_element_op = elementwise_op(), .output_element_op = elementwise_op(), diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp index 719d3cb6d5..553fb1638f 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp @@ -25,7 +25,7 @@ constexpr ConvTraits instance_to_conv_traits() .spatial_dim = InstTraits::kSpatialDim, .direction = conv_direction(), .layout = fwd_conv_layout(), - .data_type = conv_data_type(), + .data_type = conv_data_type(), .input_element_op = elementwise_op(), .weight_element_op = elementwise_op(), .output_element_op = elementwise_op(), diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_helpers.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_helpers.hpp index 6644c2ddd8..381d32016f 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_traits_helpers.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_traits_helpers.hpp @@ -433,11 +433,11 @@ constexpr auto bwd_wei_conv_layout() // ---------------------------------------------------------------------------- /// @brief Helper function to report unsupported data type with a clear error message. -template +template [[noreturn]] consteval void report_unsupported_data_type_error() { throw "Unsupported data type detected!\n" - "The ADataType is not recognized.\n" + "The DataTypeFromInstance is not recognized.\n" "Supported types are: ck::half_t (FP16), ck::Tuple (FP16_FP16), " "ck::bhalf_t (BF16), ck::Tuple (BF16_BF16), float (FP32), " "ck::Tuple (FP32_FP32), double (FP64), ck::f8_t (FP8), ck::bf8_fnuz_t " @@ -448,43 +448,42 @@ template /// @brief Derives the data type from a device kernel `Instance` type. /// Returns a `builder::DataType` enum value (e.g., FP16, BF16, FP32, BF8). -template +// Note: maybe move to types.hpp? +template constexpr builder::DataType conv_data_type() - requires HasFwdConvDataTypes> + { - using InstTraits = InstanceTraits; - using ADataType = typename InstTraits::ADataType; using enum builder::DataType; - if constexpr(std::is_same_v) + if constexpr(std::is_same_v) return FP16; - else if constexpr(std::is_same_v>) + else if constexpr(std::is_same_v>) return FP16_FP16; - else if constexpr(std::is_same_v) + else if constexpr(std::is_same_v) return BF16; - else if constexpr(std::is_same_v>) + else if constexpr(std::is_same_v>) return BF16_BF16; - else if constexpr(std::is_same_v) + else if constexpr(std::is_same_v) return FP32; - else if constexpr(std::is_same_v>) + else if constexpr(std::is_same_v>) return FP32_FP32; - else if constexpr(std::is_same_v) + else if constexpr(std::is_same_v) return FP64; - else if constexpr(std::is_same_v) + else if constexpr(std::is_same_v) return FP8; - else if constexpr(std::is_same_v) + else if constexpr(std::is_same_v) return BF8; - else if constexpr(std::is_same_v) + else if constexpr(std::is_same_v) return BF8; - else if constexpr(std::is_same_v) + else if constexpr(std::is_same_v) return I8; - else if constexpr(std::is_same_v>) + else if constexpr(std::is_same_v>) return I8_I8; - else if constexpr(std::is_same_v) + else if constexpr(std::is_same_v) return U8; else { - report_unsupported_data_type_error(); + report_unsupported_data_type_error(); return FP32; // Unreachable } } @@ -685,12 +684,11 @@ constexpr auto get_pipeline_scheduler() // ============================================================================ template -constexpr InputTileTransferInfo conv_traits_a_transfer_params(int _k1) +constexpr InputTileTransferInfo +conv_traits_a_transfer_params(int _k1, int kPerBlock = InstTraits::kKPerBlock) { return InputTileTransferInfo{ - .tile_dimensions = {.k0 = InstTraits::kKPerBlock / _k1, - .m_or_n = InstTraits::kMPerBlock, - .k1 = _k1}, + .tile_dimensions = {.k0 = kPerBlock / _k1, .m_or_n = InstTraits::kMPerBlock, .k1 = _k1}, .transfer_params = {.k1 = _k1, .thread_cluster_dims = InstTraits::kAThreadClusterLengths, .thread_cluster_order = InstTraits::kAThreadClusterArrangeOrder, @@ -703,12 +701,11 @@ constexpr InputTileTransferInfo conv_traits_a_transfer_params(int _k1) } template -constexpr InputTileTransferInfo conv_traits_b_transfer_params(int _k1) +constexpr InputTileTransferInfo +conv_traits_b_transfer_params(int _k1, int kPerBlock = InstTraits::kKPerBlock) { return InputTileTransferInfo{ - .tile_dimensions = {.k0 = InstTraits::kKPerBlock / _k1, - .m_or_n = InstTraits::kNPerBlock, - .k1 = _k1}, + .tile_dimensions = {.k0 = kPerBlock / _k1, .m_or_n = InstTraits::kNPerBlock, .k1 = _k1}, .transfer_params = {.k1 = _k1, .thread_cluster_dims = InstTraits::kBThreadClusterLengths, .thread_cluster_order = InstTraits::kBThreadClusterArrangeOrder, diff --git a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp index 29b0d3f0d3..13e25e5ebe 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_bwd_weight_xdl_cshuffle.hpp @@ -172,7 +172,7 @@ struct InstanceTraits::value; + static constexpr auto kAThreadClusterArrangeOrder = + detail::SequenceToArray::value; + static constexpr auto kABlockTransferSrcAccessOrder = + detail::SequenceToArray::value; + static constexpr ck::index_t kABlockTransferSrcVectorDim = ABlockTransferSrcVectorDim; static constexpr ck::index_t kABlockTransferSrcScalarPerVector = ABlockTransferSrcScalarPerVector; - static constexpr ck::index_t kABlockTransferDstScalarPerVector_K1 = + static constexpr ck::index_t kABlockTransferDstScalarPerVectorK1 = ABlockTransferDstScalarPerVector_K1; - static constexpr bool kABlockLdsAddExtraM = ABlockLdsAddExtraM; + static constexpr bool kABlockLdsExtraM = ABlockLdsAddExtraM; using BBlockTransferThreadClusterLengths_K0_N_K1 = BBlockTransferThreadClusterLengths_K0_N_K1_; using BBlockTransferThreadClusterArrangeOrder = BBlockTransferThreadClusterArrangeOrder_; using BBlockTransferSrcAccessOrder = BBlockTransferSrcAccessOrder_; + + // B block transfer thread cluster dimensions (converted to std::array) + static constexpr auto kBThreadClusterLengths = + detail::SequenceToArray::value; + static constexpr auto kBThreadClusterArrangeOrder = + detail::SequenceToArray::value; + static constexpr auto kBBlockTransferSrcAccessOrder = + detail::SequenceToArray::value; static constexpr ck::index_t kBBlockTransferSrcVectorDim = BBlockTransferSrcVectorDim; static constexpr ck::index_t kBBlockTransferSrcScalarPerVector = BBlockTransferSrcScalarPerVector; - static constexpr ck::index_t kBBlockTransferDstScalarPerVector_K1 = + static constexpr ck::index_t kBBlockTransferDstScalarPerVectorK1 = BBlockTransferDstScalarPerVector_K1; - static constexpr bool kBBlockLdsAddExtraN = BBlockLdsAddExtraN; + static constexpr bool kBBlockLdsExtraN = BBlockLdsAddExtraN; static constexpr ck::index_t kCShuffleMXdlPerWavePerShuffle = CShuffleMXdlPerWavePerShuffle; static constexpr ck::index_t kCShuffleNXdlPerWavePerShuffle = CShuffleNXdlPerWavePerShuffle; @@ -250,30 +266,30 @@ struct InstanceTraits(); // 22. oss << "," << detail::sequence_name(); // 23. oss << "," << detail::sequence_name(); // 24. oss << "," << kABlockTransferSrcVectorDim; // 25. oss << "," << kABlockTransferSrcScalarPerVector; // 26. - oss << "," << kABlockTransferDstScalarPerVector_K1; // 27. - oss << "," << (kABlockLdsAddExtraM ? "true" : "false"); // 28. + oss << "," << kABlockTransferDstScalarPerVectorK1; // 27. + oss << "," << (kABlockLdsExtraM ? "true" : "false"); // 28. oss << "," << detail::sequence_name(); // 29. oss << "," << detail::sequence_name(); // 30. oss << "," << detail::sequence_name(); // 31. oss << "," << kBBlockTransferSrcVectorDim; // 32. oss << "," << kBBlockTransferSrcScalarPerVector; // 33. - oss << "," << kBBlockTransferDstScalarPerVector_K1; // 34. - oss << "," << (kBBlockLdsAddExtraN ? "true" : "false"); // 35. + oss << "," << kBBlockTransferDstScalarPerVectorK1; // 34. + oss << "," << (kBBlockLdsExtraN ? "true" : "false"); // 35. oss << "," << kCShuffleMXdlPerWavePerShuffle; // 36. oss << "," << kCShuffleNXdlPerWavePerShuffle; // 37. oss << "," diff --git a/experimental/builder/test/conv/ck/test_conv_traits.cpp b/experimental/builder/test/conv/ck/test_conv_traits.cpp index 33f9bce8cf..e0288d4cc2 100644 --- a/experimental/builder/test/conv/ck/test_conv_traits.cpp +++ b/experimental/builder/test/conv/ck/test_conv_traits.cpp @@ -87,7 +87,7 @@ TEST_F(ConvTraitsTest, ConvBwdWeightMultipleDTraitsExtraction) // Verify signature information EXPECT_EQ(traits.spatial_dim, 2); - EXPECT_EQ(traits.direction, ConvDirection::FORWARD); + EXPECT_EQ(traits.direction, ConvDirection::BACKWARD_WEIGHT); EXPECT_THAT(traits.layout, ElementsAre(TensorLayout::GNHWC, TensorLayout::GKYXC, TensorLayout::GNHWK)); EXPECT_EQ(traits.data_type, DataType::FP16); @@ -97,8 +97,8 @@ TEST_F(ConvTraitsTest, ConvBwdWeightMultipleDTraitsExtraction) // Verify specializations EXPECT_EQ(traits.gemm_padding, ck_tile::builder::GemmPadding::DEFAULT); - EXPECT_EQ(std::get(traits.conv_specialization), - ck_tile::builder::ConvFwdSpecialization::DEFAULT); + EXPECT_EQ(std::get(traits.conv_specialization), + ck_tile::builder::ConvBwdWeightSpecialization::DEFAULT); // Verify algorithm information EXPECT_EQ(traits.thread_block_size, 256); @@ -147,7 +147,7 @@ TEST_F(ConvTraitsTest, ConvBwdWeightMultipleDTraitsExtraction) EXPECT_EQ(traits.c_tile_transfer.scalar_per_vector, 8); // Verify pipeline configuration - EXPECT_EQ(traits.pipeline_scheduler, PipelineScheduler::INTRAWAVE); + EXPECT_EQ(traits.pipeline_scheduler, PipelineScheduler::DEFAULT); EXPECT_EQ(traits.pipeline_version, PipelineVersion::V1); }