ggml : check cuda and metal argsort limits and add test

[CISC]

CUDA and Metal support only a very limited column width for argsort, check so we can fall back to CPU.

Edit: Looks like Vulkan is limited too, but already checks:
https://github.com/ggml-org/llama.cpp/actions/runs/18081472255/job/51445345624?pr=16323#step:3:10561

OpenCL checks:

llama.cpp/ggml/src/ggml-opencl/ggml-opencl.cpp

Line 2967 in 3ecb2f6

return cols <= max_workgroup_size && op->src[0]->type == GGML_TYPE_F32;

@qnixsynapse @NeoZhangJianyu SYCL is probably limited too, but does not check.

[jeffbolznv]

Is the model really using 16k for ne[0]? That's significantly bigger than what I've seen before and at that size we wouldn't be able to use shared memory on all hardware.

[CISC]

Is the model really using 16k for ne[0]? That's significantly bigger than what I've seen before and at that size we wouldn't be able to use shared memory on all hardware.

Depends on the input, it will be n_tokens * (n_expert / 8).

[CISC]

Is the model really using 16k for ne[0]? That's significantly bigger than what I've seen before and at that size we wouldn't be able to use shared memory on all hardware.

Depends on the input, it will be n_tokens * (n_expert / 8).

I'm thinking I should bypass the first top_k somehow (finding the top 2 tokens in each group) in the group selection as this could get quite slow...

[jeffbolznv]

If you only need the indices of the top 2, then yeah a full sort is overkill. Seems like you'd want a modified argmax...

[CISC]

If you only need the indices of the top 2, then yeah a full sort is overkill. Seems like you'd want a modified argmax...

Hmmm, an argmax2 or even argmaxn could be nice.

[JohannesGaessler]

Preferably keep the order of ggml ops in switch statements consistent with the order in which they're being declared in ggml.h.

[CISC]

Didn't reflect much on the order, just didn't want to unnecessarily break up the fall-throughs, I'll keep it in mind for the future.