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

Cleaned up the code.

Browse Source
This commit is contained in:
apoorva
2025-07-01 15:06:32 +00:00
parent 30378587a6
commit feca919c25
5 changed files with 15 additions and 55 deletions
Download Patch File Download Diff File Expand all files Collapse all files

21
example/68_gemm_add/CMakeLists.txt
View File

@@ -1,18 +1,19 @@
add_custom_target(example_gemm_add_xdl)
add_library(example_gemm_add_xdl_fp16 gemm_add_xdl_fp16.cpp)
add_example_executable(example_gemm_add_xdl_fp16 gemm_add_xdl_fp16.cpp)
add_library(example_gemm_add_xdl_bf16 gemm_add_xdl_bf16.cpp)
add_example_executable(example_gemm_add_xdl_bf16 gemm_add_xdl_bf16.cpp)
add_custom_target(example_gemm_add_wmma)
add_example_executable(example_gemm_add_wmma_bf16 gemm_add_wmma_bf16.cpp)
add_example_executable(example_gemm_add_wmma_fp16 gemm_add_wmma_fp16.cpp)
add_example_executable(example_gemm_add_wmma_v3_fp16 gemm_add_wmma_v3_fp16.cpp)
add_example_executable(example_gemm_add_wmma_v3_bf16 gemm_add_wmma_v3_bf16.cpp)
add_custom_target(example_gemm_add_xdl)
set_source_files_properties(example_gemm_add_xdl_fp16/gemm_add_xdl_fp16.cpp PROPERTIES COMPILE_OPTIONS ";-mllvm;-greedy-reverse-local-assignment=1")
add_library(example_gemm_add_xdl_fp16 gemm_add_xdl_fp16.cpp)
add_example_executable(example_gemm_add_xdl_fp16 gemm_add_xdl_fp16.cpp)
add_example_dependencies(example_gemm_add_xdl example_gemm_add_xdl_fp16)
set_source_files_properties(example_gemm_add_xdl_bf16/gemm_add_xdl_bf16.cpp PROPERTIES COMPILE_OPTIONS ";-mllvm;-greedy-reverse-local-assignment=1")
add_library(example_gemm_add_xdl_bf16 gemm_add_xdl_bf16.cpp)
add_example_executable(example_gemm_add_xdl_bf16 gemm_add_xdl_bf16.cpp)
add_example_dependencies(example_gemm_add_xdl example_gemm_add_xdl_bf16)

10
example/68_gemm_add/common.hpp
View File

@@ -13,17 +13,11 @@
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "ck/tensor_operation/gpu/device/gemm_specialization.hpp"
#ifndef CK_USE_XDL
#include "ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle.hpp"
#endif
#ifndef CK_USE_MULTIPLE_D_WMMA
#include "ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_wmma_cshuffle.hpp"
#endif
#ifndef CK_USE_WMMA_V3
#include "ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_wmma_cshuffle_v3.hpp"
#endif
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/utility/data_type.hpp"
@@ -48,6 +42,10 @@ using BF16 = ck::bhalf_t;
using F16 = ck::half_t;
using F32 = float;
using Row_Tuple = ck::Tuple<Row>;
using F16_Tuple = ck::Tuple<F16>;
using BF16_Tuple = ck::Tuple<BF16>;
struct ProblemSize final
{
ck::index_t M = 3840;

11
example/68_gemm_add/gemm_add_wmma_fp16.cpp
View File

@@ -3,17 +3,6 @@
#include "common.hpp"
template <ck::index_t... Is>
using S = ck::Sequence<Is...>;
using F16 = ck::half_t;
using F32 = float;
using Row = ck::tensor_layout::gemm::RowMajor;
using Col = ck::tensor_layout::gemm::ColumnMajor;
using PassThrough = ck::tensor_operation::element_wise::PassThrough;
using ADataType = F16;
using BDataType = F16;
using AccDataType = F32;

13
example/68_gemm_add/gemm_add_wmma_v3_bf16.cpp
View File

@@ -3,19 +3,6 @@
#include "common.hpp"
template <ck::index_t... Is>
using S = ck::Sequence<Is...>;
using BF16 = ck::bhalf_t;
using F32 = float;
using Row = ck::tensor_layout::gemm::RowMajor;
using Col = ck::tensor_layout::gemm::ColumnMajor;
using PassThrough = ck::tensor_operation::element_wise::PassThrough;
using BF16_Tuple = ck::Tuple<BF16>;
using ADataType = BF16;
using BDataType = BF16;
using AccDataType = F32;

15
example/68_gemm_add/gemm_add_wmma_v3_fp16.cpp
View File

@@ -3,19 +3,6 @@
#include "common.hpp"
template <ck::index_t... Is>
using S = ck::Sequence<Is...>;
using F16 = ck::half_t;
using F32 = float;
using Row = ck::tensor_layout::gemm::RowMajor;
using Col = ck::tensor_layout::gemm::ColumnMajor;
using PassThrough = ck::tensor_operation::element_wise::PassThrough;
using F16_Tuple = ck::Tuple<F16>;
using ADataType = F16;
using BDataType = F16;
using AccDataType = F32;
@@ -24,8 +11,6 @@ using DDataType = F16;
using DsDataType = F16_Tuple;
using EDataType = F16;
using Row_Tuple = ck::Tuple<Row>;
using ALayout = Row;
using BLayout = Row;
using DLayout = Row;