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

added datatype and use clang-format-12

Browse Source 
(cherry picked from commit ae4e853682ef1bb27784b2f965b4a66b3751ceec)

Co-authored-by: Cenxuan <cenxuan@streamhpc.com>
This commit is contained in:
Apoorva Kalyani
2025-05-26 21:39:41 +00:00
committed by apoorva
parent 455275de80
commit 1fda4990ad
4 changed files with 15 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
include/ck/utility/dtype_vector.hpp
View File

@@ -2124,6 +2124,14 @@ using int32x16_t = typename vector_type<int32_t, 16>::type;
using int32x32_t = typename vector_type<int32_t, 32>::type;
using int32x64_t = typename vector_type<int32_t, 64>::type;
// i16
using int16x2_t = typename vector_type<int16_t, 2>::type;
using int16x4_t = typename vector_type<int16_t, 4>::type;
using int16x8_t = typename vector_type<int16_t, 8>::type;
using int16x16_t = typename vector_type<int16_t, 16>::type;
using int16x32_t = typename vector_type<int16_t, 32>::type;
using int16x64_t = typename vector_type<int16_t, 64>::type;
// i8
using int8x2_t = typename vector_type<int8_t, 2>::type;
using int8x4_t = typename vector_type<int8_t, 4>::type;

7
library/include/ck/library/tensor_operation_instance/gpu/gemm_add.hpp
View File

@@ -146,15 +146,16 @@ struct DeviceOperationInstanceFactory<
#endif
#if defined(CK_ENABLE_BF16)
// TODO:
// here for WMMA, currently BDataType and ADataType must be the same
// TODO:
// here for WMMA, currently BDataType and ADataType must be the same
if constexpr(is_same_v<ADataType, ck::bhalf_t> && is_same_v<BDataType, bhalf_t> &&
is_same_v<D0DataType, ck::bhalf_t> && is_same_v<EDataType, ck::bhalf_t>)
{
if constexpr(is_same_v<ALayout, Row> && is_same_v<BLayout, Row> &&
is_same_v<D0Layout, Row> && is_same_v<ELayout, Row>)
{
add_device_gemm_add_wmma_c_shuffle_bf16_f16_bf16_bf16_mk_kn_mn_mn_instances(op_ptrs);
add_device_gemm_add_wmma_c_shuffle_bf16_f16_bf16_bf16_mk_kn_mn_mn_instances(
op_ptrs);
}
}
#endif

3
library/src/tensor_operation_instance/gpu/gemm_add/device_gemm_add_wmma_c_shuffle_bf16_f16_bf16_bf16_mk_kn_mn_mn_instance.cpp
View File

@@ -60,7 +60,8 @@ void add_device_gemm_add_wmma_c_shuffle_bf16_f16_bf16_bf16_mk_kn_mn_mn_instances
Add>>>& instances)
{
add_device_operation_instances(
instances, device_gemm_add_wmma_c_shuffle_bf16_f16_bf16_bf16_mk_kn_mn_mn_generic_instances{});
instances,
device_gemm_add_wmma_c_shuffle_bf16_f16_bf16_bf16_mk_kn_mn_mn_generic_instances{});
add_device_operation_instances(
instances, device_gemm_add_wmma_c_shuffle_bf16_f16_bf16_bf16_mk_kn_mn_mn_instances{});
}

2
library/src/tensor_operation_instance/gpu/gemm_add/device_gemm_add_wmma_c_shuffle_f16_f16_f16_f16_mk_kn_mn_mn_instance.cpp
View File

@@ -43,7 +43,7 @@ using device_gemm_add_wmma_c_shuffle_f16_f16_f16_f16_mk_kn_mn_mn_instances = std
DeviceGemmMultipleD_Wmma_CShuffle< Row, Row, Row_Tuple, Row, F16, F16, F32, F32,F16_Tuple, F16, PassThrough, PassThrough, Add, GemmMNKPadding, 1, 256, 16, 128, 32, 8, 16, 16, 1, 2, S<4, 16, 4>, S<1, 0, 2>, S<1, 0, 2>, 2, 2, 2, 1, S<4, 64, 1>, S<0, 2, 1>, S<0, 2, 1>, 1, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1, LoopScheduler::Default, PipelineVersion::v1>,
DeviceGemmMultipleD_Wmma_CShuffle< Row, Row, Row_Tuple, Row, F16, F16, F32, F32,F16_Tuple, F16, PassThrough, PassThrough, Add, GemmMNKPadding, 1, 64, 16, 16, 32, 8, 16, 16, 1, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<0, 2, 1>, S<0, 2, 1>, 1, 1, 1, 1, 1, 1, S<1, 16, 1, 4>, 1, LoopScheduler::Default, PipelineVersion::v1>,
DeviceGemmMultipleD_Wmma_CShuffle< Row, Row, Row_Tuple, Row, F16, F16, F32, F32,F16_Tuple, F16, PassThrough, PassThrough, Add, GemmMNKPadding, 1, 64, 16, 16, 64, 8, 16, 16, 1, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<0, 2, 1>, S<0, 2, 1>, 1, 1, 1, 1, 1, 1, S<1, 16, 1, 4>, 1, LoopScheduler::Default, PipelineVersion::v1>
// clang-format on
// clang-format on
>;
void add_device_gemm_add_wmma_c_shuffle_f16_f16_f16_f16_mk_kn_mn_mn_instances(