ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-14 19:18:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

update ck_tile script

Browse Source
This commit is contained in:
Lin, Qun
2025-05-19 21:02:23 -05:00
committed by Qun Lin
parent 2f0015e788
commit 94ae9ae73f
22 changed files with 238 additions and 146 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
example/ck_tile/03_gemm/script/benchmark_basic_bf16.sh
View File

@@ -1,12 +1,12 @@
#!/bin/sh
EXE="$(find . -name tile_example_gemm_basic -type f | head -n 1)"
VALID=1
VALID=2
for b_matrix_layout in "C"; do
for m in "64" "512" "1024" "2048"; do
for n in "512" "1024" "2048"; do
for k in "64" "512" "1024" "2048"; do
for m in "64" "512" "1024" "8192"; do
for n in "512" "1024" "8192"; do
for k in "64" "512" "1024" "8192"; do
$EXE -prec=bf16 -m=$m -n=$n -k=$k -a_layout="R" -b_layout="$b_matrix_layout" -c_layout="R" -v=$VALID
done
done

8
example/ck_tile/03_gemm/script/benchmark_basic_bf8.sh
View File

@@ -1,12 +1,12 @@
#!/bin/sh
EXE="$(find . -name tile_example_gemm_basic -type f | head -n 1)"
VALID=1
VALID=2
for b_matrix_layout in "C"; do
for m in "64" "512" "1024" "2048"; do
for n in "512" "1024" "2048"; do
for k in "64" "512" "1024" "2048"; do
for m in "64" "512" "1024" "8192"; do
for n in "512" "1024" "8192"; do
for k in "64" "512" "1024" "8192"; do
$EXE -prec=bf8 -m=$m -n=$n -k=$k -a_layout="R" -b_layout="$b_matrix_layout" -c_layout="R" -v=$VALID
done
done

8
example/ck_tile/03_gemm/script/benchmark_basic_fp16.sh
View File

@@ -1,11 +1,11 @@
#!/bin/sh
EXE="$(find . -name tile_example_gemm_basic -type f | head -n 1)"
VALID=1
VALID=2
for b_matrix_layout in "C"; do
for m in "64" "512" "1024" "2048"; do
for n in "512" "1024" "2048"; do
for k in "64" "512" "1024" "2048"; do
for m in "64" "512" "1024" "8192"; do
for n in "512" "1024" "8192"; do
for k in "64" "512" "1024" "8192"; do
$EXE -prec=fp16 -m=$m -n=$n -k=$k -a_layout="R" -b_layout="$b_matrix_layout" -c_layout="R" -v=$VALID
done
done

8
example/ck_tile/03_gemm/script/benchmark_basic_fp8.sh
View File

@@ -1,12 +1,12 @@
#!/bin/sh
EXE="$(find . -name tile_example_gemm_basic -type f | head -n 1)"
VALID=1
VALID=2
for b_matrix_layout in "C"; do
for m in "64" "512" "1024" "2048"; do
for n in "512" "1024" "2048"; do
for k in "64" "512" "1024" "2048"; do
for m in "64" "512" "1024" "8192"; do
for n in "512" "1024" "8192"; do
for k in "64" "512" "1024" "8192"; do
$EXE -prec=fp8 -m=$m -n=$n -k=$k -a_layout="R" -b_layout="$b_matrix_layout" -c_layout="R" -v=$VALID
done
done

8
example/ck_tile/03_gemm/script/benchmark_mem_pipeline_bf16.sh
View File

@@ -1,11 +1,11 @@
#!/bin/sh
EXE="$(find . -name tile_example_gemm_universal -type f | head -n 1)"
VALID=1
VALID=2
for b_matrix_layout in "C"; do
for m in "512" "1024" "2048" "4096"; do
for n in "512" "1024" "2048"; do
for k in "512" "1024" "2048"; do
for m in "512" "1024" "2048" "8192"; do
for n in "512" "1024" "8192"; do
for k in "512" "1024" "8192"; do
$EXE -prec=bf16 -m=$m -n=$n -k=$k -a_layout="R" -b_layout="$b_matrix_layout" -c_layout="R" -v=$VALID
done
done

8
example/ck_tile/03_gemm/script/benchmark_mem_pipeline_bf8.sh
View File

