fix output ci check bug

example/ck_tile/13_moe_sorting/moe_sorting.cpp

@@ -25,7 +25,6 @@ auto create_args(int argc, char* argv[])
         .insert("e", "8", "number of experts")
         .insert("k", "4", "topk")
         .insert("unit", "32", "unit_size")
-        .insert("st_i", "-1", "row stride of input, -1 means same as topk")
         .insert("seed", "-1", "seed to be used, -1 means random every time")
         .insert("kname", "0", "when set to 1 it will print kernel name")
         .insert("warmup", "5", "number of iterations before benchmark the kernel")
@@ -69,17 +68,11 @@ bool test_moe_sorting(ck_tile::ArgParser args)
     int experts             = args.get_int("e");
     int topk                = args.get_int("k");
     int seed                = args.get_int("seed");
-    int stride_input        = args.get_int("st_i");
     int unit_size           = args.get_int("unit");
     int kname               = args.get_int("kname");
     int warmup              = args.get_int("warmup");
     int repeat              = args.get_int("repeat");
-    int max_output_ids      = (topk * tokens * experts + (unit_size - 1)) / unit_size * unit_size;
-    if(stride_input < 0)
-    {
-        stride_input = topk;
-    }
-    assert(stride_input >= topk);
+    int max_output_ids      = ck_tile::integer_least_multiple(topk * tokens, unit_size) * experts;
 
     if(seed < 0)
     {
@@ -95,8 +88,8 @@ bool test_moe_sorting(ck_tile::ArgParser args)
     }
 
     // tokens already considered batch size
-    ck_tile::HostTensor<IndexType> topk_ids_host({tokens, topk}, {stride_input, 1});
-    ck_tile::HostTensor<WeightType> weights_host({tokens, topk}, {stride_input, 1});
+    ck_tile::HostTensor<IndexType> topk_ids_host({tokens, topk}, {topk, 1});
+    ck_tile::HostTensor<WeightType> weights_host({tokens, topk}, {topk, 1});
     ck_tile::HostTensor<IndexType> sorted_ids_host({max_output_ids}, {1});
     ck_tile::HostTensor<WeightType> sorted_weights_host({max_output_ids}, {1});
     ck_tile::HostTensor<IndexType> expert_ids_host({max_output_ids / unit_size}, {1});
@@ -134,13 +127,12 @@ bool test_moe_sorting(ck_tile::ArgParser args)
                               warmup,
                               repeat};
     auto ms = moe_sorting(trait, karg, sc);
-    printf("[%s|%s]tokens:%d, experts:%d, topk:%d, st_i:%d,  ms:%f , ",
+    printf("[%s|%s]tokens:%d, experts:%d, topk:%d,  ms:%f , ",
            index_prec.c_str(),
            weight_prec.c_str(),
            tokens,
            experts,
            topk,
-           stride_input,
            ms);
     if(ms < 0)
         printf("not supported\n");

example/ck_tile/13_moe_sorting/moe_sorting_api.cpp

@@ -7,6 +7,11 @@ float moe_sorting(moe_sorting_trait t, moe_sorting_kargs a, ck_tile::stream_conf
 {
     if(t.weight_type == "fp32" && t.index_type == "int32")
     {
+        if(t.experts > 127)
+        {
+            printf("lds size exceed, only support experts <127 \n");
+            return -1;
+        }
         using index_t        = ck_tile::index_t;
         using ms_weight_type = float;
         using ms_problem     = ck_tile::MoeSortingProblem<index_t, ms_weight_type>;

example/ck_tile/13_moe_sorting/script/smoke_test.sh

 # #!/bin/sh
 
-# EXE=./build/bin/tile_example_moe_sorting
+EXE=./build/bin/tile_example_moe_sorting
 
-# for pr_i in "fp16" "bf16" ; do
-# $EXE -pr_i=$pr_i -t=80 -e=17
-# $EXE -pr_i=$pr_i -t=111 -e=117
-# $EXE -pr_i=$pr_i -t=1000 -e=55
-# $EXE -pr_i=$pr_i -t=99 -e=180
-# $EXE -pr_i=$pr_i -t=175 -e=64 -k=8
-# $EXE -pr_i=$pr_i -t=65 -e=8 -k=2
-# $EXE -pr_i=$pr_i -t=1 -e=25
-# $EXE -pr_i=$pr_i -t=31 -e=19 -k=15
-# $EXE -pr_i=$pr_i -t=81 -e=37 -k=7
-# $EXE -pr_i=$pr_i -t=199 -e=128 -k=13
-# $EXE -pr_i=$pr_i -t=23 -e=1 -k=1
-# $EXE -pr_i=$pr_i -t=127 -e=99 -k=19 -st_i=233 -st_o=31
-# $EXE -pr_i=$pr_i -t=71 -e=11 -k=11 -st_i=30 -st_o=12
-# $EXE -pr_i=$pr_i -t=1 -e=1 -k=1
-# $EXE -pr_i=$pr_i -t=99 -e=2 -k=1 -st_i=11 -st_o=5
-# $EXE -pr_i=$pr_i -t=333 -e=99 -k=13 -st_i=191 -st_o=17
-# done
+$EXE -t=80 -e=17
+$EXE -t=111 -e=117
+$EXE -t=1000 -e=55
+$EXE -t=99 -e=120
+$EXE -t=175 -e=64 -k=8
+$EXE -t=65 -e=8 -k=2
+$EXE -t=1 -e=25
+$EXE -t=31 -e=19 -k=15
+$EXE -t=81 -e=37 -k=7
+$EXE -t=23 -e=1 -k=1
+$EXE -t=127 -e=99 -k=19
+$EXE -t=71 -e=11 -k=11
+$EXE -t=1 -e=1 -k=1
+$EXE -t=99 -e=2 -k=1
+$EXE -t=333 -e=99 -k=13
\ No newline at end of file