Misc. bug: CUDA ggml_top_k() implementation crashes for large tensor shapes

[fairydreaming]

Name and Version

./bin/llama-cli --version
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
version: 8580 (7c20367)
built with GNU 13.3.0 for Linux x86_64

Operating systems

Linux

Which llama.cpp modules do you know to be affected?

Other (Please specify in the next section)

Command line

./bin/test-backend-ops -o TOP_K

Problem description & steps to reproduce

How to reproduce the problem:

Step 1: Apply the following patch to test-backend-ops.cpp:

diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 6a4f9b634..096ec58b1 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -8461,6 +8461,10 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     //    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {i, 2, 1, 3}, rand() % i + 1));
     //}
 
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {512, 512, 1, 1}, 512));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {1024, 512, 1, 1}, 1024));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {1536, 512, 1, 1}, 1536));
+
     for (ggml_scale_mode mode : {GGML_SCALE_MODE_NEAREST, GGML_SCALE_MODE_BILINEAR, GGML_SCALE_MODE_BICUBIC, ggml_scale_mode(GGML_SCALE_MODE_BILINEAR | GGML_SCALE_FLAG_ANTIALIAS)}) {
         test_cases.emplace_back(new test_upscale(GGML_TYPE_F32, {512, 512, 3, 2}, 2, mode));
         test_cases.emplace_back(new test_upscale(GGML_TYPE_F32, {512, 512, 3, 2}, 2, mode, true));

Step 2: Make a CUDA build of llama.cpp. CUDA version I used:

nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Fri_Feb_21_20:23:50_PST_2025
Cuda compilation tools, release 12.8, V12.8.93
Build cuda_12.8.r12.8/compiler.35583870_0

Step 3: Run TOP_K op tests:

./bin/test-backend-ops -o TOP_K

Observed result: Test crashes

ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[512,512,1,1],k=512,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[1024,512,1,1],k=1024,ties=0): OK
CUDA error: an illegal memory access was encountered
  current device: 0, in function ggml_backend_cuda_synchronize at /home/phm/projects/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2924

Full logs below.

Expected result: Test finishes successfully

Rationale: #21149 needs this working for tensor shapes up to {context size, ubatch size, 1, B}.

Note: Everything works correctly if you build llama.cpp with NVIDIA CUDA CCCL library (version >3.2) and enabled GGML_CUDA_USE_CUB. Tested on CUDA 13.2 and CCCL 13.2.27.

First Bad Commit

No idea if this ever worked correctly for tensors this big.

Relevant log output

Logs

$ ./bin/test-backend-ops -o TOP_K
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
Testing 2 devices