@@ -1,11 +1,11 @@
#!/bin/sh
EXE="$(find . -name tile_example_gemm_universal -type f | head -n 1)"
VALID=1
VALID=2
for b_matrix_layout in "C"; do
for m in "512" "1024" "2048" "4096"; do
for n in "512" "1024" "2048"; do
for k in "512" "1024" "2048"; do
for m in "512" "1024" "2048" "8192"; do
for n in "512" "1024" "8192"; do
for k in "512" "1024" "8192"; do
$EXE -prec=bf8 -m=$m -n=$n -k=$k -a_layout="R" -b_layout="$b_matrix_layout" -c_layout="R" -v=$VALID
done
done

8
example/ck_tile/03_gemm/script/benchmark_mem_pipeline_fp16.sh
View File

@@ -1,11 +1,11 @@
#!/bin/sh
EXE="$(find . -name tile_example_gemm_universal -type f | head -n 1)"
VALID=1
VALID=2
for b_matrix_layout in "C"; do
for m in "512" "1024" "2048" "4096"; do
for n in "512" "1024" "2048"; do
for k in "512" "1024" "2048"; do
for m in "512" "1024" "2048" "8192"; do
for n in "512" "1024" "8192"; do
for k in "512" "1024" "8192"; do
$EXE -prec=fp16 -m=$m -n=$n -k=$k -a_layout="R" -b_layout="$b_matrix_layout" -c_layout="R" -v=$VALID
done
done

8
example/ck_tile/03_gemm/script/benchmark_mem_pipeline_fp8.sh
View File

@@ -1,11 +1,11 @@
#!/bin/sh
EXE="$(find . -name tile_example_gemm_universal -type f | head -n 1)"
VALID=1
VALID=2
for b_matrix_layout in "C"; do
for m in "512" "1024" "2048" "4096"; do
for n in "512" "1024" "2048"; do
for k in "512" "1024" "2048"; do
for m in "512" "1024" "2048" "8192"; do
for n in "512" "1024" "8192"; do
for k in "512" "1024" "8192"; do
$EXE -prec=fp8 -m=$m -n=$n -k=$k -a_layout="R" -b_layout="$b_matrix_layout" -c_layout="R" -v=$VALID
done
done

14
example/ck_tile/04_img2col/script/smoke_test.sh Executable file
View File

@@ -0,0 +1,14 @@
#!/bin/sh
# TODO: run this script from CK root
BUILD=build
EXE=$BUILD/bin/tile_example_img2col
N=256
###### Dim G N__ K___ C___ Y X Hi__ Wi__ Strides Dilations LeftPads RightPads
$EXE 1 2 1 2 1 $N 256 64 1 1 64 64 1 1 1 1 0 0 0 0
$EXE 1 2 1 2 1 $N 512 64 1 1 64 64 1 1 1 1 0 0 0 0
$EXE 1 2 1 2 1 $N 128 64 3 3 64 64 1 1 1 1 1 1 1 1

0
example/ck_tile/06_permute/script/smoke_test.sh Normal file → Executable file
View File

0
example/ck_tile/09_topk_softmax/script/smoke_test.sh Normal file → Executable file
View File

66
example/ck_tile/10_rmsnorm2d/script/perf_test.sh
View File

@@ -1,37 +1,37 @@
#!/bin/sh
EXE="$(find . -name tile_rmsnorm2d_fwd -type f | head -n 1)"
$EXE -m=1 -n=1 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=80 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=128 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=144 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=168 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=184 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=256 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=288 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=344 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=376 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=448 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=512 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=924 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=1024 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=1078 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=1996 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=4080 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=1 -n=1 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=80 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=128 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=144 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=168 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=184 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=256 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=288 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=344 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=376 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=448 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=512 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=924 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=1024 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=1078 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=1996 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=4080 -e=1e-12 -v=1 -prec_i=bf16 -repeat=1000
$EXE -m=700 -n=80 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=128 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=144 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=168 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=184 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=256 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=288 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=344 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=376 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=448 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=512 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=924 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=1024 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=1078 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=1996 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=4080 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=80 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=128 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=144 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=168 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=184 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=256 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=288 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=344 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=376 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=448 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=512 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=924 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=1024 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=1078 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=1996 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000
$EXE -m=700 -n=4080 -e=1e-12 -v=1 -prec_i=fp16 -repeat=1000

