9070XT can't not use amd llama.cpp

[k2210000]

I think It can't use now?
Can we support it or not.

[Peter0x44]

Support what? Both HIP and Vulkan should work well for that GPU.

[k2210000]

Support what? Both HIP and Vulkan should work well for that GPU.

When I use HIP, It will give me this error.

ROCm error: invalid device function

But vulkan can work.

[Tech-no-1]

I can confirm this issue.

I've got a 9070XT as well, Windows 11, I've also got the latest AMD HIP SDK 6.4 installed (don't know if this is relevant).

I've used this build/release: llama-b6346-bin-win-hip-radeon-x64. The gpu gets detected with --list-devices, but when I try to load a model I get this error:

ggml_cuda_compute_forward:` SCALE failed
ROCm error: invalid device function
D:/a/llama.cpp/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:87: ROCm error
  current device: 0, in function ggml_cuda_compute_forward at D:/a/llama.cpp/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2519
  err

[radomeit]

I can confirm the issue too with Radeon RX 9070 XT, Win11, newest llama-b6451-bin-win-hip-radeon-x64.zip, having the latest AMD HIP SDK installed:

ggml_cuda_compute_forward: MUL_MAT failed ROCm error: invalid device function

Details:

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 9070 XT, gfx1201 (0x1201), VMM: no, Wave Size: 32
load_backend: loaded ROCm backend from C:\llama-hip\ggml-hip.dll
load_backend: loaded RPC backend from C:\llama-hip\ggml-rpc.dll
load_backend: loaded CPU backend from C:\llama-hip\ggml-cpu-haswell.dll
build: 6451 (360d6533) with clang version 19.1.5 for x86_64-pc-windows-msvc
main: llama backend init
main: load the model and apply lora adapter, if any
llama_model_load_from_file_impl: using device ROCm0 (AMD Radeon RX 9070 XT) (0000:0a:00.0) - 16152 MiB free
llama_model_loader: loaded meta data with 33 key-value pairs and 459 tensors from C:\...\gpt-oss-20b-MXFP4.gguf (version GGUF V3 (latest))
llama_model_loader: Dumping metadata keys/values. Note: KV overrides do not apply in this output.
llama_model_loader: - kv   0:                       general.architecture str              = gpt-oss
llama_model_loader: - kv   1:                               general.type str              = model
llama_model_loader: - kv   2:                               general.name str              = Openai_Gpt Oss 20b
llama_model_loader: - kv   3:                           general.basename str              = openai_gpt-oss
llama_model_loader: - kv   4:                         general.size_label str              = 20B
llama_model_loader: - kv   5:                        gpt-oss.block_count u32              = 24
llama_model_loader: - kv   6:                     gpt-oss.context_length u32              = 131072
llama_model_loader: - kv   7:                   gpt-oss.embedding_length u32              = 2880
llama_model_loader: - kv   8:                gpt-oss.feed_forward_length u32              = 2880
llama_model_loader: - kv   9:               gpt-oss.attention.head_count u32              = 64
llama_model_loader: - kv  10:            gpt-oss.attention.head_count_kv u32              = 8
llama_model_loader: - kv  11:                     gpt-oss.rope.freq_base f32              = 150000.000000
llama_model_loader: - kv  12:   gpt-oss.attention.layer_norm_rms_epsilon f32              = 0.000010
llama_model_loader: - kv  13:                       gpt-oss.expert_count u32              = 32
llama_model_loader: - kv  14:                  gpt-oss.expert_used_count u32              = 4
llama_model_loader: - kv  15:               gpt-oss.attention.key_length u32              = 64
llama_model_loader: - kv  16:             gpt-oss.attention.value_length u32              = 64
llama_model_loader: - kv  17:           gpt-oss.attention.sliding_window u32              = 128
llama_model_loader: - kv  18:         gpt-oss.expert_feed_forward_length u32              = 2880
llama_model_loader: - kv  19:                  gpt-oss.rope.scaling.type str              = yarn
llama_model_loader: - kv  20:                gpt-oss.rope.scaling.factor f32              = 32.000000
llama_model_loader: - kv  21: gpt-oss.rope.scaling.original_context_length u32              = 4096
llama_model_loader: - kv  22:                       tokenizer.ggml.model str              = gpt2
llama_model_loader: - kv  23:                         tokenizer.ggml.pre str              = gpt-4o
llama_model_loader: - kv  24:                      tokenizer.ggml.tokens arr[str,201088]  = ["!", "\"", "#", "$", "%", "&", "'", ...
llama_model_loader: - kv  25:                  tokenizer.ggml.token_type arr[i32,201088]  = [1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, ...
llama_model_loader: - kv  26:                      tokenizer.ggml.merges arr[str,446189]  = ["Ġ Ġ", "Ġ ĠĠĠ", "ĠĠ ĠĠ", "...
llama_model_loader: - kv  27:                tokenizer.ggml.bos_token_id u32              = 199998
llama_model_loader: - kv  28:                tokenizer.ggml.eos_token_id u32              = 200002
llama_model_loader: - kv  29:            tokenizer.ggml.padding_token_id u32              = 199999
llama_model_loader: - kv  30:                    tokenizer.chat_template str              = {#-\n  In addition to the normal input...
llama_model_loader: - kv  31:               general.quantization_version u32              = 2
llama_model_loader: - kv  32:                          general.file_type u32              = 38
llama_model_loader: - type  f32:  289 tensors
llama_model_loader: - type q8_0:   98 tensors
llama_model_loader: - type mxfp4:   72 tensors
print_info: file format = GGUF V3 (latest)
print_info: file type   = MXFP4 MoE
print_info: file size   = 11.27 GiB (4.63 BPW)
load: printing all EOG tokens:
load:   - 199999 ('<|endoftext|>')
load:   - 200002 ('<|return|>')
load:   - 200007 ('<|end|>')
load:   - 200012 ('<|call|>')
load: special_eog_ids contains both '<|return|>' and '<|call|>' tokens, removing '<|end|>' token from EOG list
load: special tokens cache size = 21
load: token to piece cache size = 1.3332 MB
print_info: arch             = gpt-oss
print_info: vocab_only       = 0
print_info: n_ctx_train      = 131072
print_info: n_embd           = 2880
print_info: n_layer          = 24
print_info: n_head           = 64
print_info: n_head_kv        = 8
print_info: n_rot            = 64
print_info: n_swa            = 128
print_info: is_swa_any       = 1
print_info: n_embd_head_k    = 64
print_info: n_embd_head_v    = 64
print_info: n_gqa            = 8
print_info: n_embd_k_gqa     = 512
print_info: n_embd_v_gqa     = 512
print_info: f_norm_eps       = 0.0e+00
print_info: f_norm_rms_eps   = 1.0e-05
print_info: f_clamp_kqv      = 0.0e+00
print_info: f_max_alibi_bias = 0.0e+00
print_info: f_logit_scale    = 0.0e+00
print_info: f_attn_scale     = 0.0e+00
print_info: n_ff             = 2880
print_info: n_expert         = 32
print_info: n_expert_used    = 4
print_info: causal attn      = 1
print_info: pooling type     = 0
print_info: rope type        = 2
print_info: rope scaling     = yarn
print_info: freq_base_train  = 150000.0
print_info: freq_scale_train = 0.03125
print_info: n_ctx_orig_yarn  = 4096
print_info: rope_finetuned   = unknown
print_info: model type       = 20B
print_info: model params     = 20.91 B
print_info: general.name     = Openai_Gpt Oss 20b
print_info: n_ff_exp         = 2880
print_info: vocab type       = BPE
print_info: n_vocab          = 201088
print_info: n_merges         = 446189
print_info: BOS token        = 199998 '<|startoftext|>'
print_info: EOS token        = 200002 '<|return|>'
print_info: EOT token        = 199999 '<|endoftext|>'
print_info: PAD token        = 199999 '<|endoftext|>'
print_info: LF token         = 198 'Ċ'
print_info: EOG token        = 199999 '<|endoftext|>'
print_info: EOG token        = 200002 '<|return|>'
print_info: EOG token        = 200012 '<|call|>'
print_info: max token length = 256
load_tensors: loading model tensors, this can take a while... (mmap = true)
load_tensors: offloading 24 repeating layers to GPU
load_tensors: offloading output layer to GPU
load_tensors: offloaded 25/25 layers to GPU
load_tensors:        ROCm0 model buffer size = 10949.38 MiB
load_tensors:   CPU_Mapped model buffer size =   586.82 MiB
................................................................................
llama_context: constructing llama_context
llama_context: n_seq_max     = 1
llama_context: n_ctx         = 8192
llama_context: n_ctx_per_seq = 8192
llama_context: n_batch       = 2048
llama_context: n_ubatch      = 512
llama_context: causal_attn   = 1
llama_context: flash_attn    = disabled
llama_context: kv_unified    = false
llama_context: freq_base     = 150000.0
llama_context: freq_scale    = 0.03125
llama_context: n_ctx_per_seq (8192) < n_ctx_train (131072) -- the full capacity of the model will not be utilized
llama_context:  ROCm_Host  output buffer size =     0.77 MiB
llama_kv_cache_iswa: creating non-SWA KV cache, size = 8192 cells
llama_kv_cache:      ROCm0 KV buffer size =   192.00 MiB
llama_kv_cache: size =  192.00 MiB (  8192 cells,  12 layers,  1/1 seqs), K (f16):   96.00 MiB, V (f16):   96.00 MiB
llama_kv_cache_iswa: creating     SWA KV cache, size = 640 cells
llama_kv_cache:      ROCm0 KV buffer size =    15.00 MiB
llama_kv_cache: size =   15.00 MiB (   640 cells,  12 layers,  1/1 seqs), K (f16):    7.50 MiB, V (f16):    7.50 MiB
llama_context:      ROCm0 compute buffer size =  1072.51 MiB
llama_context:  ROCm_Host compute buffer size =    26.89 MiB
llama_context: graph nodes  = 1446
llama_context: graph splits = 2
common_init_from_params: added <|endoftext|> logit bias = -inf
common_init_from_params: added <|return|> logit bias = -inf
common_init_from_params: added <|call|> logit bias = -inf
common_init_from_params: setting dry_penalty_last_n to ctx_size = 8192
common_init_from_params: warming up the model with an empty run - please wait ... (--no-warmup to disable)
ggml_cuda_compute_forward: MUL_MAT failed
ROCm error: invalid device function
D:/a/llama.cpp/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:87: ROCm error
  current device: 0, in function ggml_cuda_compute_forward at D:/a/llama.cpp/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2527

[slaren]

The current ROCm releases do not have support for this GPU. I will try to add it to the build, but I will need some help to test that it works, since I don't have an AMD GPU.

[slaren]

Please check if this release works with your GPU: https://github.com/slaren/llama.cpp/releases/tag/b6459

[Tech-no-1]

Please check if this release works with your GPU: https://github.com/slaren/llama.cpp/releases/tag/b6459

Unfortunately, it doesn't seem to work. Now the GPU is no longer recognized (with --list-devices) and the CPU is used instead.

./llama-server --list-devices
load_backend: loaded RPC backend from C:\Users\User\llama-b6459-bin-win-hip-radeon-x64\ggml-rpc.dll
load_backend: loaded CPU backend from C:\Users\User\llama-b6459-bin-win-hip-radeon-x64\ggml-cpu-icelake.dll
Available devices:

I got Ollama working by replacing the files in rocblas\library with those from the AMD HIP SDK 6.4, which includes gfx1201, but that alone doesn't work for llama.cpp. With the standard release the GPU is detected, but then I get the same "invalid device function" error.

Edit: When I also added the new hipblas and rocblas.dll files, the GPU was no longer recognized. I then added the entire bin folder from the HIP SDK as a test. After that, the GPU was recognized again, but the "invalid device function" error occurred when loading a model.

[slaren]

@Tech-no-1 Thanks for testing, I think that the problem was a missing dll hipblaslt.dll that is necessary for this version of the HIP SDK, but wasn't previously (this is built using SDK 6.4.2). I have added this file now to the releases, please check if this works:
https://github.com/slaren/llama.cpp/releases/tag/b6460

[channeladam]

FYI, I have been trying to use the rocm docker image with my 9070XT, and also got that error:

common_init_from_params: warming up the model with an empty run - please wait ... (--no-warmup to disable)
set_warmup: value = 1
/app/ggml/src/ggml-cuda/ggml-cuda.cu:87: ROCm error
ggml_cuda_compute_forward: MUL failed
ROCm error: invalid device function
  current device: 0, in function ggml_cuda_compute_forward at /app/ggml/src/ggml-cuda/ggml-cuda.cu:2527
  err

I was able to overcome it by modifying the .devops/rocm.Dockerfile and append ,gfx1200,gfx1201 to the ARG ROCM_DOCKER_ARCH= string, and then rebuild the docker image.

NOTE: this should be perfectly valid because the docker image uses the base image rocm/dev-ubuntu-24.04 which is rocm 6.4.3, and official support for the 9070XT and 9060XT was added in rocm 6.4.0 but only documented in 6.4.1 (see https://rocm.docs.amd.com/projects/install-on-linux/en/docs-6.4.1/reference/system-requirements.html => AMD Radeon tab)

@slaren is that a related code change you can officially update and push through?

[Tech-no-1]

@Tech-no-1 Thanks for testing, I think that the problem was a missing dll hipblaslt.dll that is necessary for this version of the HIP SDK, but wasn't previously (this is built using SDK 6.4.2). I have added this file now to the releases, please check if this works: https://github.com/slaren/llama.cpp/releases/tag/b6460

Yes, it works! The GPU is now recognized and the model is loading without the previous error. Thank you very much.

Just one more small thing, now I get this in the logs:

rocblaslt error: Cannot read C:\User\llama-b6460-bin-win-hip-radeon-x64\/hipblaslt/library/TensileLibrary_lazy_gfx1201.dat: No such file or directory

rocblaslt error: Could not load C:\User\llama-b6460-bin-win-hip-radeon-x64\/hipblaslt/library/TensileLibrary_lazy_gfx1201.dat
hipBLASLt error: Heuristic Fetch Failed!
This message will be only be displayed once, unless the ROCBLAS_VERBOSE_HIPBLASLT_ERROR environment variable is set.

rocBLAS warning: hipBlasLT failed, falling back to tensile.
This message will be only be displayed once, unless the ROCBLAS_VERBOSE_TENSILE_ERROR environment variable is set.

Adding the hipblaslt\library from the SDK fixes this.

[slaren]

@slaren is that a related code change you can officially update and push through?

Yes, open a PR to add these archs to the dockerfile and I will merge it.