Backend 1/2: CUDA0
  Device description: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition
  Device memory: 97247 MB (96640 MB free)

  TOP_K(type=f32,ne=[1,2,1,3],k=1,ties=1): OK
  TOP_K(type=f32,ne=[2,2,1,3],k=1,ties=1): OK
  TOP_K(type=f32,ne=[2,2,1,3],k=2,ties=1): OK
  TOP_K(type=f32,ne=[3,2,1,3],k=1,ties=1): OK
  TOP_K(type=f32,ne=[3,2,1,3],k=2,ties=1): OK
  TOP_K(type=f32,ne=[3,2,1,3],k=3,ties=1): OK
  TOP_K(type=f32,ne=[4,2,1,3],k=1,ties=1): OK
  TOP_K(type=f32,ne=[4,2,1,3],k=2,ties=1): OK
  TOP_K(type=f32,ne=[4,2,1,3],k=3,ties=1): OK
  TOP_K(type=f32,ne=[4,2,1,3],k=4,ties=1): OK
  TOP_K(type=f32,ne=[1,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[12,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[12,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[2,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[13,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[13,1,2,1],k=1,ties=1): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[2,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[13,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[13,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[4,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[15,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[15,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[4,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[15,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[15,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[4,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[15,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[15,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[8,1,1,1],k=1,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[19,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[19,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[8,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[19,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[19,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[8,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[19,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[19,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[8,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[19,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[19,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[16,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[27,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[27,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[16,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[27,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[27,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[16,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[27,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[27,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[16,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[27,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[27,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[16,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[27,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[27,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[32,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[43,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[43,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[32,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[43,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[43,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[32,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[43,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[43,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[32,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[43,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[43,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[32,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[43,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[43,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[64,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[75,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[75,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[64,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[75,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[75,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[64,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[75,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[75,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[64,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[75,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[75,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[64,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[75,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[75,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[128,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[128,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[128,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[128,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[128,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[128,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[139,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[256,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[256,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[256,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[256,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[256,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[256,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[267,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[512,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[512,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[512,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[512,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[512,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[512,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[512,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[523,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[1024,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[1024,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[1024,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[1024,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[1024,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[1024,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[1024,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[1024,1,1,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[1035,1,2,1],k=1023,ties=1): OK
  TOP_K(type=f32,ne=[2048,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[2048,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[2048,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[2048,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[2048,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[2048,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[2048,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[2048,1,1,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[2059,1,2,1],k=1023,ties=1): OK
  TOP_K(type=f32,ne=[4096,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[4096,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[4096,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[4096,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[4096,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[4096,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[4096,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[4096,1,1,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[4107,1,2,1],k=1023,ties=1): OK
  TOP_K(type=f32,ne=[8192,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[8192,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[8192,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[8192,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[8192,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[8192,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[8192,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[8192,1,1,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[8203,1,2,1],k=1023,ties=1): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=1023,ties=1): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[16395,1,2,1],k=9999,ties=1): OK
  TOP_K(type=f32,ne=[32768,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[32768,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[32768,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[32768,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[32768,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[32768,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[32768,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[32768,1,1,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=1023,ties=1): OK
  TOP_K(type=f32,ne=[32768,1,1,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[32779,1,2,1],k=9999,ties=1): OK
  TOP_K(type=f32,ne=[65536,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[65536,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[65536,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[65536,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[65536,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[65536,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[65536,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[65536,1,1,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=1023,ties=1): OK
  TOP_K(type=f32,ne=[65536,1,1,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[65547,1,2,1],k=9999,ties=1): OK
  TOP_K(type=f32,ne=[131072,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[131072,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[131072,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[131072,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[131072,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[131072,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[131072,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[131072,1,1,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=1023,ties=1): OK
  TOP_K(type=f32,ne=[131072,1,1,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[131083,1,2,1],k=9999,ties=1): OK
  TOP_K(type=f32,ne=[262144,1,1,1],k=1,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[262155,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[262144,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[262144,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[262144,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[262144,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[262144,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[262144,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[262144,1,1,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=1023,ties=1): OK
  TOP_K(type=f32,ne=[262144,1,1,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[262155,1,2,1],k=9999,ties=1): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[524288,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=1,ties=1): OK
  TOP_K(type=f32,ne=[524288,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=2,ties=1): OK
  TOP_K(type=f32,ne=[524288,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=3,ties=1): OK
  TOP_K(type=f32,ne=[524288,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=7,ties=1): OK
  TOP_K(type=f32,ne=[524288,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=15,ties=1): OK
  TOP_K(type=f32,ne=[524288,1,1,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=100,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=100,ties=1): OK
  TOP_K(type=f32,ne=[524288,1,1,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=500,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=500,ties=1): OK
  TOP_K(type=f32,ne=[524288,1,1,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=1023,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=1023,ties=1): OK
  TOP_K(type=f32,ne=[524288,1,1,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=9999,ties=0): OK
  TOP_K(type=f32,ne=[524299,1,2,1],k=9999,ties=1): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[16,10,10,10],k=1,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[60,10,10,10],k=1,ties=0): OK
  TOP_K(type=f32,ne=[1023,2,1,3],k=1,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[1024,2,1,3],k=1,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[1025,2,1,3],k=1,ties=0): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=1,ties=0): OK
  TOP_K(type=f32,ne=[2047,2,1,3],k=1,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[2048,2,1,3],k=1,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[2049,2,1,3],k=1,ties=0): OK
  TOP_K(type=f32,ne=[16,10,10,10],k=2,ties=0): OK
  TOP_K(type=f32,ne=[60,10,10,10],k=2,ties=0): OK
  TOP_K(type=f32,ne=[1023,2,1,3],k=2,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[1024,2,1,3],k=2,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[1025,2,1,3],k=2,ties=0): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=2,ties=0): OK
  TOP_K(type=f32,ne=[2047,2,1,3],k=2,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[2048,2,1,3],k=2,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[2049,2,1,3],k=2,ties=0): OK
  TOP_K(type=f32,ne=[16,10,10,10],k=3,ties=0): OK
  TOP_K(type=f32,ne=[60,10,10,10],k=3,ties=0): OK
  TOP_K(type=f32,ne=[1023,2,1,3],k=3,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[1024,2,1,3],k=3,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[1025,2,1,3],k=3,ties=0): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=3,ties=0): OK
  TOP_K(type=f32,ne=[2047,2,1,3],k=3,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[2048,2,1,3],k=3,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[2049,2,1,3],k=3,ties=0): OK
  TOP_K(type=f32,ne=[16,10,10,10],k=7,ties=0): OK
  TOP_K(type=f32,ne=[60,10,10,10],k=7,ties=0): OK
  TOP_K(type=f32,ne=[1023,2,1,3],k=7,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[1024,2,1,3],k=7,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[1025,2,1,3],k=7,ties=0): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=7,ties=0): OK
  TOP_K(type=f32,ne=[2047,2,1,3],k=7,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[2048,2,1,3],k=7,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[2049,2,1,3],k=7,ties=0): OK
  TOP_K(type=f32,ne=[16,10,10,10],k=15,ties=0): OK
  TOP_K(type=f32,ne=[60,10,10,10],k=15,ties=0): OK
  TOP_K(type=f32,ne=[1023,2,1,3],k=15,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[1024,2,1,3],k=15,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[1025,2,1,3],k=15,ties=0): OK
  TOP_K(type=f32,ne=[16384,1,1,1],k=15,ties=0): OK
  TOP_K(type=f32,ne=[2047,2,1,3],k=15,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[2048,2,1,3],k=15,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[2049,2,1,3],k=15,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
  TOP_K(type=f32,ne=[512,512,1,1],k=512,ties=0): OK
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
  TOP_K(type=f32,ne=[1024,512,1,1],k=1024,ties=0): OK
CUDA error: an illegal memory access was encountered
  current device: 0, in function ggml_backend_cuda_synchronize at /home/phm/projects/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2924
  cudaStreamSynchronize(cuda_ctx->stream())
/home/phm/projects/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:98: CUDA error
[New LWP 32805]
[New LWP 32804]
[New LWP 32803]
[New LWP 32802]
[New LWP 32801]
[New LWP 32797]
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
0x00007a1f50310813 in __GI___wait4 (pid=32806, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
warning: 30	../sysdeps/unix/sysv/linux/wait4.c: No such file or directory
#0  0x00007a1f50310813 in __GI___wait4 (pid=32806, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
30	in ../sysdeps/unix/sysv/linux/wait4.c
#1  0x00007a1f509d97b3 in ggml_print_backtrace () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-base.so.0
#2  0x00007a1f509d995b in ggml_abort () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-base.so.0
#3  0x00007a1f4df8b627 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#4  0x00007a1f4df8cad8 in ggml_backend_cuda_synchronize(ggml_backend*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#5  0x00007a1f509f0c7c in ggml_backend_graph_compute () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-base.so.0
#6  0x00007a1f509f5be1 in ggml_backend_compare_graph_backend () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-base.so.0
#7  0x000058fda0bda438 in test_case::eval(ggml_backend*, ggml_backend*, char const*, printer*) ()
#8  0x000058fda0ba110e in test_backend(ggml_backend*, test_mode, char const*, char const*, printer*, char const*) ()
#9  0x000058fda0b70ec0 in main ()
[Inferior 1 (process 32796) detached]
Aborted (core dumped)

[CISC]

Right, clearly missing a bounds check there as well as support-check when there's no CUB.

cc/ @am17an

[am17an]

I had a quick look at this and didn't see the problem immediately. cc @danbev @ORippler who recently worked on the top-k kernel

[ORippler]

@fairydreaming How exactly did you configure your CUDA build? On CTK 13.1 and CCCL 3.1, tests pass for me.

[am17an]

@ORippler try CUDA 12.9. I could replicate the failure

[fairydreaming]

@fairydreaming How exactly did you configure your CUDA build? On CTK 13.1 and CCCL 3.1, tests pass for me.

@ORippler I tested:
CUDA 12.8, no CCCL -> crash
CUDA 13.2 and CCCL 13.2.27 -> all OK

[ORippler]

Feel free to try out #21181 :

.build/bin/test-backend-ops -o 'ARGSORT(type=f32,ne=[1025,256,1,1],order=0)'

The test should pass on older CCCL versions now

[fairydreaming]

@ORippler Looks good, I tortured it up to:

  ...
  TOP_K(type=f32,ne=[164352,2048,1,1],k=2048,ties=0): OK

and there are no crashes.

[fairydreaming]

@ORippler I was profiling my implementation and when tried to measure TOP_K performance in test-backend-ops (with perf option) noticed that it still crashes for large tensors (CUDA 12.8):

$ ./bin/test-backend-ops perf -o "TOP_K(type=f32,ne=[2048,512,1,1],k=2048,ties=0)"
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
Testing 2 devices

Backend 1/2: CUDA0
  Device description: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition
  Device memory: 97247 MB (96640 MB free)

ggml_backend_cuda_graph_compute: CUDA graph warmup complete
CUDA error: operation failed due to a previous error during capture
  current device: 0, in function ggml_cuda_op_top_k at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/top-k.cu:86
  cudaMemcpy2DAsync(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), nrows, cudaMemcpyDeviceToDevice, stream)
/home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/ggml-cuda.cu:100: CUDA error
[New LWP 48233]
[New LWP 48232]
[New LWP 48228]
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
0x000073f340f10813 in __GI___wait4 (pid=48234, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
warning: 30	../sysdeps/unix/sysv/linux/wait4.c: No such file or directory
#0  0x000073f340f10813 in __GI___wait4 (pid=48234, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
30	in ../sysdeps/unix/sysv/linux/wait4.c
#1  0x000073f3416917b3 in ggml_print_backtrace () from /home/phm/projects/llama.cpp-deepseek-dsa/build-cuda/bin/libggml-base.so.0
#2  0x000073f34169195b in ggml_abort () from /home/phm/projects/llama.cpp-deepseek-dsa/build-cuda/bin/libggml-base.so.0
#3  0x000073f33e9a72b7 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from /home/phm/projects/llama.cpp-deepseek-dsa/build-cuda/bin/libggml-cuda.so.0
#4  0x000073f33eb12215 in ggml_cuda_op_top_k(ggml_backend_cuda_context&, ggml_tensor*) () from /home/phm/projects/llama.cpp-deepseek-dsa/build-cuda/bin/libggml-cuda.so.0
#5  0x000073f33e9b98fc in ggml_cuda_compute_forward(ggml_backend_cuda_context&, ggml_tensor*) () from /home/phm/projects/llama.cpp-deepseek-dsa/build-cuda/bin/libggml-cuda.so.0
#6  0x000073f33e9be97b in ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context*, ggml_cgraph*, bool, bool, void const*) () from /home/phm/projects/llama.cpp-deepseek-dsa/build-cuda/bin/libggml-cuda.so.0
#7  0x000073f33e9c062e in ggml_backend_cuda_graph_compute(ggml_backend*, ggml_cgraph*) () from /home/phm/projects/llama.cpp-deepseek-dsa/build-cuda/bin/libggml-cuda.so.0
#8  0x000073f3416a9062 in ggml_backend_graph_compute () from /home/phm/projects/llama.cpp-deepseek-dsa/build-cuda/bin/libggml-base.so.0
#9  0x00005c0c98ce8efe in test_backend(ggml_backend*, test_mode, char const*, char const*, printer*, char const*) ()
#10 0x00005c0c98cb6f00 in main ()
[Inferior 1 (process 48227) detached]
Aborted (core dumped)

here's a detailed log from a debug build:

$ ./bin/test-backend-ops perf -o "TOP_K(type=f32,ne=[2048,512,1,1],k=2048,ties=0)"
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
register_backend: registered backend CUDA (1 devices)
register_device: registered device CUDA0 (NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition)
register_backend: registered backend CPU (1 devices)
register_device: registered device CPU (AMD EPYC 9374F 32-Core Processor)
load_backend: failed to find ggml_backend_init in /home/phm/projects/llama.cpp-deepseek-dsa/build-cuda-debug/bin/libggml-cuda.so
load_backend: failed to find ggml_backend_init in /home/phm/projects/llama.cpp-deepseek-dsa/build-cuda-debug/bin/libggml-cpu.so
Testing 2 devices

Backend 1/2: CUDA0
  Device description: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition
  Device memory: 97247 MB (96640 MB free)

ggml_backend_cuda_graph_compute: CUDA graph warmup complete
CUDA error: operation failed due to a previous error during capture
  current device: 0, in function ggml_cuda_op_top_k at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/top-k.cu:86
  cudaMemcpy2DAsync(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), nrows, cudaMemcpyDeviceToDevice, stream)
/home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/ggml-cuda.cu:100: CUDA error
[New LWP 49509]
[New LWP 49508]
[New LWP 49504]
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
0x00007427bef10813 in __GI___wait4 (pid=49510, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
warning: 30	../sysdeps/unix/sysv/linux/wait4.c: No such file or directory
#0  0x00007427bef10813 in __GI___wait4 (pid=49510, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
30	in ../sysdeps/unix/sysv/linux/wait4.c
#1  0x00007427c490eea7 in ggml_print_backtrace () at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml.c:219
219	        waitpid(child_pid, NULL, 0);
#2  0x00007427c490f05b in ggml_abort (file=0x7427c06bdbd0 "/home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/ggml-cuda.cu", line=100, fmt=0x7427c06bdbc5 "CUDA error") at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml.c:253
253	        ggml_print_backtrace();
#3  0x00007427bf882991 in ggml_cuda_error (stmt=0x7427c07b3810 "cudaMemcpy2DAsync(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), nrows, cudaMemcpyDeviceToDevice, stream)", func=0x7427c07b37fa "ggml_cuda_op_top_k", file=0x7427c07b3780 "/home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/top-k.cu", line=86, msg=0x7427bea93f90 "operation failed due to a previous error during capture") at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/ggml-cuda.cu:100
100	    GGML_ABORT(GGML_CUDA_NAME " error");
#4  0x00007427bfa96e52 in ggml_cuda_op_top_k (ctx=..., dst=0x6467dd2c36d0) at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/top-k.cu:86
86	    CUDA_CHECK(cudaMemcpy2DAsync(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), nrows,
#5  0x00007427bf88d508 in ggml_cuda_compute_forward (ctx=..., dst=0x6467dd2c36d0) at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/ggml-cuda.cu:2779
2779	            ggml_cuda_op_top_k(ctx, dst);
#6  0x00007427bf893ae5 in ggml_cuda_graph_evaluate_and_capture (cuda_ctx=0x6467dc0f17c0, cgraph=0x6467dd2c3840, use_cuda_graph=true, cuda_graph_update_required=true, graph_key=0x6467dd2c36d0) at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/ggml-cuda.cu:4047
4047	                bool ok = ggml_cuda_compute_forward(*cuda_ctx, node);
#7  0x00007427bf8942c6 in ggml_backend_cuda_graph_compute (backend=0x6467dbf45e90, cgraph=0x6467dd2c3840) at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-cuda/ggml-cuda.cu:4167
4167	    ggml_cuda_graph_evaluate_and_capture(cuda_ctx, cgraph, use_cuda_graph, cuda_graph_update_required, graph_key);
#8  0x00007427c492c5cf in ggml_backend_graph_compute_async (backend=0x6467dbf45e90, cgraph=0x6467dd2c3840) at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-backend.cpp:366
366	    return backend->iface.graph_compute(backend, cgraph);
#9  0x00007427c492c55b in ggml_backend_graph_compute (backend=0x6467dbf45e90, cgraph=0x6467dd2c3840) at /home/phm/projects/llama.cpp-deepseek-dsa/ggml/src/ggml-backend.cpp:359
359	    enum ggml_status err = ggml_backend_graph_compute_async(backend, cgraph);
#10 0x00006467d5c495ae in test_case::eval_perf (this=0x6467dd2647e0, backend=0x6467dbf45e90, op_names_filter=0x7ffd35af697f "TOP_K(type=f32,ne=[2048,512,1,1],k=2048,ties=0)", output_printer=0x6467dc0221e0) at /home/phm/projects/llama.cpp-deepseek-dsa/tests/test-backend-ops.cpp:1556
1556	            ggml_status status = ggml_backend_graph_compute(backend, gf);
#11 0x00006467d5c37aae in test_backend (backend=0x6467dbf45e90, mode=MODE_PERF, op_names_filter=0x7ffd35af697f "TOP_K(type=f32,ne=[2048,512,1,1],k=2048,ties=0)", params_filter=0x0, output_printer=0x6467dc0221e0, test_file_path=0x0) at /home/phm/projects/llama.cpp-deepseek-dsa/tests/test-backend-ops.cpp:9310
9310	            test->eval_perf(backend, op_names_filter, output_printer);
#12 0x00006467d5c39ae3 in main (argc=4, argv=0x7ffd35af4c38) at /home/phm/projects/llama.cpp-deepseek-dsa/tests/test-backend-ops.cpp:9562
9562	        bool ok = test_backend(backend, mode, op_names_filter, params_filter, output_printer.get(), test_file_path);
[Inferior 1 (process 49503) detached]
Aborted (core dumped)

But "CUDA error: operation failed due to a previous error during capture" doesn't sound very helpful.

Note that the same test passes when ran without perf:

$ ./bin/test-backend-ops -o "TOP_K(type=f32,ne=[2048,512,1,1],k=2048,ties=0)"
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
Testing 2 devices

Backend 1/2: CUDA0
  Device description: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition
  Device memory: 97247 MB (96640 MB free)

  TOP_K(type=f32,ne=[2048,512,1,1],k=2048,ties=0): OK
  1/1 tests passed
  Backend CUDA0: OK
Backend 2/2: CPU
  Skipping CPU backend
2/2 backends passed
OK

and it even works with perf if I disable CUDA graphs:

$ GGML_CUDA_DISABLE_GRAPHS=1 ./bin/test-backend-ops perf -o "TOP_K(type=f32,ne=[2048,512,1,1],k=2048,ties=0)"
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
Testing 2 devices

Backend 1/2: CUDA0
  Device description: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition
  Device memory: 97247 MB (96640 MB free)

  TOP_K(type=f32,ne=[2048,512,1,1],k=2048,ties=0):             16388 runs -    79.81 us/run -     8192 kB/run -   97.88 GB/s
  Backend CUDA0: OK
Backend 2/2: CPU
  Skipping CPU backend
2/2 backends passed
OK

Any clue?

[ORippler]

Does #21635 resolve these issues? There was a bug in graph-reuse logic that this PR fixed.

Generally, pulling the top-2048 elements per row out of a 2048 elements-sized vector is a no-op. What are you trying to achieve exactly?

[fairydreaming]

Does #21635 resolve these issues? There was a bug in graph-reuse logic that this PR fixed.

No, the crash still happens in the current master:

$ ./bin/llama-cli --version
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
version: 8728 (5e9c63546)
built with GNU 13.3.0 for Linux x86_64

$ ./bin/test-backend-ops perf -o "TOP_K(type=f32,ne=[2048,512,1,1],k=2048,ties=0)"
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
Testing 2 devices

Backend 1/2: CUDA0
  Device description: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition
  Device memory: 97247 MB (96640 MB free)

ggml_backend_cuda_graph_compute: CUDA graph warmup complete
CUDA error: operation failed due to a previous error during capture
  current device: 0, in function ggml_cuda_op_top_k at /home/phm/projects/llama.cpp/ggml/src/ggml-cuda/top-k.cu:86
  cudaMemcpy2DAsync(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), nrows, cudaMemcpyDeviceToDevice, stream)
/home/phm/projects/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:97: CUDA error
[New LWP 11201]
[New LWP 11200]
[New LWP 11196]
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
0x00007fa636b10813 in __GI___wait4 (pid=11202, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
warning: 30	../sysdeps/unix/sysv/linux/wait4.c: No such file or directory
#0  0x00007fa636b10813 in __GI___wait4 (pid=11202, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
30	in ../sysdeps/unix/sysv/linux/wait4.c
#1  0x00007fa6372577f3 in ggml_print_backtrace () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-base.so.0
#2  0x00007fa63725799b in ggml_abort () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-base.so.0
#3  0x00007fa6345bb887 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#4  0x00007fa634724575 in ggml_cuda_op_top_k(ggml_backend_cuda_context&, ggml_tensor*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#5  0x00007fa6345cded0 in ggml_cuda_compute_forward(ggml_backend_cuda_context&, ggml_tensor*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#6  0x00007fa6345d35ab in ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context*, ggml_cgraph*, bool, bool, void const*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#7  0x00007fa6345d523e in ggml_backend_cuda_graph_compute(ggml_backend*, ggml_cgraph*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#8  0x00007fa63726ecf2 in ggml_backend_graph_compute () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-base.so.0
#9  0x00005e7a4933a0be in test_backend(ggml_backend*, test_mode, char const*, char const*, printer*, char const*) ()
#10 0x00005e7a49308e80 in main ()
[Inferior 1 (process 11195) detached]
Aborted (core dumped)

Generally, pulling the top-2048 elements per row out of a 2048 elements-sized vector is a no-op. What are you trying to achieve exactly?

Just wanted to benchmark these OPs for my DeepSeek DSA implementation. Sure it's a no-op, but not a crash-op.

Besides, TOP_K(type=f32,ne=[4096,512,1,1],k=2048,ties=0) is not a no-op, but it crashes too.

$ ./bin/test-backend-ops perf -o "TOP_K(type=f32,ne=[4096,512,1,1],k=2048,ties=0)"
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97247 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97247 MiB
Testing 2 devices

Backend 1/2: CUDA0
  Device description: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition
  Device memory: 97247 MB (96640 MB free)

ggml_backend_cuda_graph_compute: CUDA graph warmup complete
CUDA error: operation failed due to a previous error during capture
  current device: 0, in function ggml_cuda_op_top_k at /home/phm/projects/llama.cpp/ggml/src/ggml-cuda/top-k.cu:86
  cudaMemcpy2DAsync(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), nrows, cudaMemcpyDeviceToDevice, stream)
/home/phm/projects/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:97: CUDA error
[New LWP 11225]
[New LWP 11224]
[New LWP 11220]
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
0x00007a79b3310813 in __GI___wait4 (pid=11226, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
warning: 30	../sysdeps/unix/sysv/linux/wait4.c: No such file or directory
#0  0x00007a79b3310813 in __GI___wait4 (pid=11226, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
30	in ../sysdeps/unix/sysv/linux/wait4.c
#1  0x00007a79b39347f3 in ggml_print_backtrace () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-base.so.0
#2  0x00007a79b393499b in ggml_abort () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-base.so.0
#3  0x00007a79b0bbb887 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#4  0x00007a79b0d24575 in ggml_cuda_op_top_k(ggml_backend_cuda_context&, ggml_tensor*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#5  0x00007a79b0bcded0 in ggml_cuda_compute_forward(ggml_backend_cuda_context&, ggml_tensor*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#6  0x00007a79b0bd35ab in ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context*, ggml_cgraph*, bool, bool, void const*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#7  0x00007a79b0bd523e in ggml_backend_cuda_graph_compute(ggml_backend*, ggml_cgraph*) () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-cuda.so.0
#8  0x00007a79b394bcf2 in ggml_backend_graph_compute () from /home/phm/projects/llama.cpp/build-cuda/bin/libggml-base.so.0
#9  0x000063d641a3a0be in test_backend(ggml_backend*, test_mode, char const*, char const*, printer*, char const*) ()
#10 0x000063d641a08e80 in main ()
[Inferior 1 (process 11219) detached]
Aborted (core dumped)

Here's my whole torture set if you want to try:

diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index c7129d47a..5f1ab5c21 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -8398,6 +8398,13 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
         test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {2, 8, 8192, 1}, order)); // bailingmoe2 (group selection)
     }
 
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {2048, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {4096, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {8192, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {16384, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {32768, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {65536, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+
     for (int n = 1; n < 5; ++n) {
         for (int k = 1; k <= n; ++k) {
             test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {n, 2, 1, 3}, k, true));
@@ -8429,6 +8436,13 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     //    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {i, 2, 1, 3}, rand() % i + 1));
     //}
 
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {2048, 512, 1, 1}, 2048));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {4096, 512, 1, 1}, 2048));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {8192, 512, 1, 1}, 2048));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {16384, 512, 1, 1}, 2048));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {32768, 512, 1, 1}, 2048));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {65536, 512, 1, 1}, 2048));
+
     for (ggml_scale_mode mode : {GGML_SCALE_MODE_NEAREST, GGML_SCALE_MODE_BILINEAR, GGML_SCALE_MODE_BICUBIC, ggml_scale_mode(GGML_SCALE_MODE_BILINEAR | GGML_SCALE_FLAG_ANTIALIAS)}) {
         test_cases.emplace_back(new test_upscale(GGML_TYPE_F32, {512, 512, 3, 2}, 2, mode));
         test_cases.emplace_back(new test_upscale(GGML_TYPE_F32, {512, 512, 3, 2}, 2, mode, true));
@@ -8913,6 +8927,13 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
     test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {200000, 1,  1, 1}));
     test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {200000, 16, 1, 1}));
 
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {2048, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {4096, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {8192, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {16384, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {32768, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+    test_cases.emplace_back(new test_argsort(GGML_TYPE_F32, {65536, 512, 1, 1}, GGML_SORT_ORDER_DESC));
+
     test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {2, 1, 1, 1}, 1));
     for (auto k : {1, 10, 40, 400}) {
         for (auto nrows : {1, 16}) {
@@ -8922,6 +8943,13 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
         }
     }
 
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {2048, 512, 1, 1}, 2048));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {4096, 512, 1, 1}, 2048));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {8192, 512, 1, 1}, 2048));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {16384, 512, 1, 1}, 2048));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {32768, 512, 1, 1}, 2048));
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {65536, 512, 1, 1}, 2048));
+
     for (auto nrows : {1, 4, 8, 16}) {
         for (auto cols : {128, 1024, 4096, 8192, 16384, 32768, 65536, 131072, 200000, 2000000}) {
             test_cases.emplace_back(new test_cumsum(GGML_TYPE_F32, {cols, nrows, 1, 1}));