2
example/ck_tile/11_add_rmsnorm2d_rdquant/script/perf_test.sh
View File

@@ -1,6 +1,6 @@
#!/bin/sh
EXE="$(find . -name tile_add_rmsnorm2d_rdquant_fwd -type f | head -n 1)"
echo $EXE
$EXE -m=1 -n=1 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=80 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=128 -e=1e-12 -v=1 -prec=bf16 -repeat=1000

66
example/ck_tile/12_smoothquant/script/perf_test.sh
View File

@@ -1,37 +1,37 @@
EXE="$(find . -name tile_smoothquant -type f | head -n 1)"
$EXE -m=1 -n=1 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=80 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=128 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=144 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=168 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=184 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=256 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=288 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=344 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=376 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=448 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=512 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=924 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=1024 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=1078 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=1996 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=4080 -e=1e-12 -v=1 -prec=bf16 -repeat=1000
$EXE -m=1 -n=1 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=80 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=128 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=144 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=168 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=184 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=256 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=288 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=344 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=376 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=448 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=512 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=924 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=1024 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=1078 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=1996 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=4080 -v=1 -prec=bf16 -repeat=1000
$EXE -m=700 -n=80 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=128 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=144 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=168 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=184 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=256 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=288 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=344 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=376 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=448 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=512 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=924 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=1024 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=1078 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=1996 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=4080 -e=1e-12 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=80 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=128 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=144 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=168 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=184 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=256 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=288 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=344 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=376 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=448 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=512 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=924 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=1024 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=1078 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=1996 -v=1 -prec=fp16 -repeat=1000
$EXE -m=700 -n=4080 -v=1 -prec=fp16 -repeat=1000

0
example/ck_tile/13_moe_sorting/script/smoke_test.sh Normal file → Executable file
View File

66
example/ck_tile/14_moe_smoothquant/script/perf_test.sh
View File

@@ -1,37 +1,37 @@
EXE=build/bin/tile_example_moe_smoothquant
$EXE -t=1 -h=1 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=80 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=128 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=144 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=168 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=184 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=256 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=288 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=344 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=376 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=448 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=512 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=924 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=1024 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=1078 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=1996 -v=1 -prec=bf16 -repeat=1000
$EXE -t=700 -h=4080 -v=1 -prec=bf16 -repeat=1000
$EXE -t=1 -h=1 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=80 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=128 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=144 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=168 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=184 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=256 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=288 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=344 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=376 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=448 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=512 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=924 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=1024 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=1078 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=1996 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=4080 -v=1 -prec_i=bf16 -repeat=1000
$EXE -t=700 -h=80 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=128 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=144 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=168 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=184 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=256 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=288 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=344 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=376 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=448 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=512 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=924 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=1024 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=1078 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=1996 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=4080 -v=1 -prec=fp16 -repeat=1000
$EXE -t=700 -h=80 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=128 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=144 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=168 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=184 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=256 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=288 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=344 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=376 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=448 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=512 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=924 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=1024 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=1078 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=1996 -v=1 -prec_i=fp16 -repeat=1000
$EXE -t=700 -h=4080 -v=1 -prec_i=fp16 -repeat=1000

24
example/ck_tile/16_batched_gemm/script/perf_test.sh Executable file
View File

