mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-03-22 08:07:38 +00:00
715671e419
* support bf16*mxfp4 gemm * rebase bf16*fp4 example to develop branch * Clean up commented debug code in GEMM kernel * rename example folder * support bf16*mxfp4 gemm * rebase bf16*fp4 example to develop branch * Clean up commented debug code in GEMM kernel * rename example folder * rebase to new develop * fix clang format * update code according to reviewer's comment * Update README.md * update code according to reviewer's comment * update code according to reviewer's comment * Update CMakeLists.txt * Update README.md * Update CMakeLists.txt * Delete files * Delete files * Add unit tests * Update test_gemm_quant_base.hpp * merge bf16*fp4 example to develop branch * fix clang format * fix clang format * Update CMakeLists.txt * fix ci test * fix clang format * resolve conflicts --------- Co-authored-by: eliotwang <charyang@smci355-ccs-aus-m10-29.cs-aus.dcgpu> Co-authored-by: ShaoChunLee <Shao-Chun.Lee@amd.com> Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com> Co-authored-by: illsilin_amdeng <Illia.Silin@amd.com> Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>
52 lines
2.2 KiB
C++
52 lines
2.2 KiB
C++
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
|
|
#pragma once
|
|
|
|
#include <iostream>
|
|
#include <string>
|
|
|
|
#include "ck_tile/core.hpp"
|
|
|
|
namespace ck_tile {
|
|
|
|
// clang-format off
|
|
template <typename T> struct DataTypeTraits;
|
|
template <> struct DataTypeTraits<float> { static constexpr const char * name = "fp32"; };
|
|
template <> struct DataTypeTraits<double> { static constexpr const char * name = "fp64"; };
|
|
template <> struct DataTypeTraits<int32_t> { static constexpr const char * name = "int32"; };
|
|
template <> struct DataTypeTraits<fp16_t> { static constexpr const char * name = "fp16"; };
|
|
template <> struct DataTypeTraits<bf16_t> { static constexpr const char * name = "bf16"; };
|
|
template <> struct DataTypeTraits<fp8_t> { static constexpr const char * name = "fp8"; };
|
|
template <> struct DataTypeTraits<bf8_t> { static constexpr const char * name = "bf8"; };
|
|
template <> struct DataTypeTraits<int8_t> { static constexpr const char * name = "int8"; };
|
|
template <> struct DataTypeTraits<pk_int4_t> { static constexpr const char * name = "pk_int4"; };
|
|
template <> struct DataTypeTraits<pk_fp4_t> { static constexpr const char * name = "pk_fp4"; };
|
|
template <> struct DataTypeTraits<pk_fp4_raw_t> { static constexpr const char * name = "pk_fp4_raw"; };
|
|
|
|
template <memory_operation_enum MemOp> struct memOpToStr;
|
|
template <> struct memOpToStr<memory_operation_enum::set> { static constexpr const char * name = "set"; };
|
|
template <> struct memOpToStr<memory_operation_enum::atomic_add> { static constexpr const char * name = "atomic_add"; };
|
|
template <> struct memOpToStr<memory_operation_enum::atomic_max> { static constexpr const char * name = "atomic_max"; };
|
|
template <> struct memOpToStr<memory_operation_enum::add> { static constexpr const char * name = "add"; };
|
|
// clang-format on
|
|
|
|
template <typename ADataType_, typename BDataType_>
|
|
std::string gemm_prec_str()
|
|
{
|
|
std::string base_str = std::string(DataTypeTraits<ADataType_>::name);
|
|
if(!std::is_same_v<ADataType_, BDataType_>)
|
|
{
|
|
base_str += "_" + std::string(DataTypeTraits<BDataType_>::name);
|
|
}
|
|
return base_str;
|
|
}
|
|
|
|
template <memory_operation_enum MemOp_>
|
|
std::string mem_op_string()
|
|
{
|
|
return std::string(memOpToStr<MemOp_>::name);
|
|
}
|
|
|
|
} // namespace ck_tile
|