mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-12 17:26:00 +00:00

Files

Anthony Chang 7db48f9008 Tune & add conflict-free LDS gemm kernels (#159 )

* retune & add conflict-free bf16/fp16 c-shuffle gemm instances

amend wrong K1 value in some fp16/bf16 kernel instances

* make gemm cshuffle's timing behavior consistent with all other functions

* clang-format

* retune & add conflict-free fp32 c-shuffle gemm instances

* retune & add conflict-free int8 c-shuffle gemm instances

* update the underlying gridwise gemm of all c-shuffle gemm kernels

* typo

2022-03-31 12:58:41 -05:00

CMakeLists.txt

Reorganize files, Part 1 (#119 )

2022-03-08 21:46:36 -06:00

gemm_xdl_bf16.cpp

Batched gemm and reduction (#156 )

2022-03-30 11:21:18 -05:00

gemm_xdl_fp16.cpp

Tune & add conflict-free LDS gemm kernels (#159 )

2022-03-31 12:58:41 -05:00

gemm_xdl_int8.cpp

Batched gemm and reduction (#156 )

2022-03-30 11:21:18 -05:00

README.md

Compile for gfx908 and gfx90a (#130 )

2022-03-31 12:33:34 -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`