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

Modified some code as per gemm_universal_examples

Browse Source
This commit is contained in:
apoorva
2025-05-13 09:25:54 +00:00
committed by Anton Gorenko
parent b0aa9339e9
commit c016164040
4 changed files with 27 additions and 37 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
example/01_gemm/CMakeLists.txt
View File

@@ -110,9 +110,9 @@ 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)
add_example_executable(example_gemm_wmma_bf16_v3 gemm_wmma_bf16_v3.cpp)
add_example_dependencies(example_gemm_wmma example_gemm_wmma_bf16_v3)
add_example_executable(example_gemm_wmma_f8_bf16_v3 gemm_wmma_f8_bf16_v3.cpp)
add_example_dependencies(example_gemm_wmma example_gemm_wmma_f8_bf16_v3)
add_example_executable(example_gemm_wmma_f16_v3 gemm_wmma_f16_v3.cpp)
add_example_dependencies(example_gemm_wmma example_gemm_wmma_f16_v3)

17
example/01_gemm/gemm_wmma_bf16_bf16_bf16_v3.cpp → example/01_gemm/gemm_wmma_bf16_v3.cpp
View File

@@ -5,14 +5,11 @@
#include "ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3.hpp"
using BF16 = bhalf_t;
using F32 = float;
using ADataType = BF16;
using BDataType = BF16;
using AccDataType = F32;
using CShuffleDataType = BF16;
using CDataType = BF16;
using ADataType = ck::bhalf_t;
using BDataType = ck::bhalf_t;
using AccDataType = float;
using CShuffleDataType = ck::bhalf_t;
using CDataType = ck::bhalf_t;
using ALayout = Col;
using BLayout = Row;
@@ -22,13 +19,13 @@ using AElementOp = PassThrough;
using BElementOp = PassThrough;
using CElementOp = PassThrough;
static constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecialization::GemmSpec;
static constexpr auto GemmDefault = ck::tensor_operation::device::GemmSpecialization::GemmSpec;
// clang-format off
using DeviceGemmInstance = ck::tensor_operation::device::DeviceGemm_Wmma_CShuffleV3<
ALayout, BLayout, CLayout,
ADataType, BDataType, CDataType, AccDataType, CShuffleDataType,
PassThrough, PassThrough, PassThrough, GemmSpec,
PassThrough, PassThrough, PassThrough, GemmDefault,
256,
128, 128,
32, 8, 8,

17
example/01_gemm/gemm_wmma_f16_f16_f16_v3.cpp → example/01_gemm/gemm_wmma_f16_v3.cpp
View File

@@ -5,14 +5,11 @@
#include "ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3.hpp"
using F16 = half_t;
using F32 = float;
using ADataType = F16;
using BDataType = F16;
using AccDataType = F32;
using CShuffleDataType = F16;
using CDataType = F16;
using ADataType = ck::half_t;
using BDataType = ck::half_t;
using AccDataType = float;
using CShuffleDataType = ck::half_t;
using CDataType = ck::half_t;
using ALayout = Col;
using BLayout = Row;
@@ -22,13 +19,13 @@ using AElementOp = PassThrough;
using BElementOp = PassThrough;
using CElementOp = PassThrough;
static constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecialization::GemmSpec;
static constexpr auto GemmDefault = ck::tensor_operation::device::GemmSpecialization::GemmSpec;
// clang-format off
using DeviceGemmInstance = ck::tensor_operation::device::DeviceGemm_Wmma_CShuffleV3<
ALayout, BLayout, CLayout,
ADataType, BDataType, CDataType, AccDataType, CShuffleDataType,
PassThrough, PassThrough, PassThrough, GemmSpec,
PassThrough, PassThrough, PassThrough, GemmDefault,
128,
128, 64,
64, 8, 8,

18
example/01_gemm/gemm_wmma_f8_f8_bf16_v3.cpp → example/01_gemm/gemm_wmma_f8_bf16_v3.cpp
View File

@@ -5,15 +5,11 @@
#include "ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3.hpp"
using F8 = f8_t;
using BF16 = bhalf_t;
using F32 = float;
using ADataType = F8;
using BDataType = F8;
using AccDataType = F32;
using CShuffleDataType = BF16;
using CDataType = BF16;
using ADataType = ck::f8_t;
using BDataType = ck::f8_t;
using AccDataType = float;
using CShuffleDataType = ck::bhalf_t;
using CDataType = ck::bhalf_t;
using ALayout = Col;
using BLayout = Row;
@@ -23,13 +19,13 @@ using AElementOp = PassThrough;
using BElementOp = PassThrough;
using CElementOp = PassThrough;
static constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecialization::GemmSpec;
static constexpr auto GemmDefault = ck::tensor_operation::device::GemmSpecialization::GemmSpec;
// clang-format off
using DeviceGemmInstance = ck::tensor_operation::device::DeviceGemm_Wmma_CShuffleV3<
ALayout, BLayout, CLayout,
ADataType, BDataType, CDataType, AccDataType, CShuffleDataType,
PassThrough, PassThrough, PassThrough, GemmSpec,
PassThrough, PassThrough, PassThrough, GemmDefault,
32,
16, 16,
64, 8, 8,