AMD gfx1103 laptop GPU returning `HIPBLAS_STATUS_UNKNOWN`

[lovenemesis]

First, great work on getting AMD GPU support ready on Windows in such a good shape within such a short period. Really appreciate your work!

However, once I switched to Fedora 39, on the same Ryzen 7840U with Radeon 780M laptop, things become a bit puzzling.

At first, it complains about not finding clang++ and hipcc:

HSA_OVERRIDE_GFX_VERSION=11.0.0 ./llamafile-0.6 -ngl 35 -m mixtral-8x7b-instruct-v0.1.Q5_K_M.gguf 
initializing gpu module...
note: won't compile AMD GPU support because $HIP_PATH/bin/clang++ is missing
prebuilt binary /zip/ggml-rocm.so not found
prebuilt binary /zip/ggml-cuda.so not found
fatal error: --n-gpu-layers 35 was passed but no gpus were found

Although, I do have clang++ and hipcc available in $PATH.

sudo rpm -ql hipcc clang
/usr/bin/hipcc
/usr/bin/hipcc.pl
/usr/bin/hipconfig
/usr/bin/hipconfig.pl
/usr/share/perl5/vendor_perl/hipvars.pm
/usr/bin/clang
/usr/bin/clang++
/usr/bin/clang++-17
/usr/bin/clang-17
/usr/bin/clang-cl
/usr/bin/clang-cpp
/usr/lib/.build-id
/usr/lib/.build-id/32
/usr/lib/.build-id/32/e94d93e9ba24c19eb5ffdd7288d637d7cda793
/usr/lib/.build-id/32/e94d93e9ba24c19eb5ffdd7288d637d7cda793.1
/usr/lib/.build-id/32/e94d93e9ba24c19eb5ffdd7288d637d7cda793.2
/usr/lib/.build-id/32/e94d93e9ba24c19eb5ffdd7288d637d7cda793.3
/usr/share/licenses/clang
/usr/share/licenses/clang/LICENSE.TXT
/usr/share/man/man1/clang++-17.1.gz
/usr/share/man/man1/clang++.1.gz
/usr/share/man/man1/clang-17.1.gz
/usr/share/man/man1/clang.1.gz

Then I figured out it might need a bit of manual help, hence adding an environment variable:

