mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 18:17:44 +00:00
d5645ff481
* 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>
[ROCm/composable_kernel commit: 715671e419]
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
|