ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-11 09:40:51 +00:00
Code Issues Packages Projects Releases Wiki Activity

Provide default value for API VERSION.

Browse Source 
Create a LATEST_API_VERSION value and use that as the default. This simplifies code that shouldn't depend the API version.
This commit is contained in:
John Shumway
2025-09-14 14:16:40 +00:00
parent fe6140353f
commit f4c11b6d22
7 changed files with 12 additions and 22 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
experimental/builder/include/ck_tile/builder/conv_builder.hpp
View File

@@ -6,7 +6,7 @@
#include <ck_tile/builder/conv_algorithm.hpp>
#include <ck_tile/builder/conv_factory.hpp>
#include <ck_tile/builder/conv_signature.hpp>
#include <ck_tile/builder/versions.h>
#include <ck_tile/builder/versions.hpp>
namespace ck_tile::builder {
@@ -24,7 +24,7 @@ namespace ck_tile::builder {
*/
template <ConvSignatureDescriptor auto SIGNATURE,
ConvAlgorithmDescriptor auto ALGORITHM,
StringLiteral VERSION>
StringLiteral VERSION = LATEST_API_VERSION>
requires SupportedVersion<VERSION> && ValidConvSignature<SIGNATURE>
struct ConvBuilder
{

2
experimental/builder/include/ck_tile/builder/conv_factory.hpp
View File

@@ -7,7 +7,7 @@
#include <ck_tile/builder/conv_algorithm.hpp>
#include <ck_tile/builder/builder_utils.hpp>
#include <ck_tile/builder/types.hpp>
#include <ck_tile/builder/versions.h>
#include <ck_tile/builder/versions.hpp>
namespace ck_tile::builder {

2
experimental/builder/include/ck_tile/builder/versions.h → experimental/builder/include/ck_tile/builder/versions.hpp
View File

@@ -10,6 +10,8 @@ namespace ck_tile::builder {
static constexpr StringLiteral V0_0_0 = "0.0.0";
static constexpr StringLiteral V0_1_0 = "0.1.0";
static constexpr StringLiteral LATEST_API_VERSION = V0_1_0;
template <StringLiteral V>
concept SupportedVersion = (V == V0_0_0) || (V == V0_1_0);

5
experimental/builder/test/test_conv_builder.cpp
View File

@@ -23,14 +23,11 @@ struct DefaultAlgorithm
};
static_assert(ckb::ConvAlgorithmDescriptor<DefaultAlgorithm>);
constexpr char API_VERSION[] = "0.1.0";
static_assert(ckb::SupportedVersion<API_VERSION>);
TEST(ConvBuilderTest, TestDefaultInstance)
{
static constexpr const ConvSignature SIGNATURE;
static constexpr const DefaultAlgorithm ALGORITHM;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM, API_VERSION>;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
EXPECT_EQ(
Builder::Instance::TypeString(),
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, 4, 4, "

11
experimental/builder/test/test_conv_grp_fwd_2d.cpp
View File

@@ -23,14 +23,11 @@ struct DefaultAlgorithm
};
static_assert(ckb::ConvAlgorithmDescriptor<DefaultAlgorithm>);
constexpr char API_VERSION[] = "0.1.0";
static_assert(ckb::SupportedVersion<API_VERSION>);
TEST(ConvBuilderGrpFwd2d, TestDefaultInstance)
{
static constexpr const ConvSignature SIGNATURE;
static constexpr const DefaultAlgorithm ALGORITHM;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM, API_VERSION>;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
EXPECT_EQ(
Builder::Instance::TypeString(),
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, 4, 4, "
@@ -41,7 +38,7 @@ TEST(ConvBuilderGrpFwd2d, TestDefaultFP32Instance)
{
static constexpr const ConvSignature SIGNATURE{.data_type = ckb::DataType::FP32};
static constexpr const DefaultAlgorithm ALGORITHM;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM, API_VERSION>;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
EXPECT_NE(Builder::Instance::TypeString(), "");
// It's difficult to check the types direction on the kernel, so we instead
// check that the builder has the correct data type aliases for FP32.
@@ -56,7 +53,7 @@ TEST(ConvBuilderGrpFwd2d, TestDefaultFP16Instance)
{
static constexpr const ConvSignature SIGNATURE{.data_type = ckb::DataType::FP16};
static constexpr const DefaultAlgorithm ALGORITHM;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM, API_VERSION>;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
// Check that the builder has the correct data type aliases for FP16.
EXPECT_TRUE((std::is_same_v<typename Builder::Factory::Types::ADataType, ck::half_t>));
@@ -70,7 +67,7 @@ TEST(ConvBuilderGrpFwd2d, TestDefaultBF16Instance)
{
static constexpr const ConvSignature SIGNATURE{.data_type = ckb::DataType::BF16};
static constexpr const DefaultAlgorithm ALGORITHM;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM, API_VERSION>;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
// Check that the builder has the correct data type aliases for BF16.
EXPECT_TRUE((std::is_same_v<typename Builder::Factory::Types::ADataType, ck::bhalf_t>));

5
experimental/builder/test/test_conv_grp_fwd_3d.cpp
View File

@@ -23,14 +23,11 @@ struct DefaultAlgorithm
};
static_assert(ckb::ConvAlgorithmDescriptor<DefaultAlgorithm>);
constexpr char API_VERSION[] = "0.1.0";
static_assert(ckb::SupportedVersion<API_VERSION>);
TEST(ConvBuilderGrpFwd3d, TestDefaultInstance)
{
static constexpr const ConvSignature SIGNATURE;
static constexpr const DefaultAlgorithm ALGORITHM;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM, API_VERSION>;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
EXPECT_EQ(
Builder::Instance::TypeString(),
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, 4, 4, "

5
experimental/builder/test/test_conv_instances.cpp
View File

@@ -24,9 +24,6 @@ struct ConvSignature
};
static_assert(ckb::ConvSignatureDescriptor<ConvSignature>);
constexpr char API_VERSION[] = "0.1.0";
static_assert(ckb::SupportedVersion<API_VERSION>);
// Defines the tunable algorithmic parameters for the convolution kernel.
// This includes thread block configuration, tuning parameters, data transfer
// settings, and the GEMM pipeline version.
@@ -259,7 +256,7 @@ TYPED_TEST(ConvBuilderInstancesTest, KernelParamsConfigured)
static constexpr const FwdConvAlgorithm& ALGORITHM =
ConvBuilderInstancesTest<TypeParam>::ALGORITHM;
static constexpr const ConvSignature SIGNATURE;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM, API_VERSION>;
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
EXPECT_EQ(Builder::Instance::TypeString(), ConvBuilderInstancesTest<TypeParam>::EXPECTED_TYPE);
const auto& tp = ALGORITHM.tuning_params;
EXPECT_EQ(Builder::Factory::TUNING.ak1, tp.ak1);