mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-15 10:37:44 +00:00
6c7e2d3a7e
* start
* read for gridwise gemm
* add MakeBGridDescriptor_K0_N0_N1_N2_N3_K1
* add thread copy desc and register buffer
* add K0PerBlock dim
* add read global data
* finish gridwise gemm
* finish blockwise gemm
* add print data
* add smallest config
* add compare code for gridwis gemm
* fix NXdlPerWave
* fix k0perthread and gridewis gemm main loop
* remove b matrix lds alloc
* fix name
* add test code
* create b_grid_desc_k0_k1_k2_n0_n1_n2_n3_k3 from parameter
* add double register
* modify b_thread_desc_
* add float
* fp16 tag
* add tail for pipeline
* finish main loop
* optimize main loop
* start clear gridwise gemm
* clear code
* clear redundant code
* change file name
* change file name
* fix bug after merge develop
* fix input parameters
* using MultiK0 control b load data loop
* fix some config
* 4 buffer
* fix bug
* one can use
* change read order
* change buffer array to tuple
* change to 8 buffer
* interleave buffer load
* change to 16
* read 8 buffer
* add data buffer to template
* fix after merge develop(head file)
* format
* change to 4 buffer
* remove unnecessary lambda fun
[ROCm/composable_kernel commit: 10b3278b05]
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