@@ -0,0 +1,24 @@
#!/bin/sh
# TODO: run this script from CK root
BUILD=build
EXE=$BUILD/bin/tile_example_batched_gemm
$EXE -v=0 -prec=fp16 -m=960 -n=1024 -k=1024 -batch_count=8
$EXE -v=0 -prec=fp16 -m=1920 -n=2048 -k=2048 -batch_count=8
$EXE -v=0 -prec=fp16 -m=3840 -n=4096 -k=4096 -batch_count=4
$EXE -v=0 -prec=fp16 -m=7680 -n=8192 -k=8192 -batch_count=2
$EXE -v=0 -prec=fp16 -m=1024 -n=1024 -k=1024 -stride_a=1056 -stride_b=1056 -stride_c=1056 -batch_count=8
$EXE -v=0 -prec=fp16 -m=2048 -n=2048 -k=2048 -stride_a=2080 -stride_a=2080 -stride_c=2080 -batch_count=8
$EXE -v=0 -prec=fp16 -m=4096 -n=4096 -k=4096 -stride_a=4128 -stride_a=4128 -stride_c=4128 -batch_count=4
$EXE -v=0 -prec=fp16 -m=8192 -n=8192 -k=8192 -stride_a=8224 -stride_a=8224 -stride_c=8224 -batch_count=2
#$EXE -v=0 -prec=fp8 -m=960 -n=1024 -k=1024 -batch_count=8
#$EXE -v=0 -prec=fp8 -m=1920 -n=2048 -k=2048 -batch_count=8
#$EXE -v=0 -prec=fp8 -m=3840 -n=4096 -k=4096 -batch_count=4
#$EXE -v=0 -prec=fp8 -m=7680 -n=8192 -k=8192 -batch_count=2
#$EXE -v=0 -prec=fp8 -m=1024 -n=1024 -k=1024 -stride_a=1056 -stride_b=1056 -stride_c=1056 -batch_count=8
#$EXE -v=0 -prec=fp8 -m=2048 -n=2048 -k=2048 -stride_a=2080 -stride_a=2080 -stride_c=2080 -batch_count=8
#$EXE -v=0 -prec=fp8 -m=4096 -n=4096 -k=4096 -stride_a=4128 -stride_a=4128 -stride_c=4128 -batch_count=4
#$EXE -v=0 -prec=fp8 -m=8192 -n=8192 -k=8192 -stride_a=8224 -stride_a=8224 -stride_c=8224 -batch_count=2

10
example/ck_tile/17_grouped_gemm/script/perf_test.sh Executable file
View File

@@ -0,0 +1,10 @@
#!/bin/sh
# TODO: run this script from CK root
BUILD=build
EXE=$BUILD/bin/tile_example_grouped_gemm
$EXE -Ms=256,512,1024,768 -Ns=128,256,384,1024 -Ks=128,192,256,512
$EXE -Ms=512,768,4096,768 -Ns=128,768,512,2048 -Ks=128,192,256,512
$EXE -Ms=2048,2048,2048,2048 -Ns=2048,2048,2048,2048 -Ks=2048,2048,2048,2048
$EXE -validate=0 -Ms=4096,4096,4096,4096 -Ns=4096,4096,4096,4096 -Ks=4096,4096,4096,4096
$EXE -validate=0 -Ms=8192,8192,8192,8192 -Ns=8192,8192,8192,8192 -Ks=8192,8192,8192,8192

4
example/ck_tile/18_flatmm/flatmm_basic.hpp
View File

@@ -159,7 +159,9 @@ struct GemmConfig<DataType, 7>
static constexpr ck_tile::index_t K_Warp_Tile = sizeof(DataType) == 2 ? 32 : 128;
};
constexpr ck_tile::index_t GemmConfigNumVer = 8;
// fail to compile on gfx942 with full config
//constexpr ck_tile::index_t GemmConfigNumVer = 8;
constexpr ck_tile::index_t GemmConfigNumVer = 3;
template <typename ADataType>
struct GemmBasicTypeConfig;

35
example/ck_tile/18_flatmm/script/perf_test.sh Executable file
View File

@@ -0,0 +1,35 @@
#!/bin/bash
EXE="$(find . -name tile_example_flatmm_basic -type f | head -n 1)"
KNAME=1
export CK_WARMUP=0
export CK_REPEAT=1
COMMON_ARGS='-v=2 '
run_tests() {
for m in 128 1024 8192; do
for n in 128 1024 8192; do
for k in 128 1024 8192; do
$EXE -m=$m -n=$n -k=$k -stride_a=0 -stride_b=0 -stride_c=0 -prec=$1 $COMMON_ARGS
if [ $? -eq 0 ]; then
echo "Success: Test with m=$m, n=$n, k=$k executed successfully."
else
echo "Error: Test with m=$m, n=$n, k=$k failed to execute properly."
# Optionally, exit or break if you need to halt further execution
# exit 1
fi
done
done
done
}
set -x
run_tests "bf16"
run_tests "fp16"
run_tests "fp8"
run_tests "bf8"
set +x

3
example/ck_tile/18_flatmm/script/smoke_test_basic.sh
View File

@@ -30,5 +30,6 @@ set -x
run_tests "bf16"
run_tests "fp16"
run_tests "fp8"
run_tests "bf8"
set +x

