diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index c6c09eb6ca..a2196ad2b2 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -37,9 +37,7 @@ set(REGRESSION_TESTS test_grouped_convnd_bwd_data_xdl test_conv_tensor_rearrange test_gemm_mx - test_ck_tile_batched_transpose_fp8 - test_ck_tile_batched_transpose_fp16 - test_ck_tile_batched_transpose_bf16 + test_ck_tile_batched_transpose ) function(add_test_executable TEST_NAME) diff --git a/test/ck_tile/batched_transpose/CMakeLists.txt b/test/ck_tile/batched_transpose/CMakeLists.txt index f2ef158a4d..111b7c2bed 100644 --- a/test/ck_tile/batched_transpose/CMakeLists.txt +++ b/test/ck_tile/batched_transpose/CMakeLists.txt @@ -1,33 +1,7 @@ # Currently ck_tile is only built on gfx9 if(GPU_TARGETS MATCHES "gfx9") - - function (add_batched_transpose_test TARGET_NAME MAIN_SRC) - message(DEBUG "adding ${TARGET_NAME}") - - add_test_executable(${TARGET_NAME} ${MAIN_SRC} batched_transpose_api.cpp) - target_include_directories(${TARGET_NAME} PRIVATE ${CMAKE_CURRENT_LIST_DIR}) - - # NOTE: we turn off undefined-func-template to let source compile without explicit declare function specializations - list(APPEND EXAMPLE_BATCHED_TRANSPOSE_COMPILE_OPTIONS -Wno-undefined-func-template -Wno-float-equal) - # list(APPEND EXAMPLE_BATCHED_TRANSPOSE_COMPILE_OPTIONS -v --save-temps -Wno-gnu-line-marker) - target_compile_options(${TARGET_NAME} PRIVATE ${EXAMPLE_BATCHED_TRANSPOSE_COMPILE_OPTIONS}) - - endfunction(add_batched_transpose_test TARGET_NAME MAIN_SRC) - - set(CUSTOM_TARGET_NAME test_ck_tile_batched_transpose) - - add_custom_target(${CUSTOM_TARGET_NAME}) - - add_batched_transpose_test(test_ck_tile_batched_transpose_fp16 batched_transpose_fp16.cpp) - add_dependencies(${CUSTOM_TARGET_NAME} test_ck_tile_batched_transpose_fp16) - - add_batched_transpose_test(test_ck_tile_batched_transpose_fp8 batched_transpose_fp8.cpp) - add_dependencies(${CUSTOM_TARGET_NAME} test_ck_tile_batched_transpose_fp8) - - add_batched_transpose_test(test_ck_tile_batched_transpose_bf16 batched_transpose_bf16.cpp) - add_dependencies(${CUSTOM_TARGET_NAME} test_ck_tile_batched_transpose_bf16) - add_gtest_executable(test_batched_transpose test_batched_transpose.cpp) - set_property(TARGET test_batched_transpose PROPERTY CXX_STANDARD 20) + add_gtest_executable(test_ck_tile_batched_transpose test_batched_transpose.cpp) + set_property(TARGET test_ck_tile_batched_transpose PROPERTY CXX_STANDARD 20) else() message(DEBUG "Skipping ck_tile batched_transpose tests for current target") endif() diff --git a/test/ck_tile/batched_transpose/batched_transpose.hpp b/test/ck_tile/batched_transpose/batched_transpose.hpp deleted file mode 100644 index bd1abb1191..0000000000 --- a/test/ck_tile/batched_transpose/batched_transpose.hpp +++ /dev/null @@ -1,25 +0,0 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT -#include "ck_tile/core.hpp" -#include "ck_tile/host.hpp" -#include "ck_tile/ops/reduce.hpp" -#include "ck_tile/ops/batched_transpose.hpp" - -#include -#include - -#pragma once - -struct batched_transpose_trait -{ - std::string type; - std::string layout; -}; - -struct batched_transpose_kargs : public ck_tile::BatchedTransposeHostArgs -{ -}; - -float batched_transpose(batched_transpose_trait t, - batched_transpose_kargs a, - ck_tile::stream_config s); diff --git a/test/ck_tile/batched_transpose/batched_transpose.inc b/test/ck_tile/batched_transpose/batched_transpose.inc deleted file mode 100644 index 30084f5664..0000000000 --- a/test/ck_tile/batched_transpose/batched_transpose.inc +++ /dev/null @@ -1,283 +0,0 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT - -#include -#include -#include -#include -#include -#include -#include -#include - -#include "batched_transpose.hpp" - -// different threshold for different dtype -template -auto get_elimit(std::string /*init_method*/) -{ - double rtol = 1e-3; - double atol = 1e-3; - return ck_tile::make_tuple(rtol, atol); -} - -template <> -auto get_elimit(std::string /*init_method*/) -{ - double rtol = 1e-2; - double atol = 1e-2; - return ck_tile::make_tuple(rtol, atol); -} - -template <> -auto get_elimit(std::string init_method) -{ - if(init_method == "ui" || init_method == "ni") - { - unsigned max_rounding_point_distance = 0; - double atol = 2e-3; - return ck_tile::make_tuple(max_rounding_point_distance, atol); - } - else - { - unsigned max_rounding_point_distance = 1; - double atol = 0.0625; - return ck_tile::make_tuple(max_rounding_point_distance, atol); - } -} - -auto create_args(int argc, char* argv[], int index = 0) -{ - ck_tile::ArgParser arg_parser; - arg_parser.insert("v", "1", "whether do CPU validation or not") - .insert("pr", "fp16", "input data type. fp16/fp32 (representing 8/16/32 bit data)") - .insert("N", "1", "input batch size. ") - .insert("C", "64", "input channel size.") - .insert("H", "18", "input height size.") - .insert("W", "64", "input width size. ") - .insert("layout_in", "NCHW", "input tensor data layout - NCHW by default") - .insert("layout_out", "NHWC", "output tensor data layout - NHWC by default ") - .insert("warmup", "50", "number of iterations before benchmark the kernel") - .insert("repeat", "100", "number of iterations to benchmark the kernel") - .insert("seed", "-1", "seed to be used, -1 means random every time") - .insert("kname", "0", "t to 1 will print kernel name"); - - bool result = arg_parser.parse(argc, argv, index); - return std::make_tuple(result, arg_parser); -} - -template -bool run_batched_transpose(ck_tile::ArgParser args) -{ - int validate = args.get_int("v"); - std::string prec = args.get_str("pr"); - int N = args.get_int("N"); - int C = args.get_int("C"); - int H = args.get_int("H"); - int W = args.get_int("W"); - int n_warmup = args.get_int("warmup"); - int n_repeat = args.get_int("repeat"); - std::string layout_in = args.get_str("layout_in"); - std::string layout_out = args.get_str("layout_out"); - int seed = args.get_int("seed"); - - int dim_in[4], dim_out[4]; - int stride_dim_in[4], stride_dim_out[4]; - bool nchw2nhwc = layout_in == "NCHW" && layout_out == "NHWC"; - bool nhwc2nchw = layout_in == "NHWC" && layout_out == "NCHW"; - assert(nchw2nhwc != nhwc2nchw); - (void)nhwc2nchw; - - dim_in[0] = N; - dim_in[1] = nchw2nhwc ? C : H; - dim_in[2] = nchw2nhwc ? H : W; - dim_in[3] = nchw2nhwc ? W : C; - dim_out[0] = N; - dim_out[1] = nchw2nhwc ? H : C; - dim_out[2] = nchw2nhwc ? W : H; - dim_out[3] = nchw2nhwc ? C : W; - stride_dim_in[0] = C * H * W; - stride_dim_in[1] = nchw2nhwc ? H * W : C * W; - stride_dim_in[2] = nchw2nhwc ? W : C; - stride_dim_in[3] = 1; - stride_dim_out[0] = C * H * W; - stride_dim_out[1] = nchw2nhwc ? C * W : H * W; - stride_dim_out[2] = nchw2nhwc ? C : W; - stride_dim_out[3] = 1; - - if(seed < 0) - { - seed = std::time(nullptr); - } - - ck_tile::HostTensor x_host( - {dim_in[0], dim_in[1], dim_in[2], dim_in[3]}, - {stride_dim_in[0], stride_dim_in[1], stride_dim_in[2], stride_dim_in[3]}); - ck_tile::HostTensor y_host( - {dim_out[0], dim_out[1], dim_out[2], dim_out[3]}, - {stride_dim_out[0], stride_dim_out[1], stride_dim_out[2], stride_dim_out[3]}); - - ck_tile::FillUniformDistribution{-.5f, .5f}(x_host); - - ck_tile::DeviceMem x_dev(x_host.get_element_space_size_in_bytes()); - ck_tile::DeviceMem y_dev(y_host.get_element_space_size_in_bytes()); - - x_dev.ToDevice(x_host.data()); - - auto trait = batched_transpose_trait{prec, layout_in}; - - uint32_t height = nchw2nhwc ? C : H * W; - uint32_t width = nchw2nhwc ? H * W : C; - - batched_transpose_kargs karg = [&]() { - batched_transpose_kargs a_; - a_.p_input = x_dev.GetDeviceBuffer(); - a_.p_output = y_dev.GetDeviceBuffer(); - a_.batch = N; - a_.height = height; - a_.width = width; - return a_; - }(); - - ck_tile::stream_config sc{nullptr, true, n_warmup, n_repeat}; - - auto ms = batched_transpose(trait, karg, sc); - - std::size_t num_operations = N * C * H * (W - 1); - std::size_t num_bytes = N * C * H * W * sizeof(Type); - - float ave_time = ms * 1E-3; - float gb_per_sec = num_bytes / ms * 1.E-6; - float tflops = static_cast(num_operations) / ms * 1.E-6; - - std::cout << "Run Batched Transpose kernel with N=" << N << ", C=" << C << ", H=" << H - << ", W=" << W << ", layout_in=" << layout_in << ", layout_out=" << layout_out - << " : " << ms << " ms (" << ave_time << " ave_time), " << tflops << " TFlops" - << gb_per_sec << " GB/s, " << std::endl; - - printf("[%s]N:%d, C:%d, H:%d, W:%d, layout_in:%s, %f\n", - prec.c_str(), - N, - C, - H, - W, - layout_in.c_str(), - ms); - if(ms < 0) - printf("------------------------------------not " - "supported-------------------------------------\n"); - fflush(stdout); - - if(ms < 0) - { - return false; - } - - y_dev.FromDevice(y_host.data()); - - bool rtn = true; - if(validate) - { - // this host buffer will not copy to GPU, so no need use stride - ck_tile::HostTensor y_ref( - {dim_out[0], dim_out[1], dim_out[2], dim_out[3]}, - {stride_dim_out[0], stride_dim_out[1], stride_dim_out[2], stride_dim_out[3]}); - - ck_tile::reference_batched_transpose(x_host, y_ref, layout_in, layout_out); - - auto [rtol, atol] = get_elimit(""); - - rtn &= ck_tile::check_err( - y_host, y_ref, std::string("y Error: Incorrect results!"), rtol, atol); - } - printf("-----------------------------------------------------------------------valid:%s--------" - "--------------------------------------------------------------------\n", - rtn ? "y" : "n"); - fflush(stdout); - return rtn; -} - -template -bool run_test_case(int argc, char** argv) -{ - auto [result, args] = create_args(argc, argv); - if(!result) - return false; - - return run_batched_transpose(args); -} - -template -bool run_test_cases(std::vector>& test_cases) -{ - bool valid = true; - for(std::size_t test_idx = 0; test_idx < test_cases.size(); ++test_idx) - { - constexpr int num_args = 7; - char* argv[num_args]; - - assert(test_cases[test_idx].size() == num_args && - "invalid number of arguments in test case"); - - for(std::size_t idx = 0; idx < test_cases[test_idx].size(); ++idx) - { - argv[idx] = test_cases[test_idx][idx].data(); - } - - valid = valid && run_test_case(num_args, argv); - - if(!valid) - break; - } - - return valid; -} - -std::vector> generate_test_cases(const std::string prec) -{ - return { - {"-pr=" + prec, "-N=1", "-C=32", "-H=1", "-W=32", "-layout_in=NCHW", "-layout_out=NHWC"}, - {"-pr=" + prec, "-N=1", "-C=64", "-H=1", "-W=64", "-layout_in=NCHW", "-layout_out=NHWC"}, - {"-pr=" + prec, "-N=2", "-C=12", "-H=1", "-W=32", "-layout_in=NHWC", "-layout_out=NCHW"}, - {"-pr=" + prec, "-N=3", "-C=1334", "-H=1", "-W=37", "-layout_in=NHWC", "-layout_out=NCHW"}, - {"-pr=" + prec, "-N=4", "-C=27", "-H=1", "-W=32", "-layout_in=NCHW", "-layout_out=NHWC"}, - {"-pr=" + prec, "-N=5", "-C=1234", "-H=1", "-W=12", "-layout_in=NCHW", "-layout_out=NHWC"}, - {"-pr=" + prec, "-N=1", "-C=1", "-H=1", "-W=1", "-layout_in=NCHW", "-layout_out=NHWC"}, - {"-pr=" + prec, "-N=1", "-C=1", "-H=1", "-W=1", "-layout_in=NHWC", "-layout_out=NCHW"}, - {"-pr=" + prec, - "-N=128", - "-C=1024", - "-H=64", - "-W=64", - "-layout_in=NCHW", - "-layout_out=NHWC"}, - {"-pr=" + prec, - "-N=128", - "-C=1024", - "-H=64", - "-W=64", - "-layout_in=NHWC", - "-layout_out=NCHW"}, - {"-pr=" + prec, "-N=16", "-C=64", "-H=32", "-W=128", "-layout_in=NCHW", "-layout_out=NHWC"}, - {"-pr=" + prec, "-N=16", "-C=64", "-H=128", "-W=32", "-layout_in=NHWC", "-layout_out=NCHW"}, - {"-pr=" + prec, "-N=1", "-C=2048", "-H=1", "-W=1", "-layout_in=NCHW", "-layout_out=NHWC"}, - {"-pr=" + prec, "-N=1", "-C=2048", "-H=1", "-W=1", "-layout_in=NHWC", "-layout_out=NCHW"}, - {"-pr=" + prec, - "-N=1", - "-C=1", - "-H=1024", - "-W=1024", - "-layout_in=NCHW", - "-layout_out=NHWC"}, - {"-pr=" + prec, - "-N=1", - "-C=1", - "-H=1024", - "-W=1024", - "-layout_in=NHWC", - "-layout_out=NCHW"}, - {"-pr=" + prec, "-N=8", "-C=16", "-H=8", "-W=16", "-layout_in=NCHW", "-layout_out=NHWC"}, - {"-pr=" + prec, "-N=8", "-C=16", "-H=8", "-W=16", "-layout_in=NHWC", "-layout_out=NCHW"}, - {"-pr=" + prec, "-N=1", "-C=64", "-H=1", "-W=1024", "-layout_in=NCHW", "-layout_out=NHWC"}, - {"-pr=" + prec, "-N=1", "-C=64", "-H=1024", "-W=1", "-layout_in=NHWC", "-layout_out=NCHW"}}; -} diff --git a/test/ck_tile/batched_transpose/batched_transpose_api.cpp b/test/ck_tile/batched_transpose/batched_transpose_api.cpp deleted file mode 100644 index 973a1967f2..0000000000 --- a/test/ck_tile/batched_transpose/batched_transpose_api.cpp +++ /dev/null @@ -1,109 +0,0 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT -#include "batched_transpose.hpp" - -template -float batched_transpose_dispatch(batched_transpose_kargs& a, ck_tile::stream_config& s) -{ - uint32_t dim_stride = a.height * a.width; - - a.dim_stride = dim_stride; - a.dim_block_h = block_y; - a.dim_block_w = block_x; - - using block_tile = ck_tile::sequence; - using warp_layout = ck_tile::sequence; - - using ts_problem = - ck_tile::BatchedTransposeProblem; - using ts_pipeline = ck_tile::BatchedTransposePipeline; - - using kernel = ck_tile::BatchedTransposeKernel; - - auto kargs = kernel::MakeKargs(a); - - const dim3 grids = kernel::GridSize(a); - constexpr dim3 blocks = kernel::BlockSize(); - - printf("Grid: %u %u %u\n", grids.x, grids.y, grids.z); - printf("Block: %u %u %u\n", blocks.x, blocks.y, blocks.z); - printf("kargs: kargs.batch %d kargs.height %d kargs.width %d kargs.dim_strid %d\n", - kargs.batch, - kargs.height, - kargs.width, - kargs.dim_stride); - - printf("Launching Kernel...\n"); - - float ave_time = ck_tile::launch_kernel( - s, ck_tile::make_kernel(kernel{}, grids, blocks, 0, kargs)); - - printf("Kernel finished...\n"); - - return ave_time; -} - -// Param Comb: type_size, block_x & y, warp_x & y, thread_x & y -#define FOREACH_TRANSPOSE_PARAM(F) \ - F(fp8, ck_tile::fp8_t, 64, 64, 1, 1, true, true) \ - F(fp8, ck_tile::fp8_t, 64, 64, 1, 1, false, false) \ - F(fp16, ck_tile::fp16_t, 64, 64, 1, 1, true, true) \ - F(fp16, ck_tile::fp16_t, 64, 64, 1, 1, false, false) \ - F(bf16, ck_tile::bf16_t, 64, 64, 1, 1, true, true) \ - F(bf16, ck_tile::bf16_t, 64, 64, 1, 1, false, false) - -// Macro that defines one static function per line -#define GEN_TRANSPOSE_FN(SHORT_NAME, REAL_TYPE, BX, BY, WX, WY, PADM, PADN) \ - static float transpose_fn_##SHORT_NAME##_##BX##_##BY##_##WX##_##WY##_##PADM##_##PADN( \ - batched_transpose_kargs& a, ck_tile::stream_config& s) \ - { \ - return batched_transpose_dispatch(a, s); \ - } - -FOREACH_TRANSPOSE_PARAM(GEN_TRANSPOSE_FN) - -float batched_transpose(batched_transpose_trait t, - batched_transpose_kargs a, - ck_tile::stream_config s) -{ - if(t.type == "fp8") - { - if(a.height % 64 == 0 && a.width % 64 == 0) - { - return transpose_fn_fp8_64_64_1_1_false_false(a, s); - } - else - { - return transpose_fn_fp8_64_64_1_1_true_true(a, s); - } - } - else if(t.type == "fp16") - { - if(a.height % 64 == 0 && a.width % 64 == 0) - { - return transpose_fn_fp16_64_64_1_1_false_false(a, s); - } - else - { - return transpose_fn_fp16_64_64_1_1_true_true(a, s); - } - } - else if(t.type == "bf16") - { - if(a.height % 64 == 0 && a.width % 64 == 0) - { - return transpose_fn_bf16_64_64_1_1_false_false(a, s); - } - else - { - return transpose_fn_bf16_64_64_1_1_true_true(a, s); - } - } - return -1; -} diff --git a/test/ck_tile/batched_transpose/batched_transpose_bf16.cpp b/test/ck_tile/batched_transpose/batched_transpose_bf16.cpp deleted file mode 100644 index 42642335f6..0000000000 --- a/test/ck_tile/batched_transpose/batched_transpose_bf16.cpp +++ /dev/null @@ -1,10 +0,0 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT -#include "batched_transpose.inc" - -int main() -{ - std::vector> test_cases = generate_test_cases("bf16"); - - return !run_test_cases(test_cases); -} diff --git a/test/ck_tile/batched_transpose/batched_transpose_fp16.cpp b/test/ck_tile/batched_transpose/batched_transpose_fp16.cpp deleted file mode 100644 index 5562dd54e8..0000000000 --- a/test/ck_tile/batched_transpose/batched_transpose_fp16.cpp +++ /dev/null @@ -1,10 +0,0 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT -#include "batched_transpose.inc" - -int main() -{ - std::vector> test_cases = generate_test_cases("fp16"); - - return !run_test_cases(test_cases); -} diff --git a/test/ck_tile/batched_transpose/batched_transpose_fp8.cpp b/test/ck_tile/batched_transpose/batched_transpose_fp8.cpp deleted file mode 100644 index 45e79fb4c2..0000000000 --- a/test/ck_tile/batched_transpose/batched_transpose_fp8.cpp +++ /dev/null @@ -1,10 +0,0 @@ -// Copyright © Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT -#include "batched_transpose.inc" - -int main() -{ - std::vector> test_cases = generate_test_cases("fp8"); - - return !run_test_cases(test_cases); -}