Fix bug: fix tie breakers

Author: am17an

ggml/src/ggml-cuda/topk-moe.cu

@@ -85,7 +85,7 @@ __launch_bounds__(4 * WARP_SIZE, 1) __global__ void topk_moe_cuda(const float *
         for (int mask = WARP_SIZE / 2; mask > 0; mask /= 2) {
             const float val    = __shfl_xor_sync(0xFFFFFFFF, max_val, mask, WARP_SIZE);
             const int   expert = __shfl_xor_sync(0xFFFFFFFF, max_expert, mask, WARP_SIZE);
-            if (val > max_val) {
+            if (val > max_val || (val == max_val && expert < max_expert)) {
                 max_val    = val;
                 max_expert = expert;
             }

tests/test-backend-ops.cpp

@@ -6625,7 +6625,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
 
     test_cases.emplace_back(new test_topk_moe({8, 22, 1, 1}, 4));
     test_cases.emplace_back(new test_topk_moe({32, 22, 1, 1}, 8));
-    test_cases.emplace_back(new test_topk_moe({128, 19, 1, 1}, 16));
+    test_cases.emplace_back(new test_topk_moe({128, 1, 1, 1}, 128));
 
 #if 0
     // these tests are disabled to save execution time, sbut they can be handy for debugging