diff --git a/example/ck_tile/03_gemm/script/benchmark_basic_bf16.sh b/example/ck_tile/03_gemm/script/benchmark_basic_bf16.sh index d7e5d4640a..70ad3a0478 100755 --- a/example/ck_tile/03_gemm/script/benchmark_basic_bf16.sh +++ b/example/ck_tile/03_gemm/script/benchmark_basic_bf16.sh @@ -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 diff --git a/example/ck_tile/03_gemm/script/benchmark_basic_bf8.sh b/example/ck_tile/03_gemm/script/benchmark_basic_bf8.sh index 466f6bb4e1..40a602fd2b 100755 --- a/example/ck_tile/03_gemm/script/benchmark_basic_bf8.sh +++ b/example/ck_tile/03_gemm/script/benchmark_basic_bf8.sh @@ -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 diff --git a/example/ck_tile/03_gemm/script/benchmark_basic_fp16.sh b/example/ck_tile/03_gemm/script/benchmark_basic_fp16.sh index 64d2ddbb5c..0b64f8236e 100755 --- a/example/ck_tile/03_gemm/script/benchmark_basic_fp16.sh +++ b/example/ck_tile/03_gemm/script/benchmark_basic_fp16.sh @@ -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 diff --git a/example/ck_tile/03_gemm/script/benchmark_basic_fp8.sh b/example/ck_tile/03_gemm/script/benchmark_basic_fp8.sh index 21462616be..8baa319001 100755 --- a/example/ck_tile/03_gemm/script/benchmark_basic_fp8.sh +++ b/example/ck_tile/03_gemm/script/benchmark_basic_fp8.sh @@ -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 diff --git a/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_bf16.sh b/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_bf16.sh index 903b4a3c0f..619c6254f0 100755 --- a/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_bf16.sh +++ b/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_bf16.sh @@ -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 diff --git a/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_bf8.sh b/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_bf8.sh index 8c92c2e991..ef78b63241 100755 --- a/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_bf8.sh +++ b/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_bf8.sh @@ -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 diff --git a/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_fp16.sh b/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_fp16.sh index c4cf4ddcbf..79b4de2424 100755 --- a/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_fp16.sh +++ b/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_fp16.sh @@ -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 diff --git a/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_fp8.sh b/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_fp8.sh index e238006c7d..7ab68ff97b 100755 --- a/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_fp8.sh +++ b/example/ck_tile/03_gemm/script/benchmark_mem_pipeline_fp8.sh @@ -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 diff --git a/example/ck_tile/04_img2col/script/smoke_test.sh b/example/ck_tile/04_img2col/script/smoke_test.sh new file mode 100755 index 0000000000..4d992fc48a --- /dev/null +++ b/example/ck_tile/04_img2col/script/smoke_test.sh @@ -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 + + + diff --git a/example/ck_tile/06_permute/script/smoke_test.sh b/example/ck_tile/06_permute/script/smoke_test.sh old mode 100644 new mode 100755 diff --git a/example/ck_tile/09_topk_softmax/script/smoke_test.sh b/example/ck_tile/09_topk_softmax/script/smoke_test.sh old mode 100644 new mode 100755 diff --git a/example/ck_tile/10_rmsnorm2d/script/perf_test.sh b/example/ck_tile/10_rmsnorm2d/script/perf_test.sh index 7b9d0820fd..e26bd1d8ce 100755 --- a/example/ck_tile/10_rmsnorm2d/script/perf_test.sh +++ b/example/ck_tile/10_rmsnorm2d/script/perf_test.sh @@ -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 \ No newline at end of file +$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 \ No newline at end of file diff --git a/example/ck_tile/11_add_rmsnorm2d_rdquant/script/perf_test.sh b/example/ck_tile/11_add_rmsnorm2d_rdquant/script/perf_test.sh index d02b0bab33..c4f56a94e3 100755 --- a/example/ck_tile/11_add_rmsnorm2d_rdquant/script/perf_test.sh +++ b/example/ck_tile/11_add_rmsnorm2d_rdquant/script/perf_test.sh @@ -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 diff --git a/example/ck_tile/12_smoothquant/script/perf_test.sh b/example/ck_tile/12_smoothquant/script/perf_test.sh index 741eb32ec1..3524aee32e 100755 --- a/example/ck_tile/12_smoothquant/script/perf_test.sh +++ b/example/ck_tile/12_smoothquant/script/perf_test.sh @@ -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 \ No newline at end of file +$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 \ No newline at end of file diff --git a/example/ck_tile/13_moe_sorting/script/smoke_test.sh b/example/ck_tile/13_moe_sorting/script/smoke_test.sh old mode 100644 new mode 100755 diff --git a/example/ck_tile/14_moe_smoothquant/script/perf_test.sh b/example/ck_tile/14_moe_smoothquant/script/perf_test.sh index d1e848b930..d48919b0a1 100755 --- a/example/ck_tile/14_moe_smoothquant/script/perf_test.sh +++ b/example/ck_tile/14_moe_smoothquant/script/perf_test.sh @@ -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 \ No newline at end of file +$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 \ No newline at end of file diff --git a/example/ck_tile/16_batched_gemm/script/perf_test.sh b/example/ck_tile/16_batched_gemm/script/perf_test.sh new file mode 100755 index 0000000000..a5ae31dee5 --- /dev/null +++ b/example/ck_tile/16_batched_gemm/script/perf_test.sh @@ -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 + + diff --git a/example/ck_tile/17_grouped_gemm/script/perf_test.sh b/example/ck_tile/17_grouped_gemm/script/perf_test.sh new file mode 100755 index 0000000000..5ec464d068 --- /dev/null +++ b/example/ck_tile/17_grouped_gemm/script/perf_test.sh @@ -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 diff --git a/example/ck_tile/18_flatmm/flatmm_basic.hpp b/example/ck_tile/18_flatmm/flatmm_basic.hpp index 9da3c4c143..5da3eb0196 100644 --- a/example/ck_tile/18_flatmm/flatmm_basic.hpp +++ b/example/ck_tile/18_flatmm/flatmm_basic.hpp @@ -159,7 +159,9 @@ struct GemmConfig 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 struct GemmBasicTypeConfig; diff --git a/example/ck_tile/18_flatmm/script/perf_test.sh b/example/ck_tile/18_flatmm/script/perf_test.sh new file mode 100755 index 0000000000..0baa5ca589 --- /dev/null +++ b/example/ck_tile/18_flatmm/script/perf_test.sh @@ -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 diff --git a/example/ck_tile/18_flatmm/script/smoke_test_basic.sh b/example/ck_tile/18_flatmm/script/smoke_test_basic.sh index a3fc61cc31..22ec94df65 100755 --- a/example/ck_tile/18_flatmm/script/smoke_test_basic.sh +++ b/example/ck_tile/18_flatmm/script/smoke_test_basic.sh @@ -30,5 +30,6 @@ set -x run_tests "bf16" run_tests "fp16" - +run_tests "fp8" +run_tests "bf8" set +x diff --git a/script/run_ck_tile.sh b/script/run_ck_tile.sh index d29a0b064a..91b6832749 100755 --- a/script/run_ck_tile.sh +++ b/script/run_ck_tile.sh @@ -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