HIPBlas AMD Radeon VII under windows. Quits out after "ggml_backend_sched: failed to allocate graph, reserving"

[moozoo64]

llama.cpp 25/02/2024 "git pull" ~[b2254]
Windows 10 (latest fully updated) i7-4770 16GB ram
Radeon VII 16GB vram
I'm compiling and running miniconda
I can build and run the Vulkan version fine.
However when I try HipBlas both main and server quit out after "ggml_backend_sched: failed to allocate graph, reserving"
Is this because ggml-cuda.cu doesn't support gfx906?

Building from latest Release
$env:path = "C:\SDK\LLM\ccache-4.9.1-windows-x86_64;C:\Program Files\AMD\ROCm\5.7\bin;" + $env:path
$env:HIP_PLATFORM="amd"
HipInfo
....
gcnArchName: gfx906:sramecc-:xnack-
.....
mkdir build
cd .\build
cmake -G Ninja -DAMDGPU_TARGETS=gfx906 -DLLAMA_HIPBLAS=ON -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ ..
{no errors}
cmake --build .
{bunch of warning related to [-Wgnu-zero-variadic-macro-arguments] , [-Wlanguage-extension-token] etc}
cd .\bin
.\main.exe -m H:\LLM\models\tinyllama-1.1b-chat-v1.0.Q5_K_M.gguf -ngl 40 -sm none -n -1 --color -r "User:" --in-prefix " " -i -e -p "User: Hi\nAI: Hello. I am an AI chatbot. Would you like to talk?\nUser: Sure!\nAI: What would you like to talk about?\nUser:"

Appears to load up using Rocm.

---

Log start
main: build = 2254 (9e359a4)
main: built with for x86_64-pc-windows-msvc
main: seed = 1708864556
ggml_init_cublas: GGML_CUDA_FORCE_MMQ: no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 ROCm devices:
Device 0: AMD Radeon VII, compute capability 9.0, VMM: no
........
........

llm_load_print_meta: LF token = 13 '<0x0A>'
llm_load_tensors: ggml ctx size = 0.15 MiB
llm_load_tensors: offloading 22 repeating layers to GPU
llm_load_tensors: offloading non-repeating layers to GPU
llm_load_tensors: offloaded 23/23 layers to GPU
llm_load_tensors: ROCm0 buffer size = 702.14 MiB
llm_load_tensors: CPU buffer size = 42.97 MiB
.........................................................................................
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: ROCm0 KV buffer size = 11.00 MiB
llama_new_context_with_model: KV self size = 11.00 MiB, K (f16): 5.50 MiB, V (f16): 5.50 MiB
llama_new_context_with_model: ROCm_Host input buffer size = 6.01 MiB
ggml_gallocr_reserve_n: reallocating ROCm0 buffer from size 0.00 MiB to 66.50 MiB
ggml_gallocr_reserve_n: reallocating ROCm_Host buffer from size 0.00 MiB to 4.00 MiB
llama_new_context_with_model: ROCm0 compute buffer size = 66.50 MiB
llama_new_context_with_model: ROCm_Host compute buffer size = 4.00 MiB
llama_new_context_with_model: graph splits (measure): 3
ggml_gallocr_needs_realloc: graph has different number of nodes
ggml_gallocr_alloc_graph: cannot reallocate multi buffer graph automatically, call reserve
ggml_backend_sched: failed to allocate graph, reserving
(LLM) PS C:\SDK\LLM\llama.cpp\build\bin>

Tried different models llama2, mistral different sizes 3b,7b,13b and different F16, Q8_0, Q4_0 etc all the same.

[slaren]

The "failed to allocate graph" is not likely related the issue. My guess is that it is crashing in a call to HIP. You could run it under a debugger to see exactly where it is crashing.

[moozoo64]

Ok... well this is going to be fun.. although in truth I always wanted to learn about lldb
using CodeLLDB

inside
GGML_CALL static bool ggml_backend_cuda_graph_compute(ggml_backend_t backend, ggml_cgraph * cgraph)
when i = 2