HIP_PATH=/usr ./llamafile-0.6 -ngl 35 -m mistral-7b-instruct-v0.1.Q4_K_M.gguf 
initializing gpu module...
extracting /zip/llama.cpp/ggml.h to /home/tommy/.llamafile/ggml.h
extracting /zip/llamafile/compcap.cu to /home/tommy/.llamafile/compcap.cu
extracting /zip/llamafile/llamafile.h to /home/tommy/.llamafile/llamafile.h
extracting /zip/llamafile/tinyblas.h to /home/tommy/.llamafile/tinyblas.h
extracting /zip/llamafile/tinyblas.cu to /home/tommy/.llamafile/tinyblas.cu
extracting /zip/llama.cpp/ggml-impl.h to /home/tommy/.llamafile/ggml-impl.h
extracting /zip/llama.cpp/ggml-cuda.h to /home/tommy/.llamafile/ggml-cuda.h
extracting /zip/llama.cpp/ggml-alloc.h to /home/tommy/.llamafile/ggml-alloc.h
extracting /zip/llama.cpp/ggml-backend.h to /home/tommy/.llamafile/ggml-backend.h
extracting /zip/llama.cpp/ggml-backend-impl.h to /home/tommy/.llamafile/ggml-backend-impl.h
extracting /zip/llama.cpp/ggml-cuda.cu to /home/tommy/.llamafile/ggml-cuda.cu
/usr/bin/rocminfo
hipcc -O3 -fPIC -shared -DNDEBUG --offload-arch=gfx1103 -march=native -mtune=native -use_fast_math -DGGML_BUILD=1 -DGGML_SHARED=1 -Wno-return-type -Wno-unused-result -DGGML_USE_HIPBLAS -DGGML_CUDA_MMV_Y=1 -DGGML_CUDA_DMMV_X=32 -DIGNORE4 -DK_QUANTS_PER_ITERATION=2 -DGGML_CUDA_PEER_MAX_BATCH_SIZE=128 -DIGNORE -o /home/tommy/.llamafile/ggml-rocm.so.rzs98e /home/tommy/.llamafile/ggml-cuda.cu -lhipblas -lrocblas
/home/tommy/.llamafile/ggml-cuda.cu:4595:5: warning: loop not unrolled: the optimizer was unable to perform the requested transformation; the transformation might be disabled or specified as part of an unsupported transformation ordering [-Wpass-failed=transform-warning]
 4595 |     mul_mat_q4_K(
      |     ^
/home/tommy/.llamafile/ggml-cuda.cu:4595:5: warning: loop not unrolled: the optimizer was unable to perform the requested transformation; the transformation might be disabled or specified as part of an unsupported transformation ordering [-Wpass-failed=transform-warning]
/home/tommy/.llamafile/ggml-cuda.cu:4662:1: warning: loop not unrolled: the optimizer was unable to perform the requested transformation; the transformation might be disabled or specified as part of an unsupported transformation ordering [-Wpass-failed=transform-warning]
 4662 | mul_mat_q5_K(
      | ^
/home/tommy/.llamafile/ggml-cuda.cu:4662:1: warning: loop not unrolled: the optimizer was unable to perform the requested transformation; the transformation might be disabled or specified as part of an unsupported transformation ordering [-Wpass-failed=transform-warning]
/home/tommy/.llamafile/ggml-cuda.cu:4731:5: warning: loop not unrolled: the optimizer was unable to perform the requested transformation; the transformation might be disabled or specified as part of an unsupported transformation ordering [-Wpass-failed=transform-warning]
 4731 |     mul_mat_q6_K(
      |     ^
/home/tommy/.llamafile/ggml-cuda.cu:4731:5: warning: loop not unrolled: the optimizer was unable to perform the requested transformation; the transformation might be disabled or specified as part of an unsupported transformation ordering [-Wpass-failed=transform-warning]
6 warnings generated when compiling for gfx1103.
dynamically linking /home/tommy/.llamafile/ggml-rocm.so
GPU support successfully linked and loaded
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 ROCm/CUDA devices:
  Device 0: AMD Radeon Graphics, compute capability 11.0, VMM: no
{"timestamp":1704891476,"level":"INFO","function":"server_cli","line":2812,"message":"build info","build":1500,"commit":"a30b324"}
{"timestamp":1704891476,"level":"INFO","function":"server_cli","line":2815,"message":"system info","n_threads":8,"n_threads_batch":-1,"total_threads":16,"system_info":"AVX = 1 | AVX_VNNI = 0 | AVX2 = 1 | AVX512 = 1 | AVX512_VBMI = 1 | AVX512_VNNI = 1 | FMA = 1 | NEON = 0 | ARM_FMA = 0 | F16C = 1 | FP16_VA = 0 | WASM_SIMD = 0 | BLAS = 1 | SSE3 = 1 | SSSE3 = 1 | VSX = 0 | "}
llama_model_loader: loaded meta data with 20 key-value pairs and 291 tensors from mistral-7b-instruct-v0.1.Q4_K_M.gguf (version GGUF V2)
llama_model_loader: Dumping metadata keys/values. Note: KV overrides do not apply in this output.
llama_model_loader: - kv   0:                       general.architecture str              = llama
llama_model_loader: - kv   1:                               general.name str              = mistralai_mistral-7b-instruct-v0.1
llama_model_loader: - kv   2:                       llama.context_length u32              = 32768
llama_model_loader: - kv   3:                     llama.embedding_length u32              = 4096
llama_model_loader: - kv   4:                          llama.block_count u32              = 32
llama_model_loader: - kv   5:                  llama.feed_forward_length u32              = 14336
llama_model_loader: - kv   6:                 llama.rope.dimension_count u32              = 128
llama_model_loader: - kv   7:                 llama.attention.head_count u32              = 32
llama_model_loader: - kv   8:              llama.attention.head_count_kv u32              = 8
llama_model_loader: - kv   9:     llama.attention.layer_norm_rms_epsilon f32              = 0.000010
llama_model_loader: - kv  10:                       llama.rope.freq_base f32              = 10000.000000
llama_model_loader: - kv  11:                          general.file_type u32              = 15
llama_model_loader: - kv  12:                       tokenizer.ggml.model str              = llama
llama_model_loader: - kv  13:                      tokenizer.ggml.tokens arr[str,32000]   = ["<unk>", "<s>", "</s>", "<0x00>", "<...
llama_model_loader: - kv  14:                      tokenizer.ggml.scores arr[f32,32000]   = [0.000000, 0.000000, 0.000000, 0.0000...
llama_model_loader: - kv  15:                  tokenizer.ggml.token_type arr[i32,32000]   = [2, 3, 3, 6, 6, 6, 6, 6, 6, 6, 6, 6, ...
llama_model_loader: - kv  16:                tokenizer.ggml.bos_token_id u32              = 1
llama_model_loader: - kv  17:                tokenizer.ggml.eos_token_id u32              = 2
llama_model_loader: - kv  18:            tokenizer.ggml.unknown_token_id u32              = 0
llama_model_loader: - kv  19:               general.quantization_version u32              = 2
llama_model_loader: - type  f32:   65 tensors
llama_model_loader: - type q4_K:  193 tensors
llama_model_loader: - type q6_K:   33 tensors
llm_load_vocab: special tokens definition check successful ( 259/32000 ).
llm_load_print_meta: format           = GGUF V2
llm_load_print_meta: arch             = llama
llm_load_print_meta: vocab type       = SPM
llm_load_print_meta: n_vocab          = 32000
llm_load_print_meta: n_merges         = 0
llm_load_print_meta: n_ctx_train      = 32768
llm_load_print_meta: n_embd           = 4096
llm_load_print_meta: n_head           = 32
llm_load_print_meta: n_head_kv        = 8
llm_load_print_meta: n_layer          = 32
llm_load_print_meta: n_rot            = 128
llm_load_print_meta: n_embd_head_k    = 128
llm_load_print_meta: n_embd_head_v    = 128
llm_load_print_meta: n_gqa            = 4
llm_load_print_meta: n_embd_k_gqa     = 1024
llm_load_print_meta: n_embd_v_gqa     = 1024
llm_load_print_meta: f_norm_eps       = 0.0e+00
llm_load_print_meta: f_norm_rms_eps   = 1.0e-05
llm_load_print_meta: f_clamp_kqv      = 0.0e+00
llm_load_print_meta: f_max_alibi_bias = 0.0e+00
llm_load_print_meta: n_ff             = 14336
llm_load_print_meta: n_expert         = 0
llm_load_print_meta: n_expert_used    = 0
llm_load_print_meta: rope scaling     = linear
llm_load_print_meta: freq_base_train  = 10000.0
llm_load_print_meta: freq_scale_train = 1
llm_load_print_meta: n_yarn_orig_ctx  = 32768
llm_load_print_meta: rope_finetuned   = unknown
llm_load_print_meta: model type       = 7B
llm_load_print_meta: model ftype      = Q4_K - Medium
llm_load_print_meta: model params     = 7.24 B
llm_load_print_meta: model size       = 4.07 GiB (4.83 BPW) 
llm_load_print_meta: general.name     = mistralai_mistral-7b-instruct-v0.1
llm_load_print_meta: BOS token        = 1 '<s>'
llm_load_print_meta: EOS token        = 2 '</s>'
llm_load_print_meta: UNK token        = 0 '<unk>'
llm_load_print_meta: LF token         = 13 '<0x0A>'
llm_load_tensors: ggml ctx size       =    0.11 MiB
llm_load_tensors: using ROCm/CUDA for GPU acceleration
llm_load_tensors: system memory used  =   70.42 MiB
llm_load_tensors: VRAM used           = 4095.05 MiB
llm_load_tensors: offloading 32 repeating layers to GPU
llm_load_tensors: offloading non-repeating layers to GPU
llm_load_tensors: offloaded 33/33 layers to GPU
...............................................................................................
llama_new_context_with_model: n_ctx      = 512
llama_new_context_with_model: freq_base  = 10000.0
llama_new_context_with_model: freq_scale = 1
llama_kv_cache_init: VRAM kv self = 64.00 MB
llama_new_context_with_model: KV self size  =   64.00 MiB, K (f16):   32.00 MiB, V (f16):   32.00 MiB
llama_build_graph: non-view tensors processed: 676/676
llama_new_context_with_model: compute buffer total size = 76.19 MiB
llama_new_context_with_model: VRAM scratch buffer: 73.00 MiB
llama_new_context_with_model: total VRAM used: 4232.06 MiB (model: 4095.05 MiB, context: 137.00 MiB)
CUDA error: unknown error
  current device: 0, in function ggml_cuda_op_mul_mat_cublas at /home/tommy/.llamafile/ggml-cuda.cu:7784
  hipblasGemmEx(g_cublas_handles[id], HIPBLAS_OP_T, HIPBLAS_OP_N, row_diff, src1_ncols, ne10, &alpha_f16, src0_ptr, HIPBLAS_R_16F, ne00, src1_ptr, HIPBLAS_R_16F, ne10, &beta_f16, dst_f16.get(), HIPBLAS_R_16F, ldc, HIPBLAS_R_16F, HIPBLAS_GEMM_DEFAULT)
GGML_ASSERT: /home/tommy/.llamafile/ggml-cuda.cu:386: !"CUDA error"

This time it compiles but eventually failed due to cuda error.

May I know what additional steps I should take it get it working?

Thanks,

Originally posted by @lovenemesis in #92 (comment)

[jart]

Thanks for filing a new issue.

I noticed two things orthogonal to your issue, but thanks to your issue, that could be improved. I've done so in 15e2339 and 67d97b5.

Could you build llamafile at HEAD and try again? You should get the same error, but it'll give us something better to go on than "unknown error".

git clone https://github.com/Mozilla-Ocho/llamafile
cd llamafile
make -j16
o//llama.cpp/main/main -ngl 35 --gpu amd -m mixtral-8x7b-instruct-v0.1.Q5_K_M.gguf -p hello

Thanks!

[lovenemesis]

Here's the output. I chose to use mistral model instead, considering only 16G VRAM allocated to GPU on this 7840U laptop:

❯ o//llama.cpp/main/main -ngl 35 --gpu amd -m ../llamafile/mistral-7b-instruct-v0.1.Q4_K_M.gguf -p hello
import_cuda_impl: initializing gpu module...
get_rocm_bin_path: note: amdclang++ not found on $PATH
get_rocm_bin_path: note: $HIP_PATH/bin/amdclang++ does not exist
get_rocm_bin_path: note: /opt/rocm/bin/amdclang++ does not exist
link_cuda_dso: note: dynamically linking /home/tommy/.llamafile/ggml-rocm.so
import_cuda: GPU support successfully linked and loaded
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 ROCm/CUDA devices:
  Device 0: AMD Radeon Graphics, compute capability 11.0, VMM: no
Log start
main: llamafile version 0.6.0
main: seed  = 1704949319
llama_model_loader: loaded meta data with 20 key-value pairs and 291 tensors from ../llamafile/mistral-7b-instruct-v0.1.Q4_K_M.gguf (version GGUF V2)
llama_model_loader: Dumping metadata keys/values. Note: KV overrides do not apply in this output.
llama_model_loader: - kv   0:                       general.architecture str              = llama
llama_model_loader: - kv   1:                               general.name str              = mistralai_mistral-7b-instruct-v0.1
llama_model_loader: - kv   2:                       llama.context_length u32              = 32768
llama_model_loader: - kv   3:                     llama.embedding_length u32              = 4096
llama_model_loader: - kv   4:                          llama.block_count u32              = 32
llama_model_loader: - kv   5:                  llama.feed_forward_length u32              = 14336
llama_model_loader: - kv   6:                 llama.rope.dimension_count u32              = 128
llama_model_loader: - kv   7:                 llama.attention.head_count u32              = 32
llama_model_loader: - kv   8:              llama.attention.head_count_kv u32              = 8
llama_model_loader: - kv   9:     llama.attention.layer_norm_rms_epsilon f32              = 0.000010
llama_model_loader: - kv  10:                       llama.rope.freq_base f32              = 10000.000000
llama_model_loader: - kv  11:                          general.file_type u32              = 15
llama_model_loader: - kv  12:                       tokenizer.ggml.model str              = llama
llama_model_loader: - kv  13:                      tokenizer.ggml.tokens arr[str,32000]   = ["<unk>", "<s>", "</s>", "<0x00>", "<...
llama_model_loader: - kv  14:                      tokenizer.ggml.scores arr[f32,32000]   = [0.000000, 0.000000, 0.000000, 0.0000...
llama_model_loader: - kv  15:                  tokenizer.ggml.token_type arr[i32,32000]   = [2, 3, 3, 6, 6, 6, 6, 6, 6, 6, 6, 6, ...
llama_model_loader: - kv  16:                tokenizer.ggml.bos_token_id u32              = 1
llama_model_loader: - kv  17:                tokenizer.ggml.eos_token_id u32              = 2
llama_model_loader: - kv  18:            tokenizer.ggml.unknown_token_id u32              = 0
llama_model_loader: - kv  19:               general.quantization_version u32              = 2
llama_model_loader: - type  f32:   65 tensors
llama_model_loader: - type q4_K:  193 tensors
llama_model_loader: - type q6_K:   33 tensors
llm_load_vocab: special tokens definition check successful ( 259/32000 ).
llm_load_print_meta: format           = GGUF V2
llm_load_print_meta: arch             = llama
llm_load_print_meta: vocab type       = SPM
llm_load_print_meta: n_vocab          = 32000
llm_load_print_meta: n_merges         = 0
llm_load_print_meta: n_ctx_train      = 32768
llm_load_print_meta: n_embd           = 4096
llm_load_print_meta: n_head           = 32
llm_load_print_meta: n_head_kv        = 8
llm_load_print_meta: n_layer          = 32
llm_load_print_meta: n_rot            = 128
llm_load_print_meta: n_embd_head_k    = 128
llm_load_print_meta: n_embd_head_v    = 128
llm_load_print_meta: n_gqa            = 4
llm_load_print_meta: n_embd_k_gqa     = 1024
llm_load_print_meta: n_embd_v_gqa     = 1024
llm_load_print_meta: f_norm_eps       = 0.0e+00
llm_load_print_meta: f_norm_rms_eps   = 1.0e-05
llm_load_print_meta: f_clamp_kqv      = 0.0e+00
llm_load_print_meta: f_max_alibi_bias = 0.0e+00
llm_load_print_meta: n_ff             = 14336
llm_load_print_meta: n_expert         = 0
llm_load_print_meta: n_expert_used    = 0
llm_load_print_meta: rope scaling     = linear
llm_load_print_meta: freq_base_train  = 10000.0
llm_load_print_meta: freq_scale_train = 1
llm_load_print_meta: n_yarn_orig_ctx  = 32768
llm_load_print_meta: rope_finetuned   = unknown
llm_load_print_meta: model type       = 7B
llm_load_print_meta: model ftype      = Q4_K - Medium
llm_load_print_meta: model params     = 7.24 B
llm_load_print_meta: model size       = 4.07 GiB (4.83 BPW) 
llm_load_print_meta: general.name     = mistralai_mistral-7b-instruct-v0.1
llm_load_print_meta: BOS token        = 1 '<s>'
llm_load_print_meta: EOS token        = 2 '</s>'
llm_load_print_meta: UNK token        = 0 '<unk>'
llm_load_print_meta: LF token         = 13 '<0x0A>'
llm_load_tensors: ggml ctx size       =    0.11 MiB
llm_load_tensors: using ROCm/CUDA for GPU acceleration
llm_load_tensors: system memory used  =   70.42 MiB
llm_load_tensors: VRAM used           = 4095.05 MiB
llm_load_tensors: offloading 32 repeating layers to GPU
llm_load_tensors: offloading non-repeating layers to GPU
llm_load_tensors: offloaded 33/33 layers to GPU
...............................................................................................
llama_new_context_with_model: n_ctx      = 512
llama_new_context_with_model: freq_base  = 10000.0
llama_new_context_with_model: freq_scale = 1
llama_kv_cache_init: VRAM kv self = 64.00 MB
llama_new_context_with_model: KV self size  =   64.00 MiB, K (f16):   32.00 MiB, V (f16):   32.00 MiB
llama_build_graph: non-view tensors processed: 676/676
llama_new_context_with_model: compute buffer total size = 76.19 MiB
llama_new_context_with_model: VRAM scratch buffer: 73.00 MiB
llama_new_context_with_model: total VRAM used: 4232.06 MiB (model: 4095.05 MiB, context: 137.00 MiB)
CUDA error: HIPBLAS_STATUS_UNKNOWN
  current device: 0, in function ggml_cuda_op_mul_mat_cublas at /home/tommy/.llamafile/ggml-cuda.cu:8266
  hipblasGemmEx(g_cublas_handles[id], HIPBLAS_OP_T, HIPBLAS_OP_N, row_diff, src1_ncols, ne10, &alpha_f16, src0_ptr, HIPBLAS_R_16F, ne00, src1_ptr, HIPBLAS_R_16F, ne10, &beta_f16, dst_f16.get(), HIPBLAS_R_16F, ldc, HIPBLAS_R_16F, HIPBLAS_GEMM_DEFAULT)
GGML_ASSERT: /home/tommy/.llamafile/ggml-cuda.cu:390: !"CUDA error"

Since it seems to related with hipblas, below are my installations:

sudo rpm -ql hipblas hipblas-devel
/usr/lib/.build-id
/usr/lib/.build-id/19
/usr/lib/.build-id/19/aead64ad45551a081d5ece41628341839d0f22
/usr/lib/.build-id/19/aead64ad45551a081d5ece41628341839d0f22.1
/usr/lib/.build-id/90
/usr/lib/.build-id/90/5c38f38c531a0a65299c003f29c7706f5da7d9
/usr/lib/.build-id/cc
/usr/lib/.build-id/cc/621395ca064e1036285a414b67b19e908bebde
/usr/lib/.build-id/cc/621395ca064e1036285a414b67b19e908bebde.1
/usr/lib64/cmake/hipblas
/usr/lib64/libhipblas.so.1
/usr/lib64/libhipblas.so.1.1.0
/usr/lib64/rocm/gfx10/lib/libhipblas.so.1
/usr/lib64/rocm/gfx10/lib/libhipblas.so.1.1.0
/usr/lib64/rocm/gfx11/lib/libhipblas.so.1
/usr/lib64/rocm/gfx11/lib/libhipblas.so.1.1.0
/usr/lib64/rocm/gfx8/lib/libhipblas.so.1
/usr/lib64/rocm/gfx8/lib/libhipblas.so.1.1.0
/usr/lib64/rocm/gfx9/lib/libhipblas.so.1
/usr/lib64/rocm/gfx9/lib/libhipblas.so.1.1.0
/usr/share/licenses/hipblas
/usr/share/licenses/hipblas/LICENSE.md
/usr/include/hipblas
/usr/include/hipblas/hipblas-export.h
/usr/include/hipblas/hipblas-version.h
/usr/include/hipblas/hipblas.h
/usr/include/hipblas/hipblas_module.f90
/usr/lib64/cmake/hipblas
/usr/lib64/cmake/hipblas/hipblas-config-version.cmake
/usr/lib64/cmake/hipblas/hipblas-config.cmake
/usr/lib64/cmake/hipblas/hipblas-targets-release.cmake
/usr/lib64/cmake/hipblas/hipblas-targets.cmake
/usr/lib64/libhipblas.so
/usr/lib64/rocm/gfx10/lib/cmake/hipblas
/usr/lib64/rocm/gfx10/lib/cmake/hipblas/hipblas-config-version.cmake
/usr/lib64/rocm/gfx10/lib/cmake/hipblas/hipblas-config.cmake
/usr/lib64/rocm/gfx10/lib/cmake/hipblas/hipblas-targets-release.cmake
/usr/lib64/rocm/gfx10/lib/cmake/hipblas/hipblas-targets.cmake
/usr/lib64/rocm/gfx10/lib/libhipblas.so
/usr/lib64/rocm/gfx11/lib/cmake/hipblas
/usr/lib64/rocm/gfx11/lib/cmake/hipblas/hipblas-config-version.cmake
/usr/lib64/rocm/gfx11/lib/cmake/hipblas/hipblas-config.cmake
/usr/lib64/rocm/gfx11/lib/cmake/hipblas/hipblas-targets-release.cmake
/usr/lib64/rocm/gfx11/lib/cmake/hipblas/hipblas-targets.cmake
/usr/lib64/rocm/gfx11/lib/libhipblas.so
/usr/lib64/rocm/gfx8/lib/cmake/hipblas
/usr/lib64/rocm/gfx8/lib/cmake/hipblas/hipblas-config-version.cmake
/usr/lib64/rocm/gfx8/lib/cmake/hipblas/hipblas-config.cmake
/usr/lib64/rocm/gfx8/lib/cmake/hipblas/hipblas-targets-release.cmake
/usr/lib64/rocm/gfx8/lib/cmake/hipblas/hipblas-targets.cmake
/usr/lib64/rocm/gfx8/lib/libhipblas.so
/usr/lib64/rocm/gfx9/lib/cmake/hipblas
/usr/lib64/rocm/gfx9/lib/cmake/hipblas/hipblas-config-version.cmake
/usr/lib64/rocm/gfx9/lib/cmake/hipblas/hipblas-config.cmake
/usr/lib64/rocm/gfx9/lib/cmake/hipblas/hipblas-targets-release.cmake
/usr/lib64/rocm/gfx9/lib/cmake/hipblas/hipblas-targets.cmake
/usr/lib64/rocm/gfx9/lib/libhipblas.so
/usr/share/doc/hipblas-devel
/usr/share/doc/hipblas-devel/README.md

Let me know anything else I could provide to assist diagnostic.

Thanks,

[jart]

Earlier @cgmb said in #122 (comment) said your gfx1103 laptop gpu "will only work for programs that limit themselves to a compatible subset of the ISA". In that case I'd ask @cgmb if hipBLAS returning HIPBLAS_STATUS_UNKNOWN is indicative of us not using a compatible subset?

Also, @lovenemesis, I noticed you passed HSA_OVERRIDE_GFX_VERSION=11.0.0. Did you find the earlier comment I referenced to figure that out? Should we be specifying that flag automatically? Or should we simply just ignore gfx1103 GPUs from now on?

[lovenemesis]

Regardless of setting HSA_OVERRIDE_GFX_VERSION=11.0.0 or not, I got the same error:

CUDA error: HIPBLAS_STATUS_UNKNOWN
  current device: 0, in function ggml_cuda_op_mul_mat_cublas at /home/tommy/.llamafile/ggml-cuda.cu:8266
  hipblasGemmEx(g_cublas_handles[id], HIPBLAS_OP_T, HIPBLAS_OP_N, row_diff, src1_ncols, ne10, &alpha_f16, src0_ptr, HIPBLAS_R_16F, ne00, src1_ptr, HIPBLAS_R_16F, ne10, &beta_f16, dst_f16.get(), HIPBLAS_R_16F, ldc, HIPBLAS_R_16F, HIPBLAS_GEMM_DEFAULT)
GGML_ASSERT: /home/tommy/.llamafile/ggml-cuda.cu:390: !"CUDA error"

If I rebooted into Win11 X64 on the same laptop, I could offload to the exact gfx1103 GPU on llamafile 0.6 release without any tweaking. Hence, I assume the GPU itself should retain the capability to handle workload. It works well with PyTorch + ROCm, too.

[cgmb]

Earlier @cgmb said in #122 (comment) said your gfx1103 laptop gpu "will only work for programs that limit themselves to a compatible subset of the ISA". In that case I'd ask @cgmb if hipBLAS returning HIPBLAS_STATUS_UNKNOWN is indicative of us not using a compatible subset?

I don't know. When you run machine code that contains instructions that don't exist (or don't mean the same thing as the compiler expected) on the processor you're running it on... you've entered the land of undefined behavior. It's plausible that could be your problem, but it could be something else entirely. As far as I know, AMD does not test ROCm on gfx1103 hardware.

Also, @lovenemesis, I noticed you passed HSA_OVERRIDE_GFX_VERSION=11.0.0. Did you find the earlier comment I referenced to figure that out? Should we be specifying that flag automatically? Or should we simply just ignore gfx1103 GPUs from now on?

rocBLAS/hipBLAS was not designed to support gfx1103, regardless of whether HSA_OVERRIDE_GFX_VERSION is specified or not. That's true of many architectures and there are often simple ways to make things work regardless, but gfx1103 doesn't seem to be one of them.

In the future, I'm hoping that the 'generic' ISAs proposed for introduction into LLVM will avoid the need for ugly hacks like HSA_OVERRIDE_GFX_VERSION. With that said, there might still be issues besides the ISA compatibility that limit the usability of ROCm on APUs like the Radeon 780M.