ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-18 03:49:41 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge commit 'f36cb5b2aad0acf655173290ba672066ecfa85d1' into develop

Browse Source
This commit is contained in:
assistant-librarian[bot]
2025-08-01 22:11:57 +00:00
parent 29f8d7250c
commit 0e541583d6
15 changed files with 185 additions and 735 deletions
Download Patch File Download Diff File Expand all files Collapse all files

146
Jenkinsfile vendored
View File

@@ -438,34 +438,6 @@ def cmake_build(Map conf=[:]){
echo "could not locate the requested artifacts: ${err.getMessage()}. will skip the stashing."
}
}
if (params.RUN_CK_TILE_TRANSPOSE_TESTS){
try{
archiveArtifacts "perf_transpose_*.log"
if (arch_type == 1){
stash includes: "perf_transpose_**_gfx90a.log", name: "perf_transpose_log_gfx90a"
}
else if (arch_type == 2){
stash includes: "perf_transpose_**_gfx942.log", name: "perf_transpose_log_gfx942"
}
}
catch(Exception err){
echo "could not locate the requested artifacts: ${err.getMessage()}. will skip the stashing."
}
}
if (params.RUN_CK_TILE_GEMM_TESTS){
try{
archiveArtifacts "perf_tile_gemm_**.log"
if (arch == 1){
stash includes: "perf_tile_gemm_**_gfx90a.log", name: "perf_tile_gemm_log_gfx90a"
}
else if (arch == 2){
stash includes: "perf_tile_gemm_**_gfx942.log", name: "perf_tile_gemm_log_gfx942"
}
}
catch(Exception err){
echo "could not locate the requested artifacts: ${err.getMessage()}. will skip the stashing."
}
}
}
def buildHipClangJob(Map conf=[:]){
@@ -762,24 +734,6 @@ def process_results(Map conf=[:]){
echo "could not locate the FMHA performance logs: ${err.getMessage()}."
}
}
if (params.RUN_CK_TILE_TRANSPOSE_TESTS){
try{
unstash "perf_transpose_log_gfx942"
unstash "perf_transpose_log_gfx90a"
}
catch(Exception err){
echo "could not locate the Transpose performance logs: ${err.getMessage()}."
}
}
if (params.RUN_CK_TILE_GEMM_TESTS){
try{
unstash "perf_tile_gemm_log_gfx942"
unstash "perf_tile_gemm_log_gfx90a"
}
catch(Exception err){
echo "could not locate the GEMM performance logs: ${err.getMessage()}."
}
}
if (params.RUN_FULL_QA || params.BUILD_INSTANCES_ONLY){
// unstash deb packages
unstash "packages"
@@ -861,7 +815,7 @@ def run_aiter_tests(Map conf=[:]){
}
//launch develop branch daily jobs
CRON_SETTINGS = BRANCH_NAME == "develop" ? '''0 23 * * * % RUN_FULL_QA=true;DISABLE_DL_KERNELS=true;RUN_CK_TILE_FMHA_TESTS=true;RUN_CK_TILE_TRANSPOSE_TESTS=true;RUN_CK_TILE_GEMM_TESTS=true;RUN_TILE_ENGINE_GEMM_TESTS=true;RUN_PERFORMANCE_TESTS=true;RUN_ALL_UNIT_TESTS=true
CRON_SETTINGS = BRANCH_NAME == "develop" ? '''0 23 * * * % RUN_FULL_QA=true;DISABLE_DL_KERNELS=true;RUN_CK_TILE_FMHA_TESTS=true;RUN_TILE_ENGINE_GEMM_TESTS=true;RUN_PERFORMANCE_TESTS=true;RUN_ALL_UNIT_TESTS=true
0 21 * * * % RUN_GROUPED_CONV_LARGE_CASES_TESTS=true;hipTensor_test=true;BUILD_GFX908=true;BUILD_GFX942=true;BUILD_GFX950=true;RUN_PERFORMANCE_TESTS=true;RUN_ALL_UNIT_TESTS=true
0 19 * * * % BUILD_DOCKER=true;COMPILER_VERSION=amd-staging;BUILD_COMPILER=/llvm-project/build/bin/clang++;USE_SCCACHE=false;NINJA_BUILD_TRACE=true;RUN_ALL_UNIT_TESTS=true
0 17 * * * % BUILD_DOCKER=true;COMPILER_VERSION=amd-mainline;BUILD_COMPILER=/llvm-project/build/bin/clang++;USE_SCCACHE=false;NINJA_BUILD_TRACE=true;RUN_ALL_UNIT_TESTS=true
@@ -941,14 +895,6 @@ pipeline {
name: "RUN_CK_TILE_FMHA_TESTS",
defaultValue: false,
description: "Run the ck_tile FMHA tests (default: OFF)")
booleanParam(
name: "RUN_CK_TILE_TRANSPOSE_TESTS",
defaultValue: false,
description: "Run the ck_tile Transpose tests (default: OFF)")
booleanParam(
name: "RUN_CK_TILE_GEMM_TESTS",
defaultValue: false,
description: "Run the ck_tile GEMM tests (default: OFF)")
booleanParam(
name: "RUN_TILE_ENGINE_GEMM_TESTS",
defaultValue: false,
@@ -1198,94 +1144,6 @@ pipeline {
}
}
}
stage("Run CK_TILE_TRANSPOSE Tests")
{
parallel
{
stage("Run CK_TILE_TRANSPOSE Tests on gfx90a")
{
when {
beforeAgent true
expression { params.RUN_CK_TILE_TRANSPOSE_TESTS.toBoolean() }
}
agent{ label rocmnode("gfx90a") }
environment{
setup_args = "NO_CK_BUILD"
execute_args = """ ../script/cmake-ck-dev.sh ../ gfx90a && \
make -j64 tile_example_batched_transpose && \
cd ../ &&
example/ck_tile/35_batched_transpose/script/run_full_test.sh "CI_${params.COMPILER_VERSION}" "${env.BRANCH_NAME}" "${NODE_NAME}" gfx90a """
}
steps{
buildHipClangJobAndReboot(setup_args:setup_args, no_reboot:true, build_type: 'Release', execute_cmd: execute_args)
cleanWs()
}
}
stage("Run CK_TILE_TRANSPOSE Tests on gfx942")
{
when {
beforeAgent true
expression { params.RUN_CK_TILE_TRANSPOSE_TESTS.toBoolean() }
}
agent{ label rocmnode("gfx942") }
environment{
setup_args = "NO_CK_BUILD"
execute_args = """ ../script/cmake-ck-dev.sh ../ gfx942 && \
make -j64 tile_example_batched_transpose && \
cd ../ &&
example/ck_tile/35_batched_transpose/script/run_full_test.sh "CI_${params.COMPILER_VERSION}" "${env.BRANCH_NAME}" "${NODE_NAME}" gfx942 """
}
steps{
buildHipClangJobAndReboot(setup_args:setup_args, no_reboot:true, build_type: 'Release', execute_cmd: execute_args)
cleanWs()
}
}
}
}
stage("Run CK_TILE_GEMM Tests")
{
parallel
{
stage("Run CK_TILE_GEMM Tests on gfx90a")
{
when {
beforeAgent true
expression { params.RUN_CK_TILE_GEMM_TESTS.toBoolean() }
}
agent{ label rocmnode("gfx90a") }
environment{
setup_args = "NO_CK_BUILD"
execute_args = """ ../script/cmake-ck-dev.sh ../ gfx90a && \
make -j64 tile_example_gemm_universal && \
cd ../ &&
example/ck_tile/03_gemm/script/run_full_test.sh "CI_${params.COMPILER_VERSION}" "${env.BRANCH_NAME}" "${NODE_NAME}" gfx90a """
}
steps{
buildHipClangJobAndReboot(setup_args:setup_args, no_reboot:true, build_type: 'Release', execute_cmd: execute_args)
cleanWs()
}
}
stage("Run CK_TILE_GEMM Tests on gfx942")
{
when {
beforeAgent true
expression { params.RUN_CK_TILE_GEMM_TESTS.toBoolean() }
}
agent{ label rocmnode("gfx942") }
environment{
setup_args = "NO_CK_BUILD"
execute_args = """ ../script/cmake-ck-dev.sh ../ gfx942 && \
make -j64 tile_example_gemm_universal && \
cd ../ &&
example/ck_tile/03_gemm/script/run_full_test.sh "CI_${params.COMPILER_VERSION}" "${env.BRANCH_NAME}" "${NODE_NAME}" gfx942 """
}
steps{
buildHipClangJobAndReboot(setup_args:setup_args, no_reboot:true, build_type: 'Release', execute_cmd: execute_args)
cleanWs()
}
}
}
}
stage("Run TILE_ENGINE_GEMM Tests")
{
parallel
@@ -1492,7 +1350,7 @@ pipeline {
-DGPU_TARGETS="gfx90a" \
-DCMAKE_CXX_COMPILER="${build_compiler()}" \
-DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang \
-DCMAKE_CXX_FLAGS=" -O3 " .. && make -j """
-DCMAKE_CXX_FLAGS=" -O3 " .. && make -j 32"""
}
steps{
Build_CK_and_Reboot(setup_args: setup_args, config_targets: "install", no_reboot:true, build_type: 'Release', execute_cmd: execute_args, prefixpath: '/usr/local')

2
include/ck/tensor_operation/gpu/device/impl/device_sparse_embeddings_forward_layernorm.hpp
View File

@@ -12,7 +12,7 @@
#include "ck/utility/common_header.hpp"
#include "ck/tensor_description/tensor_descriptor.hpp"
#include "ck/tensor_description/tensor_descriptor_helper.hpp"
#if __clang_major__ == 20
#if __clang_major__ >= 20
#include "ck/tensor_operation/gpu/grid/gridwise_sparse_embeddings_forward_layernorm_builtins.hpp"
#else
#include "ck/tensor_operation/gpu/grid/gridwise_sparse_embeddings_forward_layernorm.hpp"

2
include/ck/utility/common_header.hpp
View File

@@ -33,7 +33,7 @@
#include "ck/utility/thread_group.hpp"
#include "ck/utility/debug.hpp"
#if __clang_major__ == 20
#if __clang_major__ >= 20
#include "amd_buffer_addressing_builtins.hpp"
#else
#include "amd_buffer_addressing.hpp"

2
include/ck/utility/dynamic_buffer.hpp
View File

@@ -7,7 +7,7 @@
#include "ck/utility/data_type.hpp"
#include "enable_if.hpp"
#include "c_style_pointer_cast.hpp"
#if __clang_major__ == 20
#if __clang_major__ >= 20
#include "amd_buffer_addressing_builtins.hpp"
#else
#include "amd_buffer_addressing.hpp"

2
include/ck_tile/core/config.hpp
View File

@@ -253,7 +253,7 @@
#endif
#ifndef CK_TILE_USE_BUFFER_ADDRESSING_BUILTIN
#if __clang_major__ == 20
#if __clang_major__ >= 20
#define CK_TILE_USE_BUFFER_ADDRESSING_BUILTIN 1
#else
#define CK_TILE_USE_BUFFER_ADDRESSING_BUILTIN 0

2
include/ck_tile/core/tensor/buffer_view.hpp
View File

@@ -5,7 +5,7 @@
#include "ck_tile/core/config.hpp"
#include "ck_tile/core/arch/arch.hpp"
#if __clang_major__ == 20
#if __clang_major__ >= 20
#include "ck_tile/core/arch/amd_buffer_addressing_builtins.hpp"
#else
#include "ck_tile/core/arch/amd_buffer_addressing.hpp"

4
test/CMakeLists.txt
View File

@@ -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)

30
test/ck_tile/batched_transpose/CMakeLists.txt
View File

@@ -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()

25
test/ck_tile/batched_transpose/batched_transpose.hpp
View File

@@ -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 <vector>
#include <string>
#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);

283
test/ck_tile/batched_transpose/batched_transpose.inc
View File

@@ -1,283 +0,0 @@
// Copyright © Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include <vector>
#include <iostream>
#include <numeric>
#include <cassert>
#include <cstdlib>
#include <iostream>
#include <time.h>
#include <unordered_set>
#include "batched_transpose.hpp"
// different threshold for different dtype
template <typename DataType>
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<ck_tile::bf16_t>(std::string /*init_method*/)
{
double rtol = 1e-2;
double atol = 1e-2;
return ck_tile::make_tuple(rtol, atol);
}
template <>
auto get_elimit<ck_tile::fp8_t>(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 <typename Type>
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<Type> 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<Type> 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<Type>{-.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<float>(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<Type> 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<Type>(x_host, y_ref, layout_in, layout_out);
auto [rtol, atol] = get_elimit<Type>("");
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 <typename PrecType>
bool run_test_case(int argc, char** argv)
{
auto [result, args] = create_args(argc, argv);
if(!result)
return false;
return run_batched_transpose<PrecType>(args);
}
template <typename PrecType>
bool run_test_cases(std::vector<std::vector<std::string>>& 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<PrecType>(num_args, argv);
if(!valid)
break;
}
return valid;
}
std::vector<std::vector<std::string>> 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"}};
}

109
test/ck_tile/batched_transpose/batched_transpose_api.cpp
View File

@@ -1,109 +0,0 @@
// Copyright © Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include "batched_transpose.hpp"
template <typename ts_type,
ck_tile::index_t block_x,
ck_tile::index_t block_y,
ck_tile::index_t warp_x,
ck_tile::index_t warp_y,
bool kPadM,
bool kPadN>
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<block_x, block_y>;
using warp_layout = ck_tile::sequence<warp_x, warp_y>;
using ts_problem =
ck_tile::BatchedTransposeProblem<ts_type, block_tile, warp_layout, kPadM, kPadN>;
using ts_pipeline = ck_tile::BatchedTransposePipeline<ts_problem>;
using kernel = ck_tile::BatchedTransposeKernel<ts_pipeline>;
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<blocks.x, 1>(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<REAL_TYPE, BX, BY, WX, WY, PADM, PADN>(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;
}

10
test/ck_tile/batched_transpose/batched_transpose_bf16.cpp
View File

@@ -1,10 +0,0 @@
// Copyright © Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include "batched_transpose.inc"
int main()
{
std::vector<std::vector<std::string>> test_cases = generate_test_cases("bf16");
return !run_test_cases<ck_tile::bf16_t>(test_cases);
}

10
test/ck_tile/batched_transpose/batched_transpose_fp16.cpp
View File

@@ -1,10 +0,0 @@
// Copyright © Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include "batched_transpose.inc"
int main()
{
std::vector<std::vector<std::string>> test_cases = generate_test_cases("fp16");
return !run_test_cases<ck_tile::fp16_t>(test_cases);
}

10
test/ck_tile/batched_transpose/batched_transpose_fp8.cpp
View File

@@ -1,10 +0,0 @@
// Copyright © Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include "batched_transpose.inc"
int main()
{
std::vector<std::vector<std::string>> test_cases = generate_test_cases("fp8");
return !run_test_cases<ck_tile::fp8_t>(test_cases);
}

283
tile_engine/ops/gemm/CMakeLists.txt
View File

@@ -1,148 +1,215 @@
set(GEMM_DATATYPE "fp8;fp16" CACHE STRING "List of datatypes for GEMM (semicolon-separated)")
set(GEMM_LAYOUT "rcr" CACHE STRING "List of layout for GEMM (semicolon-separated)")
# Pre-generate all kernel lists to avoid blocking during parallel builds
foreach(dt IN LISTS GEMM_DATATYPE)
foreach(l IN LISTS GEMM_LAYOUT)
set(working_path "${CMAKE_CURRENT_BINARY_DIR}/${dt}/${l}")
file(MAKE_DIRECTORY "${working_path}")
if (l STREQUAL "rcr")
set(json_blob "${CMAKE_CURRENT_LIST_DIR}/configs/default_config.json")
else()
set(json_blob "${CMAKE_CURRENT_LIST_DIR}/configs/custom_ci_config.json")
endif()
# Only run if files don't exist
if (NOT EXISTS "${working_path}/gemm_instance_blobs.txt")
execute_process(
COMMAND ${Python3_EXECUTABLE} "${CMAKE_CURRENT_LIST_DIR}/gemm_instance_builder.py"
--working_path "${working_path}"
--datatype "${dt}"
--layout "${l}"
--config_json "${json_blob}"
--list_blobs
RESULT_VARIABLE ret
)
if (NOT ret EQUAL 0)
message(FATAL_ERROR "Failed to pre-generate kernel list for ${dt} ${l}")
endif()
endif()
endforeach()
endforeach()
function(build_gemm_for_datatype datatype layout)
set(working_path "${CMAKE_CURRENT_BINARY_DIR}/${datatype}/${layout}")
# Comment this if-else block when using user_provided_config
if(layout STREQUAL "rcr")
if (layout STREQUAL "rcr")
set(json_blob "${CMAKE_CURRENT_LIST_DIR}/configs/default_config.json")
else()
set(json_blob "${CMAKE_CURRENT_LIST_DIR}/configs/custom_ci_config.json")
endif()
# uncomment this if you want to use user_provided_config.json
# Uncomment to override:
# set(json_blob "${CMAKE_CURRENT_LIST_DIR}/configs/user_provided_config.json")
# Generate kernel list
execute_process(
COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/gemm_instance_builder.py
--working_path ${working_path}
--datatype ${datatype}
--layout ${layout}
--config_json ${json_blob}
--list_blobs
RESULT_VARIABLE ret
)
if(NOT ret EQUAL 0)
message(FATAL_ERROR "Failed to list kernels for ${datatype} ${layout}: ${ret}")
endif()
# Read pre-generated kernel lists
file(STRINGS "${working_path}/gemm_instance_blobs.txt" codegen_blobs)
file(STRINGS "${working_path}/gemm_instance_blobs_range.txt" codegen_blobs_range)
# Generate the blobs
add_custom_command(
OUTPUT ${codegen_blobs}
COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/gemm_instance_builder.py
COMMAND ${Python3_EXECUTABLE} "${CMAKE_CURRENT_LIST_DIR}/gemm_instance_builder.py"
--working_path "${working_path}"
--datatype ${datatype}
--layout ${layout}
--datatype "${datatype}"
--layout "${layout}"
--config_json "${json_blob}"
--gen_blobs
COMMENT "Generating GEMM instance sources for ${datatype} ${layout}"
)
add_custom_target(gemm_gen_${datatype}_${layout} DEPENDS ${codegen_blobs})
set(intermediate_libs)
list(LENGTH codegen_blobs codegen_blobs_len)
# Parse ranges to identify unique trait names
set(unique_traits)
foreach(range_line IN LISTS codegen_blobs_range)
string(STRIP "${range_line}" stripped_line)
separate_arguments(split_line UNIX_COMMAND "${stripped_line}")
list(GET split_line 0 trait_name)
list(APPEND unique_traits "${trait_name}")
endforeach()
list(REMOVE_DUPLICATES unique_traits)
foreach(blob IN LISTS codegen_blobs_range)
string(STRIP "${blob}" stripped_blob)
separate_arguments(spilit_blob UNIX_COMMAND "${stripped_blob}")
# Each line is: <trait_name> <first_index_inclusive> <last_index_exclusive>
list(GET spilit_blob 0 name)
list(GET spilit_blob 1 first)
list(GET spilit_blob 2 last)
math(EXPR total_files "${last} - ${first}")
if(total_files EQUAL 0)
continue() # nothing for this trait
endif()
# Build each trait separately
foreach(trait IN LISTS unique_traits)
set(trait_files)
foreach(range_line IN LISTS codegen_blobs_range)
string(STRIP "${range_line}" stripped_line)
separate_arguments(split_line UNIX_COMMAND "${stripped_line}")
list(GET split_line 0 name)
if (name STREQUAL trait)
list(GET split_line 1 first)
list(GET split_line 2 last)
math(EXPR total_files "${last} - ${first}")
if (total_files GREATER 0)
foreach(j RANGE ${first} ${last}-1)
list(LENGTH codegen_blobs blobs_len)
if (j LESS blobs_len)
list(GET codegen_blobs ${j} f)
list(APPEND trait_files "${f}")
endif()
endforeach()
endif()
endif()
endforeach()
# Object libraries (chunked) per trait
set(sub_intermediate_libs)
set(chunk_size 3)
math(EXPR num_chunks "( ${total_files} + ${chunk_size} - 1 ) / ${chunk_size}")
math(EXPR num_chunks_minus_1 "${num_chunks} - 1")
foreach(i RANGE 0 ${num_chunks_minus_1})
math(EXPR start "${first} + ${i} * ${chunk_size} ")
math(EXPR end "${start} + ${chunk_size} - 1")
if (trait_files)
# Create object libraries with chunking
set(chunk_size 3) # adjust as needed for memory vs parallelism
list(LENGTH trait_files num_files)
math(EXPR num_chunks "( ${num_files} + ${chunk_size} - 1 ) / ${chunk_size}")
set(chunk_files)
foreach(j RANGE ${start} ${end})
if(j LESS ${last} AND j LESS ${codegen_blobs_len})
list(GET codegen_blobs ${j} f)
list(APPEND chunk_files "${f}")
set(trait_obj_libs)
foreach(i RANGE 0 ${num_chunks}-1)
math(EXPR start "${i} * ${chunk_size}")
math(EXPR end "${start} + ${chunk_size} - 1")
set(chunk_files)
foreach(j RANGE ${start} ${end})
if (j LESS ${num_files})
list(GET trait_files ${j} f)
list(APPEND chunk_files "${f}")
endif()
endforeach()
if (chunk_files)
set(obj_lib_name "gemm_obj_${trait}_${i}_${datatype}_${layout}")
add_library(${obj_lib_name} OBJECT ${chunk_files})
add_dependencies(${obj_lib_name} gemm_gen_${datatype}_${layout})
target_compile_options(${obj_lib_name} PRIVATE
-Wno-undefined-func-template
-Wno-float-equal
--offload-compress
-O3
-fno-exceptions
)
set_target_properties(${obj_lib_name} PROPERTIES
UNITY_BUILD ON
UNITY_BUILD_BATCH_SIZE 2
)
list(APPEND trait_obj_libs "${obj_lib_name}")
endif()
endforeach()
#list(LENGTH chunk_files chunk_files_len)
#if(chunk_files_len AND chunk_files_len GREATER 1)
if(chunk_files)
set(sub_intermediate_lib_name "gemm_objlib_${name}_${i}_${datatype}_${layout}")
add_library(${sub_intermediate_lib_name} OBJECT ${chunk_files})
list(APPEND sub_intermediate_libs ${sub_intermediate_lib_name})
# Static library for this trait
if (trait_obj_libs)
set(trait_lib_name "gemm_lib_${trait}_${datatype}_${layout}")
set(obj_exprs)
foreach(objlib IN LISTS trait_obj_libs)
list(APPEND obj_exprs "$<TARGET_OBJECTS:${objlib}>")
endforeach()
add_library(${trait_lib_name} STATIC ${obj_exprs})
add_dependencies(${trait_lib_name} gemm_gen_${datatype}_${layout})
# Trait-specific executable
set(exec_name "benchmark_gemm_${datatype}_${layout}_${trait}")
add_executable(${exec_name} benchmark_gemm.cpp)
target_link_libraries(${exec_name} PRIVATE ${trait_lib_name})
target_include_directories(${exec_name} PRIVATE
"${CMAKE_CURRENT_LIST_DIR}"
"${working_path}"
)
target_compile_definitions(${exec_name} PRIVATE
GEMM_TRAIT_FILTER="${trait}"
)
target_compile_options(${exec_name} PRIVATE
-Wno-undefined-func-template
-Wno-float-equal
--offload-compress
)
endif()
endforeach()
# ------------------ Bundle the object libs into one static lib ---------
#list(LENGTH sub_intermediate_libs sub_intermediate_libs_len)
#if(sub_intermediate_libs AND sub_intermediate_libs_len GREATER 1)
if(sub_intermediate_libs)
set(intermediate_lib_name "gemm_staticlib_${name}_${datatype}_${layout}")
# Collect the $<TARGET_OBJECTS:...> expressions
set(obj_exprs)
foreach(objlib IN LISTS sub_intermediate_libs)
list(APPEND obj_exprs $<TARGET_OBJECTS:${objlib}>)
endforeach()
add_library(${intermediate_lib_name} STATIC ${obj_exprs})
add_dependencies(${intermediate_lib_name} gemm_gen_${datatype}_${layout})
#foreach(objlib IN LISTS sub_intermediate_libs)
# target_sources(${intermediate_lib_name} PRIVATE $<TARGET_OBJECTS:${objlib}>)
#endforeach()
list(APPEND intermediate_libs ${intermediate_lib_name})
endif()
endforeach()
# Interface library for instances
add_library(gemm_template_instances_${datatype}_${layout} INTERFACE)
add_dependencies(gemm_template_instances_${datatype}_${layout} gemm_gen_${datatype}_${layout})
target_link_libraries(gemm_template_instances_${datatype}_${layout} INTERFACE ${intermediate_libs})
target_include_directories(gemm_template_instances_${datatype}_${layout} INTERFACE
${CMAKE_CURRENT_LIST_DIR}
"${working_path}"
)
set_target_properties(gemm_template_instances_${datatype}_${layout} PROPERTIES LINKER_LANGUAGE CXX)
# Host API interface library
add_library(gemm_host_api_${datatype}_${layout} INTERFACE)
target_link_libraries(gemm_host_api_${datatype}_${layout} INTERFACE gemm_template_instances_${datatype}_${layout})
target_include_directories(gemm_host_api_${datatype}_${layout} INTERFACE
${CMAKE_CURRENT_LIST_DIR}
"${working_path}"
)
# Executable per datatype
set(exec_name "benchmark_gemm_${datatype}_${layout}")
add_executable(${exec_name} benchmark_gemm.cpp)
target_link_libraries(${exec_name} PRIVATE gemm_host_api_${datatype}_${layout})
target_compile_options(${exec_name} PRIVATE
-Wno-undefined-func-template
-Wno-float-equal
--offload-compress
)
# Master executable including all traits
set(all_trait_libs)
foreach(trait IN LISTS unique_traits)
if (TARGET gemm_lib_${trait}_${datatype}_${layout})
list(APPEND all_trait_libs "gemm_lib_${trait}_${datatype}_${layout}")
endif()
endforeach()
if (all_trait_libs)
add_executable(benchmark_gemm_${datatype}_${layout} benchmark_gemm.cpp)
target_link_libraries(benchmark_gemm_${datatype}_${layout} PRIVATE ${all_trait_libs})
target_include_directories(benchmark_gemm_${datatype}_${layout} PRIVATE
"${CMAKE_CURRENT_LIST_DIR}"
"${working_path}"
)
target_compile_options(benchmark_gemm_${datatype}_${layout} PRIVATE
-Wno-undefined-func-template
-Wno-float-equal
--offload-compress
)
endif()
endfunction()
# Process each datatype in isolation
# Process each datatype/layout
foreach(dt IN LISTS GEMM_DATATYPE)
foreach(l IN LISTS GEMM_LAYOUT)
build_gemm_for_datatype(${dt} ${l})
build_gemm_for_datatype("${dt}" "${l}")
endforeach()
endforeach()
# Master target for parallel builds
set(ALL_GEMM_TARGETS)
foreach(dt IN LISTS GEMM_DATATYPE)
foreach(l IN LISTS GEMM_LAYOUT)
list(APPEND ALL_GEMM_TARGETS "benchmark_gemm_${dt}_${l}")
endforeach()
endforeach()
add_custom_target(benchmark_gemm_all DEPENDS ${ALL_GEMM_TARGETS})
# Use faster linker if available
find_program(LLD_LINKER "ld.lld")
find_program(MOLD_LINKER "mold")
if (MOLD_LINKER)
message(STATUS "Using mold linker for faster linking")
add_link_options(-fuse-ld=mold)
elseif (LLD_LINKER)
message(STATUS "Using lld linker for faster linking")
add_link_options(-fuse-ld=lld)
endif()