Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
22 commits
Select commit Hold shift + click to select a range
f8ef4c7
added reflection for conv_fwd_multiple_d_wmma_cshuffle.hpp
kabraham-streamhpc Jan 16, 2026
fbd0277
added reflection for device_grouped_conv_bwd_weight_xdl_cshuffle
kabraham-streamhpc Jan 16, 2026
640cb6b
added reflection for device_grouped_conv_bwd_weight_xdl_cshuffle v3
kabraham-streamhpc Jan 16, 2026
5da6fc2
added reflection of max_transpose parameters
kabraham-streamhpc Jan 16, 2026
67cbf4d
fix printing of std optional parameters
kabraham-streamhpc Jan 16, 2026
6382912
fix use of undefined ck::index
kabraham-streamhpc Jan 16, 2026
d16d945
added conv traits for device_grouped_conv_bwd_weight_multiple_d_xdl_c…
kabraham-streamhpc Jan 19, 2026
43b44d8
added xdl two stage instance to reflection
kabraham-streamhpc Jan 19, 2026
072ef16
added additional variables
kabraham-streamhpc Jan 19, 2026
2ab0c98
added reflection for grouped_conv_bwd_weight_multiple_d_wmma_cshuffle…
kabraham-streamhpc Jan 20, 2026
7b9ec3b
added reflection for device_grouped_conv_bwd_weigh_wmma_cshuffle_v3
kabraham-streamhpc Jan 20, 2026
f1b6789
added reflection for bwd_weight_wmma_cshuffle
kabraham-streamhpc Jan 20, 2026
9657d0e
added comments back in
kabraham-streamhpc Jan 20, 2026
6b8e425
add printed output for optional parameters
kabraham-streamhpc Jan 21, 2026
ef08ebd
update README
kabraham-streamhpc Jan 21, 2026
662dee9
fix typo
kabraham-streamhpc Jan 21, 2026
bd66848
added num_gemm_k_prefetch_stage and small fixes
kabraham-streamhpc Jan 21, 2026
28b4f74
modified test string due to reflection of new parameter
kabraham-streamhpc Jan 22, 2026
9e9ef0e
use ConvTraits directly in reflection
kabraham-streamhpc Jan 26, 2026
68b5bab
added num_gemm_k_prefetch_stage param
kabraham-streamhpc Jan 26, 2026
c982b71
added tests for optional parameters to ConvDescription and fixed prin…
kabraham-streamhpc Jan 26, 2026
48a50c9
fixed wrong rename of DeviceInstance
kabraham-streamhpc Jan 27, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 10 additions & 9 deletions experimental/builder/include/ck_tile/builder/reflect/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,7 @@ See the [main builder documentation](../README.md) for an overview.
The reflection system works by extracting properties from a convolution kernel *type* and formatting them into a string. This is useful for debugging, performance tuning, and generating documentation.

1. **Trait Extraction**: The `ConvTraits` template (in `conv_traits.hpp`) is specialized for each kernel instance. It extracts low-level details like tile sizes, data layouts, and pipeline versions from the kernel's type definition.
This template is common for xld and wmma, fwd and backwards weight kernels. std::optional is used for parameters that are only used by some kernels

2. **Description Generation**: The `describe<Instance>()` function (in `conv_description.hpp`) uses `ConvTraits` to populate a `ConvDescription` (`Description`) object.

Expand Down Expand Up @@ -48,6 +49,15 @@ The reflection system (`ckr::describe`) currently supports the following convolu
- **Standard XDL Forward Convolution** (`DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle`)
- **Large Tensor XDL Forward Convolution** (`DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor`)
- **V3 XDL Forward Convolution** (`DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3`)
- **V3 WMMA Forward Convolution** (`DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffleV3`)
- **XDL Backward Weight Convolution** (`DeviceGroupedConvBwdWeight_Xdl_CShuffle`)
- **V3 XDL Backward Weight Convolution** (`DeviceGroupedConvBwdWeight_Xdl_CShuffleV3`)
- **XDL Multiple D Backward Weight Convolution** (`DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle`)
- **Two Stage XDL Backward Weight Convolution** (`DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle`)
- **V3 Two Stage XDL Backward Weight Convolution** (`DeviceGroupedConvBwdWeightTwoStage_Wmma_CShuffleV3`)
- **Wmma Backward Weight Convolution** (`DeviceGroupedConvBwdWeight_Wmma_CShuffle`)
- **V3 Wmma Backward Weight Convolution** (`DeviceGroupedConvBwdWeight_Wmma_CShuffleV3`)
- **V3 Wmma Multiple D Backward Weight Convolution** (`DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffleV3`)

These variants all share similar template parameter structures and are compatible with the current `ConvTraits` implementation.

Expand All @@ -59,15 +69,6 @@ The following instance types are **not yet supported** by the reflection system:
- Uses different internal structure with parameters like `K0PerBlock`, `K1`, `M1PerThread`, etc.
- Missing standard members like `kKPerBlock`, `kMPerXDL`, `kAK1`

- **WMMA Variants** (`DeviceGroupedConvFwdMultipleD_Wmma_CShuffle`)
- Uses WMMA-specific parameters like `MPerWmma`, `NPerWmma`, `MRepeat`, `NRepeat`
- Different tile transfer structure incompatible with current `ConvTraits`

- **Backward Weight Convolution** (`DeviceGroupedConvBwdWeight_Xdl_CShuffle`)
- Uses different layout naming: `InLayout`, `WeiLayout`, `OutLayout` instead of `ALayout`, `BLayout`, `ELayout`
- Different specialization type: `ConvBackwardWeightSpecialization` vs `ConvForwardSpecialization`
- Missing several members expected by forward convolution traits

### Future Work

To support these additional instance types, the reflection system would need:
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -29,30 +29,7 @@ conv::ConvDescription describe()
const auto traits = conv::instance_to_conv_traits<Instance>();

return conv::ConvDescription(
conv::ConvSignatureInfo{
.spatial_dim = traits.spatial_dim,
.direction = traits.direction,
.input_layout = traits.layout[0],
.weight_layout = traits.layout[1],
.output_layout = traits.layout[2],
.data_type = traits.data_type,
.input_element_op = traits.input_element_op,
.weight_element_op = traits.weight_element_op,
.output_element_op = traits.output_element_op,
},
conv::GemmAlgorithmInfo{
.thread_block_size = traits.thread_block_size,
.tile_dims = traits.tile_dims,
.warp_gemm = traits.warp_gemm,
.a_tile_transfer = traits.a_tile_transfer,
.b_tile_transfer = traits.b_tile_transfer,
.c_tile_transfer = traits.c_tile_transfer,
.pipeline_version = traits.pipeline_version,
.pipeline_scheduler = traits.pipeline_scheduler,
.conv_specialization = traits.conv_specialization,
.padding = traits.gemm_padding,
},
[]<typename T = Instance>() { return reflect::instance_string<T>(); });
traits, []<typename T = Instance>() { return reflect::instance_string<T>(); });
}

} // namespace ck_tile::reflect
Loading