mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-12 09:16:52 +00:00

Files

Bartlomiej Wroblewski 37a8c1f756 Redesign the DPP8 GEMM kernel to use warp-wise component (#863 )

* Redesign the DPP8 GEMM kernel to use warp-wise component

* Review: Improve error messages

* Review: Remove unnecessary empty lines

* Review: Fix M, N per thread names

* Review: Rename mfma_input_type to dpp_input_type

* Review: Fix tensor adaptor; remove unnecessary element

* Review: Remove calls to dpp_gemm's MakeCDescriptor

* Review: Add blockwise doc, change function names to include dimension names

* Review: Remove duplicated code; Move Block2CtileMap alias to the top of the file

* Review: Add __restrict__ keywords

* Review: Use MatrixPadder for padding A, B, C matrices

* Review: Remove hardcoded datatypes

* Review: Change names from FloatX to XDataType

* Review: Introduce AK0 and BK0 instead of a single K0

* Review: Remove construction of dpp_datatypes object

* Review: Rename DppInstrRunner to DppLanegroupGemm

2023-09-06 11:44:09 -05:00

CMakeLists.txt

Redesign the DPP8 GEMM kernel to use warp-wise component (#863 )

2023-09-06 11:44:09 -05:00

common.hpp

initial stream-k implementation with example (#699 )

2023-07-26 14:18:15 -05:00

gemm_dl_fp16.cpp

update copyright headers (#726 )

2023-05-31 18:46:57 -05:00

gemm_dl_fp32.cpp

update copyright headers (#726 )

2023-05-31 18:46:57 -05:00

gemm_dl_int4.cpp

update copyright headers (#726 )

2023-05-31 18:46:57 -05:00

gemm_dl_int8.cpp

update copyright headers (#726 )

2023-05-31 18:46:57 -05:00

gemm_dpp_fp16.cpp

Redesign the DPP8 GEMM kernel to use warp-wise component (#863 )

2023-09-06 11:44:09 -05:00

gemm_wmma_fp16.cpp

update copyright headers (#726 )

2023-05-31 18:46:57 -05:00

gemm_xdl_bf16_rtn.cpp

add an example of customized type convert - bfp16_rtn (#869 )

2023-08-29 12:31:24 -05:00

gemm_xdl_bf16.cpp

update copyright headers (#726 )

2023-05-31 18:46:57 -05:00

gemm_xdl_f8.cpp

Add fp8 GEMM and an example for it (#767 )

2023-07-04 20:38:49 -06:00

gemm_xdl_fp16_f8.cpp

Enable f16/f8 mixed precision mode (#820 )

2023-08-09 08:44:23 -05:00

gemm_xdl_fp16.cpp

update copyright headers (#726 )

2023-05-31 18:46:57 -05:00

gemm_xdl_fp64.cpp

update copyright headers (#726 )

2023-05-31 18:46:57 -05:00

gemm_xdl_int4.cpp

update copyright headers (#726 )

2023-05-31 18:46:57 -05:00

gemm_xdl_int8.cpp

update copyright headers (#726 )

2023-05-31 18:46:57 -05:00

gemm_xdl_skip_b_lds_fp16.cpp

Disable XDL kernels on unsupported HW Add ck::is_xdl_supported (#768 )

2023-07-26 07:19:55 -07:00

gemm_xdl_streamk.cpp

initial stream-k implementation with example (#699 )

2023-07-26 14:18:15 -05:00

gemm_xdl_wavelet_fp16.cpp

Move Device Ops implementations into impl directory. (#777 )

2023-07-06 16:15:51 +02:00

README.md

Compile for gfx908 and gfx90a (#130 )

2022-03-31 12:33:34 -05:00

run_gemm_example.inc

initial stream-k implementation with example (#699 )

2023-07-26 14:18:15 -05:00

README.md

Instructions for `example_gemm_xdl`

Run `example_gemm_xdl`

#arg1: verification (0=no, 1=yes)
#arg2: initialization (0=no init, 1=integer value, 2=decimal value)
#arg3: run kernel # of times (>1)
./bin/example_gemm_xdl 0 1 5

Result (MI100 @ 1087Mhz, 133.5TFlops peak FP16)

a_m_k: dim 2, lengths {3840, 4096}, strides {4096, 1}
b_k_n: dim 2, lengths {4096, 4096}, strides {1, 4096}
c_m_n: dim 2, lengths {3840, 4096}, strides {4096, 1}
arg.a_grid_desc_k0_m_k1_{512, 3840, 8}
arg.b_grid_desc_k0_n_k1_{512, 4096, 8}
arg.c_grid_desc_m_n_{ 3840, 4096}
launch_and_time_kernel: grid_dim {480, 1, 1}, block_dim {256, 1, 1}
Warm up
Start running 5 times...
Perf: 1.19685 ms, 107.657 TFlops, 78.8501 GB/s

README.md

Instructions for example_gemm_xdl

Run example_gemm_xdl

Instructions for `example_gemm_xdl`

Run `example_gemm_xdl`