Skip to content

Conversation

@aosewski
Copy link
Collaborator

@aosewski aosewski commented Nov 4, 2025

Proposed changes

Added:

  1. Convolution traits & unit tests
  2. Update builder enumerators to have representation of Convolution Kernels properties.
  3. Unified builder pipeline version & scheduler enumerators

@aosewski aosewski requested a review from Copilot November 4, 2025 09:17
Copy link
Contributor

Copilot AI left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull Request Overview

This PR consolidates and unifies pipeline-related enums across the builder framework, replacing multiple specialized enum types (BlockGemmPipelineVersion, GridwiseGemmPipelineVersion, BlockGemmPipelineScheduler, LoopScheduler) with two unified enums (PipelineVersion and PipelineScheduler).

Key changes:

  • Unified enum types for pipeline configuration that simplify the type system
  • Added new backward convolution specialization and GEMM padding enums to support comprehensive reflection
  • Introduced a new ConvTraits reflection infrastructure for extracting compile-time kernel properties
  • Added comprehensive test coverage for the new ConvTraits functionality
  • Included necessary header includes (<string>) to support the new reflection code

Reviewed Changes

Copilot reviewed 20 out of 20 changed files in this pull request and generated 5 comments.

Show a summary per file
File Description
experimental/builder/include/ck_tile/builder/types.hpp Replaced multiple pipeline-related enums with unified PipelineVersion and PipelineScheduler enums; added ConvBwdDataSpecialization, ConvBwdWeightSpecialization, and GemmPadding enums
experimental/builder/include/ck_tile/builder/conv_factory.hpp Updated all references from old enum types to unified PipelineVersion and PipelineScheduler
experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp Updated concept requirements to use unified enum types
experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp New file implementing comprehensive compile-time reflection for extracting kernel properties from device instances
experimental/builder/include/ck_tile/builder/reflect/instance_traits.hpp Removed redundant includes, keeping only essential headers
experimental/builder/include/ck_tile/builder/reflect/instance_traits_util.hpp Added case-insensitive string comparison utility and necessary includes
experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp Added include for instance_traits_util.hpp
experimental/builder/include/ck_tile/builder/reflect/instance_traits_device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp Added include for instance_traits_util.hpp
experimental/builder/test/impl/conv_algorithm_types.hpp Updated struct definitions to use unified enum types
experimental/builder/test/utils/ckb_conv_test_common.hpp Updated test utilities to use unified enum types
experimental/builder/test/conv/test_ckb_conv_fwd_*.cpp Updated test instantiations to use PipelineVersion instead of BlockGemmPipelineVersion
experimental/builder/test/conv/test_conv_traits.cpp New comprehensive test file for ConvTraits reflection functionality
experimental/builder/test/CMakeLists.txt Added new test target for test_conv_traits.cpp and duplicate test file entry
include/ck/tensor_operation/gpu/device/convolution_backward_weight_specialization.hpp Added missing <string> include

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

V3, // Only used in stream-K implementation
V4,
V5,
WEIGHT_ONLY
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is this WEIGHT ONLY?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If I remember correctly it is some Navi specific pipeline for inference.

@@ -0,0 +1,719 @@
// SPDX-License-Identifier: MIT
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If you make changes to this PR, it would be helpful to update this to match the new style (no year, copyright on first line).

@shumway
Copy link
Collaborator

shumway commented Nov 4, 2025

Where do we want to make the jump from specialized (templated) structs to just a standard struct?

My original thought was this workflow:

InstanceTraits -> MakeConvTraits() -> ConvTraits

The logic is:

  1. InstanceTraits are tightly coupled to the instance, so they have to be a templated type.
  2. The function MakeConvTraits needs to read the specialized instance traits, so it needs to be templated and specialized.
  3. The struct ConvTraits is back to plain runtime C++ structs. Idiomatic C++, not metaprogramming.

In this PR, the MakeConvTraits is built into the ConvTraits specializaiton, and I don't know where we end the metaprogramming. I don't think there is a runtime case for making all the reflection code build time metaprogramming. Shouldn't we get back to simple C++ design sooner rather than later?

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants