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

Fixed review comments and added executable to cmakeLists

Browse Source
This commit is contained in:
apoorva
2025-05-08 21:39:14 +00:00
committed by Anton Gorenko
parent 8f8e631355
commit 840b79dfa3
3 changed files with 34 additions and 25 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
example/01_gemm/CMakeLists.txt
View File

@@ -109,3 +109,10 @@ add_example_executable(example_gemm_wmma_bf16 gemm_wmma_bf16.cpp)
add_example_dependencies(example_gemm_wmma example_gemm_wmma_bf16)
add_example_executable(example_gemm_wmma_int8 gemm_wmma_int8.cpp)
add_example_dependencies(example_gemm_wmma example_gemm_wmma_int8)
add_example_executable(example_gemm_wmma_bf16_bf16_bf16_v3 gemm_wmma_bf16_bf16_bf16_v3.cpp)
add_example_dependencies(example_gemm_wmma example_gemm_wmma_bf16_bf16_bf16_v3)
add_example_executable(example_gemm_wmma_f8_f8_bf16_v3 gemm_wmma_f8_f8_bf16_v3.cpp)
add_example_dependencies(example_gemm_wmma example_gemm_wmma_f8_f8_bf16_v3)
add_example_executable(example_gemm_wmma_f16_f16_f16_v3 gemm_wmma_f16_f16_f16_v3.cpp)
add_example_dependencies(example_gemm_wmma example_gemm_wmma_f16_f16_f16_v3)

31
example/01_gemm/gemm_wmma_f16_f16_f16_v3.cpp
View File

@@ -5,22 +5,21 @@
#include "ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3.hpp.hpp"
using F8 = f8_t;
using BF16 = bhalf_t;
using F32 = float;
using F16 = half_t;
using F32 = float;
using Row = tensor_layout::gemm::RowMajor;
using Col = tensor_layout::gemm::ColumnMajor;
using PassThrough = element_wise::PassThrough;
using ADataType = F8;
using BDataType = F8;
using ADataType = F16;
using BDataType = F16;
using AccDataType = F32;
using CShuffleDataType = BF16;
using CDataType = BF16;
using CShuffleDataType = F16;
using CDataType = F16;
using ALayout = Row;
using ALayout = Col;
using BLayout = Row;
using CLayout = Row;
@@ -35,16 +34,16 @@ using DeviceGemmInstance = ck::tensor_operation::device::DeviceGemm_Wmma_CShuffl
ALayout, BLayout, CLayout,
ADataType, BDataType, CDataType, AccDataType, CShuffleDataType,
PassThrough, PassThrough, PassThrough, GemmSpec,
32,
16, 16,
128,
128, 64,
64, 8, 8,
16, 16,
1, 1,
S<2, 16, 1>, S<1, 0, 2>, S<1, 0, 2>,
2, 8, 8, 1,
S<2, 16, 1>, S<0, 2, 1>, S<0, 2, 1>,
1, 1, 4, 0, 1, 1,
S<1, 16, 1, 2>, 8,
4, 2,
S<4, 32, 1>, S<0, 2, 1>, S<0, 2, 1>,
1, 1, 8, 1,
S<4, 32, 1>, S<0, 2, 1>, S<0, 2, 1>,
1, 1, 8, 1, 1, 1,
S<1, 32, 1, 4>, 8,
ck::BlockGemmPipelineScheduler::Intrawave,ck::BlockGemmPipelineVersion::v3>;
// clang-format on

21
example/01_gemm/gemm_wmma_f8_f8_bf16_v3.cpp
View File

@@ -5,6 +5,7 @@
#include "ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3.hpp.hpp"
using F8 = f8_t;
using BF16 = bhalf_t;
using F32 = float;
@@ -13,8 +14,8 @@ using Col = tensor_layout::gemm::ColumnMajor;
using PassThrough = element_wise::PassThrough;
using ADataType = BF16;
using BDataType = BF16;
using ADataType = F8;
using BDataType = F8;
using AccDataType = F32;
using CShuffleDataType = BF16;
using CDataType = BF16;
@@ -34,14 +35,16 @@ using DeviceGemmInstance = ck::tensor_operation::device::DeviceGemm_Wmma_CShuffl
ALayout, BLayout, CLayout,
ADataType, BDataType, CDataType, AccDataType, CShuffleDataType,
PassThrough, PassThrough, PassThrough, GemmSpec,
256,
128, 128,
32, 8, 8,
32,
16, 16,
64, 8, 8,
16, 16,
4, 2,
S<4, 64, 1>, S<0, 2, 1>, S<0, 2, 1>,
1, 1, 8, 1, 1, 1,
S<1, 32, 1, 8>, 8,
1, 1,
S<2, 16, 1>, S<1, 0, 2>, S<1, 0, 2>,
2, 8, 8, 1,
S<2, 16, 1>, S<0, 2, 1>, S<0, 2, 1>,
1, 1, 4, 0, 1, 1,
S<1, 16, 1, 2>, 8,
ck::BlockGemmPipelineScheduler::Intrawave,ck::BlockGemmPipelineVersion::v3>;
// clang-format on