diff --git a/.gitignore b/.gitignore index c361863baf..27e53fa342 100644 --- a/.gitignore +++ b/.gitignore @@ -80,4 +80,4 @@ __pycache__/ !experimental/builder/** # Ignore the generated .cpp files for the CK Builder -experimental/builder/generated/*.cpp +experimental/builder/codegen/*.cpp diff --git a/experimental/builder/CMakeLists.txt b/experimental/builder/CMakeLists.txt index 103acbad55..3d17e140f5 100644 --- a/experimental/builder/CMakeLists.txt +++ b/experimental/builder/CMakeLists.txt @@ -1,3 +1,4 @@ if(BUILD_TESTING) add_subdirectory(test) + add_subdirectory(codegen) endif() diff --git a/experimental/builder/codegen/CMakeLists.txt b/experimental/builder/codegen/CMakeLists.txt new file mode 100644 index 0000000000..4ca1303397 --- /dev/null +++ b/experimental/builder/codegen/CMakeLists.txt @@ -0,0 +1,27 @@ +# CMakeLists.txt for ConvBuilder generated instantiations + +# Use GLOB to collect all generated instance files +file(GLOB CONV_INSTANCE_SOURCES "conv_instances_batch_*.cpp") + +# Check if any files were found +if(NOT CONV_INSTANCE_SOURCES) + message(WARNING "No conv_instances_batch_*.cpp files found in ${CMAKE_CURRENT_SOURCE_DIR}") +endif() + +# Add all generated instance files as a library +add_library(ckb_instances STATIC ${CONV_INSTANCE_SOURCES}) + +target_include_directories(ckb_instances PUBLIC + ${PROJECT_SOURCE_DIR}/include + ${PROJECT_SOURCE_DIR}/experimental/builder/include + ${CMAKE_CURRENT_SOURCE_DIR}/../test/impl +) + +target_compile_features(ckb_instances PUBLIC cxx_std_20) + +# Add necessary compile options for the builder +target_compile_options(ckb_instances PRIVATE + -Wno-global-constructors + -Wno-c++20-compat + -ftemplate-backtrace-limit=0 # Helpful for template debugging +) diff --git a/experimental/builder/examples/example_usage.cpp b/experimental/builder/examples/example_usage.cpp new file mode 100644 index 0000000000..80a62c27cb --- /dev/null +++ b/experimental/builder/examples/example_usage.cpp @@ -0,0 +1,105 @@ +// Copyright (C) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +/** + * Example: Using ConvBuilder-generated kernel instances + * + * This example demonstrates how to use the automatically generated + * convolution kernel instances from the ConvBuilder system. + */ + +#include +#include "ck_tile/builder/conv_builder.hpp" +#include "experimental/builder/test/impl/conv_signature_types.hpp" +#include "experimental/builder/test/impl/conv_algorithm_types.hpp" + +using namespace ck_tile::builder; +using namespace ck_tile::builder::test; + +int main() { + // Example 1: Define a convolution signature at compile-time + constexpr ConvSignature my_signature = { + .spatial_dim = 2, + .direction = ConvDirection::FORWARD, + .layout = GroupConvLayout{._2d = GroupConvLayout2D::GNHWC_GKYXC_GNHWK}, + .data_type = DataType::FP16, + .elementwise_operation = ElementwiseOperation::PASS_THROUGH + }; + + // Example 2: Define an algorithm at compile-time (Standard XDL) + constexpr ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle my_algorithm = { + .thread_block = { + .block_size = 256, + .tile_size = {.m = 128, .n = 128, .k = 32} + }, + .gridwise_gemm = { + .ak1 = 8, + .bk1 = 8, + .m_per_xdl = 32, + .n_per_xdl = 32, + .m_xdl_per_wave = 2, + .n_xdl_per_wave = 2 + }, + .block_transfer = { + .block_transfer_a = {.k0 = 4, .m_n = 64, .k1 = 1}, + .block_transfer_b = {.k0 = 4, .m_n = 64, .k1 = 1}, + .thread_cluster_dims_c = { + .m_block = 1, + .m_wave_per_xdl = 32, + .n_block = 1, + .n_wave_per_xdl = 8 + }, + .lds_transfer_a = { + .src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = true + }, + .lds_transfer_b = { + .src_vector_dim = 2, + .src_scalar_per_vector = 8, + .lds_dst_scalar_per_vector = 8, + .is_direct_load = false, + .lds_padding = true + }, + .epilogue_c = { + .m_per_wave_per_shuffle = 1, + .n_per_wave_per_shuffle = 1, + .scalar_per_vector = 8 + }, + .block_transfer_access_order_a = {.order = {1, 0, 2}}, + .block_transfer_access_order_b = {.order = {1, 0, 2}}, + .src_access_order_a = {.order = {1, 0, 2}}, + .src_access_order_b = {.order = {1, 0, 2}} + }, + .fwd_specialization = ConvFwdSpecialization::DEFAULT, + .gemm_specialization = GemmSpecialization::MNKPadding, + .num_gemm_k_prefetch_stages = 1, + .num_groups_to_merge = 1, + .loop_scheduler = LoopScheduler::DEFAULT + }; + + // Example 3: Instantiate ConvBuilder with signature and algorithm + using MyConvBuilder = ConvBuilder; + using MyKernelInstance = MyConvBuilder::Instance; + + std::cout << "ConvBuilder instantiation successful!\n"; + std::cout << "Kernel instance type created.\n"; + + // Example 4: Alternatively, include pre-generated instances + // #include "experimental/builder/generated/conv_instances_batch_00.cpp" + // using namespace ck_tile::builder::generated::batch_0; + // using PreGeneratedKernel = Instance_0; // Use pre-generated instance + + return 0; +} + +/* + * To build this example: + * + * mkdir -p build && cd build + * cmake .. -DBUILD_TESTING=ON + * make example_conv_builder_usage + * ./bin/example_conv_builder_usage + */ diff --git a/experimental/builder/generated/CMakeLists.txt b/experimental/builder/generated/CMakeLists.txt deleted file mode 100644 index db845a9be7..0000000000 --- a/experimental/builder/generated/CMakeLists.txt +++ /dev/null @@ -1,13 +0,0 @@ -# CMakeLists.txt for ConvBuilder generated instantiations - -# Use GLOB to collect all generated instance files -file(GLOB CONV_INSTANCE_SOURCES "conv_instances_batch_*.cpp") - -# Add all generated instance files as a library -add_library(ckb_generated_instances STATIC ${CONV_INSTANCE_SOURCES}) - -target_include_directories(ckb_generated_instances PUBLIC - ${CMAKE_SOURCE_DIR}/include - ${CMAKE_SOURCE_DIR}/experimental/builder/include - ${CMAKE_SOURCE_DIR}/experimental/builder/test/impl -) diff --git a/experimental/builder/instances/README.md b/experimental/builder/instances/README.md index 67fcce8e38..ff8e41e7e8 100644 --- a/experimental/builder/instances/README.md +++ b/experimental/builder/instances/README.md @@ -24,7 +24,7 @@ experimental/builder/ ├── instances/ # JSON instantiation database │ └── forward_conv_structured_instantiations.json │ └── generate_conv_builder_instances.py # Code generation scripts -└── generated/ # Generated C++ files (created by script) +└── codegen/ # Generated C++ files (created by script) ├── CMakeLists.txt └── conv_instances_batch_*.cpp ``` @@ -33,7 +33,7 @@ experimental/builder/ ### JSON-to-C++ Mapping -The JSON file contains 753 convolution instantiations, each with: +The JSON file contains convolution instantiations, each with: - **Signature**: Mathematical interface (spatial dim, direction, layouts, data types, elementwise op) - **Algorithm**: Implementation parameters (block sizes, tiling, pipeline configuration) @@ -83,20 +83,14 @@ python3 experimental/builder/scripts/generate_conv_builder_instances.py --help ### Build Generated Instances ```bash -# Create build directory -mkdir -p build && cd build - -# Configure CMake -cmake .. -DBUILD_TESTING=ON - # Build the generated instances library -make ck_builder_generated_instances +make ckb_instances ``` ### Use in Your Code ```cpp -#include "experimental/builder/generated/conv_instances_batch_00.cpp" +#include "experimental/builder/codegen/conv_instances_batch_00.cpp" using namespace ck_tile::builder::generated::batch_0; @@ -146,7 +140,7 @@ To add a new instantiation: 3. Rebuild To modify the generator: -1. Edit `scripts/generate_conv_builder_instances.py` +1. Edit `instances/generate_conv_builder_instances.py` 2. Test with `--batch-size 5` for quick iteration 3. Regenerate and verify build diff --git a/experimental/builder/instances/generate_conv_builder_instances.py b/experimental/builder/instances/generate_conv_builder_instances.py index 83d31de8df..f9a0dab72b 100644 --- a/experimental/builder/instances/generate_conv_builder_instances.py +++ b/experimental/builder/instances/generate_conv_builder_instances.py @@ -4,7 +4,7 @@ Generate C++ source files with ConvBuilder instantiations from JSON. This script reads forward_conv_structured_instantiations.json and generates batched C++ files containing compile-time ConvBuilder instantiations for all -753 device operation entries. +device operation entries. """ import json @@ -353,8 +353,8 @@ def generate_cpp_file(batch_idx: int, instantiations: List[Dict], output_dir: Pa // Batch {batch_idx} of convolution builder instantiations #include "ck_tile/builder/conv_builder.hpp" -#include "ck_tile/builder/test/impl/conv_signature_types.hpp" -#include "ck_tile/builder/test/impl/conv_algorithm_types.hpp" +#include "conv_signature_types.hpp" +#include "conv_algorithm_types.hpp" namespace ck_tile::builder::generated::batch_{batch_idx} {{ @@ -476,10 +476,10 @@ def main(): parser = argparse.ArgumentParser(description='Generate ConvBuilder instantiations from JSON') parser.add_argument('--json', type=str, - default='experimental/builder/instances/forward_conv_structured_instantiations.json', + default='./forward_conv_structured_instantiations.json', help='Path to JSON file') parser.add_argument('--output', type=str, - default='experimental/builder/generated', + default='../codegen', help='Output directory for generated files') parser.add_argument('--batch-size', type=int, default=50, help='Number of instantiations per C++ file') @@ -503,10 +503,6 @@ def main(): generate_cmake_file(output_dir, cpp_files) print("\n✓ Code generation complete!") - print(f"\nTo build:") - print(f" cd build") - print(f" cmake .. -DBUILD_TESTING=ON") - print(f" make ck_builder_generated_instances") return 0