30
script/run_ck_tile.sh
View File

@@ -37,12 +37,14 @@ print_log_header $layernorm2d_log $env_type $branch $host_name
./example/ck_tile/02_layernorm2d/script/perf_test.sh 2>&1 | tee -a $layernorm2d_log
#run gemm
./example/ck_tile/03_gemm/script/run_full_test.sh
export gemm_basic_log="perf_gemm_basic_$GPU_arch.log"
print_log_header $gemm_basic_log $env_type $branch $host_name
./example/ck_tile/03_gemm/script/run_full_test.sh 2>&1 | tee -a $gemm_basic_log
#run image2col
export img2col_log="perf_img2col_$GPU_arch.log"
print_log_header $layernorm2d_log $env_type $branch $host_name
./build/bin/tile_example_img2col 2>&1 | tee -a $img2col_log
print_log_header $img2col_log $env_type $branch $host_name
./example/ck_tile/04_img2col/script/smoke_test.sh 2>&1 | tee -a $img2col_log
#run reduce
export reduce_log="perf_reduce_$GPU_arch.log"
@@ -69,38 +71,42 @@ export add_rmsnorm2d_rdquant_log="perf_add_rmsnorm2d_rdquant_$GPU_arch.log"
print_log_header $add_rmsnorm2d_rdquant_log $env_type $branch $host_name
./example/ck_tile/11_add_rmsnorm2d_rdquant/script/perf_test.sh 2>&1 | tee -a $add_rmsnorm2d_rdquant_log
#run smoothquant
#run 12 smoothquant
export smoothquant_log="perf_smoothquant_$GPU_arch.log"
print_log_header $smoothquant_log $env_type $branch $host_name
./example/ck_tile/12_smoothquant/script/perf_test.sh 2>&1 | tee -a $smoothquant_log
#run moe_sorting
#run 13 moe_sorting
export moe_sorting_log="perf_moe_sorting_$GPU_arch.log"
print_log_header $moe_sorting_log $env_type $branch $host_name
./example/ck_tile/13_moe_sorting/script/smoke_test.sh 2>&1 | tee -a $moe_sorting_log
#run moe_smoothquant
#run 14 moe_smoothquant
export moe_smoothquant_log="perf_moe_smoothquant_$GPU_arch.log"
print_log_header $moe_smoothquant_log $env_type $branch $host_name
./example/ck_tile/14_moe_sorting/script/perf_test.sh 2>&1 | tee -a $moe_smoothquant_log
./example/ck_tile/14_moe_smoothquant/script/perf_test.sh 2>&1 | tee -a $moe_smoothquant_log
#run fused_moe
#run 15 fused_moe
export fused_moe_log="perf_fused_moe_$GPU_arch.log"
print_log_header $fused_moe_log $env_type $branch $host_name
#example/ck_tile/14_moe_sorting/script/perf_test.sh 2>&1 | tee -a $fused_moe_log
#run batched_gemm
export batched_gemm_log="perf_batched_gemm_$GPU_arch.log"
print_log_header $batched_gemm_log $env_type $branch $host_name
./example/ck_tile/16_batched_gemm/script/perf_test.sh 2>&1 | tee -a $batched_gemm_log
#run grouped_gemm
export grouped_gemm_log="perf_grouped_gemm_$GPU_arch.log"
print_log_header $grouped_gemm_log $env_type $branch $host_name
./example/ck_tile/17_grouped_gemm/script/perf_test.sh 2>&1 | tee -a $grouped_gemm_log
#run flatmm
export flatmm_log="perf_flatmm_$GPU_arch.log"
print_log_header $flatmm_log $env_type $branch $host_name
./example/ck_tile/18_flatmm/script/smoke_test_basic.sh 2>&1 | tee -a $flatmm_log
./example/ck_tile/18_flatmm/script/perf_test.sh 2>&1 | tee -a $flatmm_log
#run batched_transpose
export batched_transpose_log="perf_batched_transpose_$GPU_arch.log"
print_log_header $batched_transpose_log $env_type $branch $host_name
./example/ck_tile/35_batched_transpose/script/smoke_test_basic.sh 2>&1 | tee -a $batched_transpose_log
./example/ck_tile/35_batched_transpose/script/smoke_test.sh 2>&1 | tee -a $batched_transpose_log