decends into
bool ok = ggml_cuda_compute_forward(&params, node);

decends into
static void ggml_cuda_mul_mat(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
Note CUDA_USE_TENSOR_CORES is true... ->use_mul_mat_q true

decends into
ggml_cuda_op_mul_mat(src0, src1, dst, ggml_cuda_op_mul_mat_vec_q, true);

decends into
quantize_row_q8_1_cuda(dev[id].src1_ddf, dev[id].src1_ddq, ne10, nrows1, src1_padded_col_size, stream);

decends into
quantize_q8_1<<<num_blocks, block_size, 0, stream>>>(x, vy, kx, kx_padded);

leads to
@HipMemUnmap
Exception has occurred: Exception
Exception 0xc0000005 encountered at address 0x7ffd915fefe0: Access violation reading location 0x00000020

Going to see what turning off CUDA_USE_TENSOR_CORES does.

[Engininja2]

Is gfx906 supported/functional on Windows HIP? https://rocm.docs.amd.com/projects/install-on-windows/en/latest/reference/system-requirements.html has it listed as unsupported, which might indicate known unfixed bugs. It might also be a red herring meaning only that AMD doesn't want to support it officially. The rocblas in the Windows SDK is compiled for the gfx906 architecture so it's odd that AMD explicitly says it's unsupported even for the runtime.

[moozoo64]

I think that "AMD doesn't want to support it officially"
HipInfo and a few other things work fine.
Also that llama.cpp get so far in run it.
It's running most of what is inside ggml-cuda.cu to get to the point where it fails.
I tried GGML_CUDA_FORCE_MMQ , which turned off CUDA_USE_TENSOR_CORES. that made no difference.
ditto LLAMA_CUDA_FORCE_DMMV
ends up in ggml_cuda_op_mul_mat and hits the quantize_q8_1<<<num_blocks, block_size, 0, stream>>>(x, vy, kx, kx_padded);
giving the HipMemUnmap exception.

[github-actions]

This issue was closed because it has been inactive for 14 days since being marked as stale.

[gjs278]

I actually also get these error messages with a Radeon VII on linux as well.

# llama-cpp -m Meta-Llama-3-8B-Instruct-Q8_0.gguf -n 256 -ngl 33 --interactive-first --color -e --in-prefix '<|start_header_id|>user<|end_header_id|>\n\n' --in-suffix '<|eot_id|><|start_header_id|>assistant  <|end_header_id|>\n\n' -r '<|eot_id|>' --no-mmap --mlock --repeat_last_n 128                                                                                                                                    
Log start                                                                                                                                                                                                       
main: build = 2836 (8c570c9)                                                                                                                                                                                    
main: built with clang version 17.0.6 for x86_64-pc-linux-gnu                                                                                                                                                   
main: seed  = 1715897847                                                                                                                                                                                        
llama_model_loader: loaded meta data with 26 key-value pairs and 291 tensors from Meta-Llama-3-8B-Instruct-Q8_0.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              = llama                                                                                                               
llama_model_loader: - kv   1:                               general.name str              = Meta-Llama-3-8B-Instruct                                                                                            
llama_model_loader: - kv   2:                          llama.block_count u32              = 32                                                                                                                  
llama_model_loader: - kv   3:                       llama.context_length u32              = 8192                                                                                                                
llama_model_loader: - kv   4:                     llama.embedding_length u32              = 4096                                                                                                                
llama_model_loader: - kv   5:                  llama.feed_forward_length u32              = 14336                                                                                                               
llama_model_loader: - kv   6:                 llama.attention.head_count u32              = 32                                                                                                                  
llama_model_loader: - kv   7:              llama.attention.head_count_kv u32              = 8                                                                                                                   
llama_model_loader: - kv   8:                       llama.rope.freq_base f32              = 500000.000000                                                                                                       
llama_model_loader: - kv   9:     llama.attention.layer_norm_rms_epsilon f32              = 0.000010                                                                                                            
llama_model_loader: - kv  10:                          general.file_type u32              = 7                                                                                                                   
llama_model_loader: - kv  11:                           llama.vocab_size u32              = 128256                                                                                                              
llama_model_loader: - kv  12:                 llama.rope.dimension_count u32              = 128                                                                                                                 
llama_model_loader: - kv  13:                       tokenizer.ggml.model str              = gpt2                                                                                                                
llama_model_loader: - kv  14:                         tokenizer.ggml.pre str              = llama-bpe                                                                                                           
llama_model_loader: - kv  15:                      tokenizer.ggml.tokens arr[str,128256]  = ["!", "\"", "#", "$", "%", "&", "'", ...                                                                            
llama_model_loader: - kv  16:                  tokenizer.ggml.token_type arr[i32,128256]  = [1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, ...                                                                            
llama_model_loader: - kv  17:                      tokenizer.ggml.merges arr[str,280147]  = ["Ä  Ä ", "Ä  Ä Ä Ä ", "Ä Ä  Ä Ä ", "...                                                                            
llama_model_loader: - kv  18:                tokenizer.ggml.bos_token_id u32              = 128000                                                                                                              
llama_model_loader: - kv  19:                tokenizer.ggml.eos_token_id u32              = 128001                                                                                                              
llama_model_loader: - kv  20:                    tokenizer.chat_template str              = {% set loop_messages = messages %}{% ...                                                                            
llama_model_loader: - kv  21:               general.quantization_version u32              = 2                                                                                                                   
llama_model_loader: - kv  22:                      quantize.imatrix.file str              = /models/Meta-Llama-3-8B-Instruct-GGUF...                                                                            
llama_model_loader: - kv  23:                   quantize.imatrix.dataset str              = /training_data/groups_merged.txt                                                                                    
llama_model_loader: - kv  24:             quantize.imatrix.entries_count i32              = 224                                                                                                                 
llama_model_loader: - kv  25:              quantize.imatrix.chunks_count i32              = 88                                                                                                                  
llama_model_loader: - type  f32:   65 tensors                                                                                                                                                                   
llama_model_loader: - type q8_0:  226 tensors                                                                                                                                                                   
llm_load_vocab: special tokens definition check successful ( 256/128256 ).                                                                                                                                      
llm_load_print_meta: format           = GGUF V3 (latest)                                                                                                                                                        
llm_load_print_meta: arch             = llama                                                                                                                                                                   
llm_load_print_meta: vocab type       = BPE                                                                                                                                                                     
llm_load_print_meta: n_vocab          = 128256                                                                                                                                                                  
llm_load_print_meta: n_merges         = 280147                                                                                                                                                                  
llm_load_print_meta: n_ctx_train      = 8192                                                                                                                                                                    
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: f_logit_scale    = 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: causal attn      = 1                                                                                                                                                                       
llm_load_print_meta: pooling type     = 0                                                                                                                                                                       
llm_load_print_meta: rope type        = 0                                                                                                                                                                       
llm_load_print_meta: rope scaling     = linear                                                                                                                                                                  
llm_load_print_meta: freq_base_train  = 500000.0                                                                                                                                                                
llm_load_print_meta: freq_scale_train = 1                                                                                                                                                                       
llm_load_print_meta: n_yarn_orig_ctx  = 8192                                                                                                                                                                    
llm_load_print_meta: rope_finetuned   = unknown                                                                                                                                                                 
llm_load_print_meta: ssm_d_conv       = 0                                                                                                                                                                       
llm_load_print_meta: ssm_d_inner      = 0                                                                                                                                                                       
llm_load_print_meta: ssm_d_state      = 0                                                                                                                                                                       
llm_load_print_meta: ssm_dt_rank      = 0                                                                                                                                                                       
llm_load_print_meta: model type       = 8B                                                                                                                                                                      
llm_load_print_meta: model ftype      = Q8_0                                                                                                                                                                    
llm_load_print_meta: model params     = 8.03 B                                                                                                                                                                  
llm_load_print_meta: model size       = 7.95 GiB (8.50 BPW)                                                                                                                                                     
llm_load_print_meta: general.name     = Meta-Llama-3-8B-Instruct                                                                                                                                                
llm_load_print_meta: BOS token        = 128000 '<|begin_of_text|>'                                                                                                                                              
llm_load_print_meta: EOS token        = 128001 '<|end_of_text|>'                                                                                                                                                
llm_load_print_meta: LF token         = 128 'Ã                                                                                                                                                                  
                                              '
llm_load_print_meta: EOT token        = 128009 '<|eot_id|>'                                                                                                                                                     
ggml_cuda_init: GGML_CUDA_FORCE_MMQ:   no                                                                                                                                                                       
ggml_cuda_init: CUDA_USE_TENSOR_CORES: yes                                                                                                                                                                      
ggml_cuda_init: found 1 ROCm devices:                                                                                                                                                                           
  Device 0: AMD Radeon VII, compute capability 9.0, VMM: no                                                                                                                                                     
llm_load_tensors: ggml ctx size =    0.30 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                                                                                                                                                                 
llm_load_tensors:      ROCm0 buffer size =  7605.33 MiB                                                                                                                                                         
llm_load_tensors:  ROCm_Host buffer size =   532.31 MiB                                                                                                                                                         
.........................................................................................                                                                                                                       
llama_new_context_with_model: n_ctx      = 512                                                                                                                                                                  
llama_new_context_with_model: n_batch    = 512                                                                                                                                                                  
llama_new_context_with_model: n_ubatch   = 512                                                                                                                                                                  
llama_new_context_with_model: flash_attn = 0                                                                                                                                                                    
llama_new_context_with_model: freq_base  = 500000.0                                                                                                                                                             
llama_new_context_with_model: freq_scale = 1                                                                                                                                                                    
llama_kv_cache_init:      ROCm0 KV buffer size =    64.00 MiB                                                                                                                                                   
llama_new_context_with_model: KV self size  =   64.00 MiB, K (f16):   32.00 MiB, V (f16):   32.00 MiB                                                                                                           
llama_new_context_with_model:  ROCm_Host  output buffer size =     0.49 MiB                                                                                                                                     
ggml_gallocr_reserve_n: reallocating ROCm0 buffer from size 0.00 MiB to 258.50 MiB
ggml_gallocr_reserve_n: reallocating ROCm_Host buffer from size 0.00 MiB to 9.01 MiB
llama_new_context_with_model:      ROCm0 compute buffer size =   258.50 MiB
llama_new_context_with_model:  ROCm_Host compute buffer size =     9.01 MiB
llama_new_context_with_model: graph nodes  = 1030
llama_new_context_with_model: graph splits = 2

system_info: n_threads = 6 / 6 | AVX = 0 | AVX_VNNI = 0 | AVX2 = 0 | AVX512 = 0 | AVX512_VBMI = 0 | AVX512_VNNI = 0 | FMA = 0 | NEON = 0 | ARM_FMA = 0 | F16C = 0 | FP16_VA = 0 | WASM_SIMD = 0 | BLAS = 1 | SSE3 = 1 | SSSE3 = 1 | VSX = 0 | MATMUL_INT8 = 0 | LLAMAFILE = 1 | 
main: interactive mode on.
Reverse prompt: '<|eot_id|>'
Input prefix: '<|start_header_id|>user<|end_header_id|>

'
Input suffix: '<|eot_id|><|start_header_id|>assistant<|end_header_id|>

'
sampling: 
        repeat_last_n = 128, repeat_penalty = 1.000, frequency_penalty = 0.000, presence_penalty = 0.000
        top_k = 40, tfs_z = 1.000, top_p = 0.950, min_p = 0.050, typical_p = 1.000, temp = 0.800
        mirostat = 0, mirostat_lr = 0.100, mirostat_ent = 5.000
sampling order: 
CFG -> Penalties -> top_k -> tfs_z -> typical_p -> top_p -> min_p -> temperature 
generate: n_ctx = 512, n_batch = 2048, n_predict = 256, n_keep = 0

== Running in interactive mode. ==
 - Press Ctrl+C to interject at any time.
 - Press Return to return control to LLaMa.
 - To return control without starting a new line, end your input with '/'.
 - If you want to submit another line, end your input with '\'.

<|begin_of_text|><|start_header_id|>user<|end_header_id|>

how can faceless void avoid being ganked by pudge
<|eot_id|><|start_header_id|>assistant<|end_header_id|>

The eternal struggle of avoiding being ganked by Pudge!

Faceless Void's abilities can be challenging to navigate when facing Pudge's hook and ultimate, but here are some general tips to help you avoid being ganked:

1. **Stay mobile**: Keep moving and avoid standing in one spot for too long. Use your Blink ability to quickly reposition yourself and make it harder for Pudge to hook you.
2. **Use your Teleport**: If you see Pudge attempting to hook you, use your Teleport ability to instantly move away from him. This can also help you reposition quickly to avoid his hooks.
3. **Keep an eye on Pudge**: Monitor Pudge's position and movement closely. If you see him setting up for a hook, try to anticipate his attack and move out of the way in time.
4. **Positioning**: Be mindful of your positioning in team fights. Try to stay close to your allies and avoid being isolated from the rest of your team. This will make it harder for Pudge to hook you.
5. **Avoid being caught out of position**: Make sure you're not caught out of position when Pudge is roaming. If you're alone<|start_header_id|>user<|end_header_id|>


 and see Pudge approaching, try to quickly retreat to a safe location or use your Blink to reposition.
6. **Use your abilities to block**: If you see Pudge's hook coming, use your Force Staff ability to block it. This can help absorb the damage and prevent you from being hooked.
7. **Communicate with your team**: Coordinate with your team to keep an eye on Pudge's position and movement. If you see him setting up for a hook, alert your teammates so they can help you avoid it.
8. **Use your ultimate wisely**: If you have your ultimate ability available, use it to escape or reposition quickly. This can be especially useful if you're caught out of position and Pudge is approaching.

Remember, avoiding being ganked by Pudge requires a combination of quick reflexes, good positioning, and effective use of your abilities. Stay alert and stay one step ahead of Pudge to come out on top!<|eot_id|>
<|start_header_id|>user<|end_header_id|>



<|start_header_id|>user<|end_header_id|>

what if pudge hooks me
<|eot_id|><|start_header_id|>assistant<|end_header_id|>

Don't worry, getting hooked by Pudge is not the end of the world!

If Pudge hooks you, here are some tips to help you survive and getggml_gallocr_needs_realloc: graph has different number of nodes
ggml_gallocr_alloc_graph: cannot reallocate multi buffer graph automatically, call reserve
ggml_backend_sched_alloc_splits: failed to allocate graph, reserving
 back:

Force Staff: Hook!<|eot_id|>
<|start_header_id|>user<|end_header_id|>

as you can see it works for a bit, but then it errors out. after that, input and output are all nonsense. I have to restart it and it works fine.

�[37mROCk module is loaded�[0m
=====================    
HSA System Attributes    
=====================    
Runtime Version:         1.1
System Timestamp Freq.:  1000.000000MHz
Sig. Max Wait Duration:  18446744073709551615 (0xFFFFFFFFFFFFFFFF) (timestamp count)
Machine Model:           LARGE                              
System Endianness:       LITTLE                             
Mwaitx:                  DISABLED
DMAbuf Support:          YES

==========               
HSA Agents               
==========               
*******                  
Agent 1                  
*******                  
  Name:                    Intel(R) Xeon(R) CPU           X5690  @ 3.47GHz
  Uuid:                    CPU-XX                             
  Marketing Name:          Intel(R) Xeon(R) CPU           X5690  @ 3.47GHz
  Vendor Name:             CPU                                
  Feature:                 None specified                     
  Profile:                 FULL_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        0(0x0)                             
  Queue Min Size:          0(0x0)                             
  Queue Max Size:          0(0x0)                             
  Queue Type:              MULTI                              
  Node:                    0                                  
  Device Type:             CPU                                
  Cache Info:              
    L1:                      32768(0x8000) KB                   
  Chip ID:                 0(0x0)                             
  ASIC Revision:           0(0x0)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   3468                               
  BDFID:                   0                                  
  Internal Node ID:        0                                  
  Compute Unit:            6                                  
  SIMDs per CU:            0                                  
  Shader Engines:          0                                  
  Shader Arrs. per Eng.:   0                                  
  WatchPts on Addr. Ranges:1                                  
  Features:                None
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: FINE GRAINED        
      Size:                    49333376(0x2f0c480) KB             
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: KERNARG, FINE GRAINED
      Size:                    49333376(0x2f0c480) KB             
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 3                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    49333376(0x2f0c480) KB             
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
  ISA Info:                
*******                  
Agent 2                  
*******                  
  Name:                    gfx906                             
  Uuid:                    GPU-95ae416172df888e               
  Marketing Name:          AMD Radeon VII                     
  Vendor Name:             AMD                                
  Feature:                 KERNEL_DISPATCH                    
  Profile:                 BASE_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        128(0x80)                          
  Queue Min Size:          64(0x40)                           
  Queue Max Size:          131072(0x20000)                    
  Queue Type:              MULTI                              
  Node:                    1                                  
  Device Type:             GPU                                
  Cache Info:              
    L1:                      16(0x10) KB                        
    L2:                      8192(0x2000) KB                    
  Chip ID:                 26287(0x66af)                      
  ASIC Revision:           1(0x1)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   1801                               
  BDFID:                   1536                               
  Internal Node ID:        1                                  
  Compute Unit:            60                                 
  SIMDs per CU:            4                                  
  Shader Engines:          4                                  
  Shader Arrs. per Eng.:   1                                  
  WatchPts on Addr. Ranges:4                                  
  Features:                KERNEL_DISPATCH 
  Fast F16 Operation:      TRUE                               
  Wavefront Size:          64(0x40)                           
  Workgroup Max Size:      1024(0x400)                        
  Workgroup Max Size per Dimension:
    x                        1024(0x400)                        
    y                        1024(0x400)                        
    z                        1024(0x400)                        
  Max Waves Per CU:        40(0x28)                           
  Max Work-item Per CU:    2560(0xa00)                        
  Grid Max Size:           4294967295(0xffffffff)             
  Grid Max Size per Dimension:
    x                        4294967295(0xffffffff)             
    y                        4294967295(0xffffffff)             
    z                        4294967295(0xffffffff)             
  Max fbarriers/Workgrp:   32                                 
  Packet Processor uCode:: 472                                
  SDMA engine uCode::      145                                
  IOMMU Support::          None                               
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    16760832(0xffc000) KB              
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 2                   
      Segment:                 GLOBAL; FLAGS:                     
      Size:                    16760832(0xffc000) KB              
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 3                   
      Segment:                 GROUP                              
      Size:                    64(0x40) KB                        
      Allocatable:             FALSE                              
      Alloc Granule:           0KB                                
      Alloc Alignment:         0KB                                
      Accessible by all:       FALSE                              
  ISA Info:                
    ISA 1                    
      Name:                    amdgcn-amd-amdhsa--gfx906:sramecc+:xnack-
      Machine Models:          HSA_MACHINE_MODEL_LARGE            
      Profiles:                HSA_PROFILE_BASE                   
      Default Rounding Mode:   NEAR                               
      Default Rounding Mode:   NEAR                               
      Fast f16:                TRUE                               
      Workgroup Max Size:      1024(0x400)                        
      Workgroup Max Size per Dimension:
        x                        1024(0x400)                        
        y                        1024(0x400)                        
        z                        1024(0x400)                        
      Grid Max Size:           4294967295(0xffffffff)             
      Grid Max Size per Dimension:
        x                        4294967295(0xffffffff)             
        y                        4294967295(0xffffffff)             
        z                        4294967295(0xffffffff)             
      FBarrier Max Size:       32                                 
*** Done ***             

this may just be related to Radeon VII in general, I'm happy to help debug