Eval bug: ROCm error: batched GEMM not supported

[MrGibus]

Name and Version

llama built from source - latest master.

llama-cli --version
ggml_cuda_init: GGML_CUDA_FORCE_MMQ:    no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 1 ROCm devices:
  Device 0: AMD Radeon RX 7900 XTX, gfx1100 (0x1100), VMM: no, Wave Size: 32
version: 5836 (b9c3eefd)
built with cc (GCC) 14.2.1 20240912 (Red Hat 14.2.1-3) for x86_64-redhat-linux

Build command:

HIPCXX="$(hipconfig -l)/clang" HIP_PATH="$(hipconfig -R)"  cmake -S . -B build -DGGML_HIP=ON -DAMDGPU_TARGETS=gfx1100 -DCMAKE_BUILD_TYPE=Release -DLLAMA_CURL=OFF && cmake --build build --config Release -- -j 16

Operating systems

Linux

GGML backends

HIP

Hardware

AMD Ryzen 5 5600X
AMD ATI Radeon RX 7900 XTX

Models

All models tested failed
Note: Non-Quant models untested.

Problem description & steps to reproduce

llama-run --ngl 10 <MODEL>
fails with any GPU use, ngl 0 functions as expected

First Bad Commit

No response

Relevant log output

llama-run --ngl 10 codellama-7b.Q5_K_S.gguf
> hello
ROCm error: CUBLAS_STATUS_NOT_SUPPORTED
  current device: 0, in function ggml_cuda_mul_mat_batched_cublas_impl at /home/Gibus/dev/repos/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:1928
  hipblasGemmStridedBatchedEx(ctx.cublas_handle(), HIPBLAS_OP_T, HIPBLAS_OP_N, ne01, ne11, ne10, alpha, src0_ptr, cu_data_type_a, nb01/nb00, nb02/nb00, src1_ptr, cu_data_type_b, s11, s12, beta, dst_t, cu_data_type, ne0, ne1*ne0, ne12*ne13, cu_compute_type, HIPBLAS_GEMM_DEFAULT)
/home/Gibus/dev/repos/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:78: ROCm error
[New LWP 14104]
[New LWP 14103]
[New LWP 14102]
[New LWP 14096]
[New LWP 14077]

This GDB supports auto-downloading debuginfo from the following URLs:
  <https://debuginfod.fedoraproject.org/>
Enable debuginfod for this session? (y or [n]) [answered N; input not from terminal]
Debuginfod has been disabled.
To make this setting permanent, add 'set debuginfod enabled off' to .gdbinit.
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib64/libthread_db.so.1".
0x00007f60018ead93 in wait4 () from /lib64/libc.so.6
#0  0x00007f60018ead93 in wait4 () from /lib64/libc.so.6
#1  0x00007f600484702b in ggml_print_backtrace () from /home/Gibus/dev/repos/llama.cpp/build/bin/libggml-base.so
#2  0x00007f600484717e in ggml_abort () from /home/Gibus/dev/repos/llama.cpp/build/bin/libggml-base.so
#3  0x00007f6004469862 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from /home/Gibus/dev/repos/llama.cpp/build/bin/libggml-hip.so
#4  0x00007f6004474176 in ggml_cuda_mul_mat_batched_cublas(ggml_backend_cuda_context&, ggml_tensor const*, ggml_tensor const*, ggml_tensor*) () from /home/Gibus/dev/repos/llama.cpp/build/bin/libggml-hip.so
#5  0x00007f60044710b8 in ggml_cuda_mul_mat(ggml_backend_cuda_context&, ggml_tensor const*, ggml_tensor const*, ggml_tensor*) () from /home/Gibus/dev/repos/llama.cpp/build/bin/libggml-hip.so
#6  0x00007f600446f237 in ggml_backend_cuda_graph_compute(ggml_backend*, ggml_cgraph*) () from /home/Gibus/dev/repos/llama.cpp/build/bin/libggml-hip.so
#7  0x00007f600485df94 in ggml_backend_sched_graph_compute_async () from /home/Gibus/dev/repos/llama.cpp/build/bin/libggml-base.so
#8  0x00007f6004a49e61 in llama_context::graph_compute(ggml_cgraph*, bool) () from /home/Gibus/dev/repos/llama.cpp/build/bin/libllama.so
#9  0x00007f6004a4a0ab in llama_context::process_ubatch(llama_ubatch const&, llm_graph_type, llama_memory_context_i*, ggml_status&) () from /home/Gibus/dev/repos/llama.cpp/build/bin/libllama.so
#10 0x00007f6004a4f045 in llama_context::decode(llama_batch const&) () from /home/Gibus/dev/repos/llama.cpp/build/bin/libllama.so
#11 0x00007f6004a501cb in llama_decode () from /home/Gibus/dev/repos/llama.cpp/build/bin/libllama.so
#12 0x000000000041b926 in main ()
[Inferior 1 (process 14075) detached]
Aborted (core dumped)

[MrGibus]

Been poking around with this.
In ggml-cuda.cu around line 1917 there's an if-else. I removed the condition so it only runs cublasGemmBatchedEx instead of cublasGemmStridedBatchedEx but also a similar error to the above.

So I was not fully convinced that this is actually unsupported.
Managed to run a bench for gemm_strided_batched_ex in hipblas-bench

~$ /opt/rocm/bin/hipblas-bench \
    -f gemm_strided_batched_ex \
    -r s \
    --transposeA N  \
    --transposeB N \
    -m 1024 -n 1024 -k 1024 \
    --alpha 1.0 --beta 0.0  \
    --lda 1024 --ldb 1024 --ldc 1024  \
    --stride_a $((1024 * 1024))  \
    --stride_b $((1024 * 1024))  \
    --stride_c $((1024 * 1024)) \
    --batch_count 8

and got a result.
Note this works with both -r s f32 and -r h f16 precision

Query device success: there are 1 devices
Device ID 0 : AMD Radeon RX 7900 XTX ------------------------------------------------------
with 25.8 GB memory, clock rate 2482MHz @ computing capability 11.0
maxGridDimX 2147483647, sharedMemPerBlock 65.5 KB, maxThreadsPerBlock 1024, warpSize 32
-------------------------------------------------------------------------

hipblas-bench INFO: ldd < min_ldd, set ldd = 1024
a_type,c_type,compute_type,transA,transB,M,N,K,alpha,lda,ldb,beta,ldc,stride_scale,batch_count,with_flags,flags,hipblas-Gflops,hipblas-GB/s,hipblas-us,
f32_r,f32_r,f32_r,N,N,1024,1024,1024,1,1024,1024,0,1024,1,8,0,0,15847.1, 92.8543, 1084.1,

So I'm not sure what's going on here.