search

ggerganov / llama.cpp

LLM inference in C/C++
MIT License
64.57k stars 9.24k forks source link

multi-gpu inference produces broken output #3772

Open nih23 opened 10 months ago

nih23 commented 10 months ago

Prerequisites

Please answer the following questions for yourself before submitting an issue.

Expected Behavior

I am running several large language models on my small GPU cluster using the latest version of llama.cpp. The GPU cluster has multiple NVIDIA RTX 3070 GPUs. Inference on a single GPU, enforced by CUDA_VISIBLE_DEVICES=0, of different flavors of LLMs (llama, mistral, mistral german) works as expected, i.e. the model answers my prompt in the appropriate language (German/English) .

CUDA_VISIBLE_DEVICES=0 ./main -ngl 99 -m ../LLM_stack/models/llama-2-7b.Q5_K_M.gguf --color -c 1500 --temp 0.01 -p "Why is the sky blue? Answer for a 5 year old child." -n 100
[...]

Why is the sky blue? Answer for a 5 year old child.
The sky is blue because of the scattering of light by molecules in the atmosphere. The sunlight that reaches us from space has all colors mixed together, but when it passes through our atmosphere, some of its color is scattered away. Blue light scatters more than other colors, so we see a blue sky.

Current Behavior

However, the model is simply returning characters and sharps (#) once I run inference on multiple GPUs:

CUDA_VISIBLE_DEVICES=0,1 ./main -ngl 99 -m ../LLM_stack/models/llama-2-7b.Q5_K_M.gguf --color -c 1500 --temp 0.01 -p "Why is the sky blue? Answer for a 5 year old child." -n 100`

Why is the sky blue? Answer for a 5 year old child. dispos###################################################################################################

Environment and Context

Please provide detailed information about your computer setup. This is important in case the issue is not reproducible except for under certain specific conditions.

$ lscpu
Architecture:                       x86_64
CPU op-mode(s):                     32-bit, 64-bit
Address sizes:                      39 bits physical, 48 bits virtual
Byte Order:                         Little Endian
CPU(s):                             4
On-line CPU(s) list:                0-3

Vendor ID:                          GenuineIntel
Model name:                         Intel(R) Core(TM) i5-7600K CPU @ 3.80GHz
CPU family:                         6
Model:                              158
Thread(s) per core:                 1
Core(s) per socket:                 4
Socket(s):                          1
Stepping:                           9
CPU(s) scaling MHz:                 19%
CPU max MHz:                        4200.0000
CPU min MHz:                        800.0000
BogoMIPS:                           7599.80
Flags:                              fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx pdpe1gb rdtscp lm constant_tsc art arch_perfmon pebs bts rep_good nopl xtopology nonstop_tsc cpuid aperfmperf pni pclmulqdq dtes64 monitor ds_cpl vmx est tm2 ssse3 sdbg fma cx16 xtpr pdcm pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand lahf_lm abm 3dnowprefetch cpuid_fault epb invpcid_single pti ibrs ibpb stibp tpr_shadow vnmi flexpriority ept vpid ept_ad fsgsbase tsc_adjust bmi1 hle avx2 smep bmi2 erms invpcid rtm mpx rdseed adx smap clflushopt intel_pt xsaveopt xsavec xgetbv1 xsaves dtherm ida arat pln pts hwp hwp_notify hwp_act_window hwp_epp
Virtualization:                     VT-x
L1d cache:                          128 KiB (4 instances)
L1i cache:                          128 KiB (4 instances)
L2 cache:                           1 MiB (4 instances)
L3 cache:                           6 MiB (1 instance)
NUMA node(s):                       1
NUMA node0 CPU(s):                  0-3
Vulnerability Gather data sampling: Vulnerable: No microcode
Vulnerability Itlb multihit:        KVM: Mitigation: VMX disabled
Vulnerability L1tf:                 Mitigation; PTE Inversion; VMX conditional cache flushes, SMT disabled
Vulnerability Mds:                  Vulnerable: Clear CPU buffers attempted, no microcode; SMT disabled
Vulnerability Meltdown:             Mitigation; PTI
Vulnerability Mmio stale data:      Vulnerable: Clear CPU buffers attempted, no microcode; SMT disabled
Vulnerability Retbleed:             Mitigation; IBRS
Vulnerability Spec rstack overflow: Not affected
Vulnerability Spec store bypass:    Vulnerable
Vulnerability Spectre v1:           Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:           Mitigation; IBRS, IBPB conditional, STIBP disabled, RSB filling, PBRSB-eIBRS Not affected
Vulnerability Srbds:                Vulnerable: No microcode
Vulnerability Tsx async abort:      Vulnerable: Clear CPU buffers attempted, no microcode; SMT disabled
$ uname -a

Linux ml 6.1.0-13-amd64 #1 SMP PREEMPT_DYNAMIC Debian 6.1.55-1 (2023-09-29) x86_64 GNU/Linux
$ python3 --version
Python 3.9.18

$ make --version
GNU Make 4.3
Built for x86_64-pc-linux-gnu
Copyright (C) 1988-2020 Free Software Foundation, Inc.
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html>
This is free software: you are free to change and redistribute it.
There is NO WARRANTY, to the extent permitted by law.

$ g++ --version
g++ (Debian 12.2.0-14) 12.2.0
Copyright (C) 2022 Free Software Foundation, Inc.
This is free software; see the source for copying conditions.  There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Tue_Aug_15_22:02:13_PDT_2023
Cuda compilation tools, release 12.2, V12.2.140
Build cuda_12.2.r12.2/compiler.33191640_0

$ nvidia-smi
Wed Oct 25 05:57:04 2023
+---------------------------------------------------------------------------------------+
| NVIDIA-SMI 545.23.06              Driver Version: 545.23.06    CUDA Version: 12.3     |
|-----------------------------------------+----------------------+----------------------+
| GPU  Name                 Persistence-M | Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp   Perf          Pwr:Usage/Cap |         Memory-Usage | GPU-Util  Compute M. |
|                                         |                      |               MIG M. |
|=========================================+======================+======================|
|   0  NVIDIA GeForce RTX 3070        On  | 00000000:02:00.0 Off |                  N/A |
|  0%   43C    P8              22W / 220W |      2MiB /  8192MiB |      0%      Default |
|                                         |                      |                  N/A |
+-----------------------------------------+----------------------+----------------------+
|   1  NVIDIA GeForce RTX 3070        On  | 00000000:03:00.0 Off |                  N/A |
|  0%   45C    P8              15W / 220W |      2MiB /  8192MiB |      0%      Default |
|                                         |                      |                  N/A |
+-----------------------------------------+----------------------+----------------------+
|   2  NVIDIA GeForce RTX 3070        On  | 00000000:04:00.0 Off |                  N/A |
|  0%   40C    P8              20W / 220W |      2MiB /  8192MiB |      0%      Default |
|                                         |                      |                  N/A |
+-----------------------------------------+----------------------+----------------------+
[...]

Failure Information (for bugs)

The issue seems to be unrelated to the actual model as well as its size. I'm observing this issue with llama models ranging from 7B to 70B parameters. It almost doesn't depend on the choice of -ngl as the model is producing broken output for any value larger than 0. Context size -c, generated tokens -n, --no-mmap, -nommq don't resolve the issue either.

Steps to Reproduce

Please provide detailed steps for reproducing the issue. We are not sitting in front of your screen, so the more detail the better.

  1. Get code

    git clone https://github.com/ggerganov/llama.cpp
cd llama.cpp

  2. Build with CUDA support

    make LLAMA_CUBLAS=1

  3. Get model in GGUF format e.g. huggingface.co/TheBloke/Llama-2-7B-GGUF)

  4. Query model

    CUDA_VISIBLE_DEVICES=0,1 ./main -ngl 99 -m ../LLM_stack/models/llama-2-7b.Q5_K_M.gguf --color -c 1500 --temp 0.01 -p "Why is the sky blue? Answer for a 5 year old child." -n 100

Failure Logs

Verbose console output for inference of llama-2 7B: output.log

Make log: make.log

ggerganov commented 10 months ago

Did it work before? If it did, can you bisect where it stopped working? Can you check if going back before 2b4ea35e56792064598e922e46d081e02bc96b94 fixes it?

nih23 commented 10 months ago

Yes! I just tested different commits to narrow down the issue: Multi-gpu inference has worked fine even on 8 GPUs until (including) 8b428c9bc84be6887d904600d1298b28baffd552. It seems that from 111163e2463171891680feed94371eb9becd9817 something has broken (@JohannesGaessler).

JohannesGaessler commented 10 months ago

I cannot reproduce the issue using 3x P40. Are you running llama.cpp inside a virtual machine or WSL?

nih23 commented 10 months ago

Nope, the linux system and llama.cpp are directly accessing the hardware, i.e. no virtualisation is involved.

kotee4ko commented 10 months ago

The thing is that in ggml_op_mul_mat we using cudaMemcpyD2Async to put data from active >1 gpu to main gpu, but this is possible only when gpus has cross-dma feature, which is true only for large bar pci systems.

The solution is to use hip/cudaMemcpyDtoDAsync in a loop to fill slices in dst from data in src0 -- this allow to bypass crashes on devices without p2p access.

But there are second sort of bug present. I can't say, yet, if it is AMD specific.

https://github.com/ggerganov/ggml/issues/590

kotee4ko commented 10 months ago

Maybe, anybody could explain what are we expecting to achive in dst and in which form? I think I can fix system code and make it right on both cuda/hip device/devices with/without p2p.

But I just cant understand math. Especially about src0 transposition -- call to ggml_is_transposed(src0) return False!

@ggerganov @JohannesGaessler @slaren @FSSRepo

Upd: When op() is called, and control flow reach ggml_cuda_op_mul_mat_cublas() the next thing is taking place:

On solo gpu: Convert src0 and src1 to f16, mul using hip/cublasGemmEx, convert to f32, ret.

On multi-gpu: Almost same, but mul with hip/cublasSgemm, and return WITHOUT convertion to f32.

Maybe, this could lead to sigsegv later on memcpy2d.

mgolub2 commented 9 months ago

I’m also having this issue with 2x 4090s - it actually corrupts the model files when I use 2 GPUs. Both work fine by themselves using CUDA_VISIBLE_DEVICES , and both pass gpu_burn for an hour without issue too.

Tried CUDA 12.3, 12.1, rocky linux, and ubuntu.

dji-transpire commented 9 months ago

The STRANGEST part is that it works beautifully on my old box a dual Xeon DELL 7610 with two 1080TI and one M6000 but it produces only garbage on my newly built box an ASUS X99-e WS build with two 3090 24GB founders edition.

I copied the source and recompiled with the same make LLAMA_CUBLAS=1 and on the DELL with older 1080TI and even older M6000 24GB the 13B Llama 2 produces nice output, pretty decent speed, but on the ASUS with two 3090 it produces garbage. It works if I take one of the 3090 out, but what is the point, I want to use the Q4 70B model.

wookayin commented 9 months ago

@dji-transpire Can you check the versions for CUDA, CUDNN, CUBLAS, NVIDIA driver versions, or any relevant SDKs -- were they the same? The model (generation) of GPU are already different so this also might be one factor even if everything else is the same.

dji-transpire commented 9 months ago

Thanks!!! You nailed it! The old box is running the 535 driver, the new box runs the latest 545 driver.

Downgrading nvidia-dkms nvidia-utils and lib32-nividia--utils to 535 and putting these on the IgnorePkg list solved the issue. Now both 3090 founders edition cards play nicely with LLAMA 13B and Q4 70B.

So: Be careful with the 545 version of the Nvidia driver and multiple GPUs????

peteygao commented 9 months ago

@dji-transpire Also running into the same issue with 3x 1080Ti, running driver version 545.29.08. Which exact version of 535 did you revert to? Was it to the latest 535.129.03?

wookayin commented 9 months ago

Note: A workaround for this bug is to use the CMake flag -DLLAMA_CUDA_PEER_MAX_BATCH_SIZE=0 when building llama.cpp, as done in jmorganca/ollama#1261. Or more simply:

make LLAMA_CUBLAS=1 LLAMA_CUDA_PEER_MAX_BATCH_SIZE=0

This will disable CUDA peer access completely and produce correct output when multiple GPUs are used.

Alumniminium commented 8 months ago

is there a way to make this work in textgeneration-webui without downgrading nvidia drivers?

morphles commented 5 months ago

Any news on this? For dual 7900 XTX I'm still getting garbage with hipBLAS build, regardless of model. But on single card it works. I tried the -DLLAMA_CUDA_PEER_MAX_BATCH_SIZE=0 option, but as it's CUDA did not have high hopes for it, and it did not help. Is there similar var for HIP maybe?

slaren commented 5 months ago

https://rocm.docs.amd.com/projects/radeon/en/latest/docs/limitations.html

morphles commented 5 months ago

@slaren yeah I know that, and I have no hope of it being fixed on AMD side soon, so I have very little hope in using pytorch with dual cards. Yet llama.cpp is much much better imo :) and flexible. I already have them working via vulkan, just mixtral on vk is still missing, but I know 0cc4m is working on it. But even without it I think llama.cpp already does some "manual workarounds" for what underlying libs do not provide, thus it that one 8x is on CPU other via chipset problem that I have, can be worked around via some slower "manual" data copying would still be nice :) . In any case for now vulkan seems like my best bet, so I'll be waiting for updates from 0cc4m :)

slaren commented 5 months ago

Somebody with access to dual 7900 XTX would need to diagnose the issue. AFAIK nobody who is working on the CUDA/HIP backend at the moment has access to this hardware.

morphles commented 5 months ago

Yeah, understandable :) for now I'm mostly happy with vulkan, and when mixstral is supported, I think I'll have basically no need for HIP build. Still if this somehow progresses, will also be nice to know. Thanks!

slaren commented 5 months ago

Can you test if it works with this change? (do not use -sm row).

diff --git a/ggml-cuda.cu b/ggml-cuda.cu
index 04c6f5d0..06af740e 100644
--- a/ggml-cuda.cu
+++ b/ggml-cuda.cu
@@ -797,7 +797,7 @@ static ggml_backend_buffer_i ggml_backend_cuda_buffer_interface = {
     /* .init_tensor     = */ ggml_backend_cuda_buffer_init_tensor,
     /* .set_tensor      = */ ggml_backend_cuda_buffer_set_tensor,
     /* .get_tensor      = */ ggml_backend_cuda_buffer_get_tensor,
-    /* .cpy_tensor      = */ ggml_backend_cuda_buffer_cpy_tensor,
+    /* .cpy_tensor      = */ NULL,//ggml_backend_cuda_buffer_cpy_tensor,
     /* .clear           = */ ggml_backend_cuda_buffer_clear,
     /* .reset           = */ NULL,
 };
@@ -11584,7 +11584,7 @@ static ggml_backend_i ggml_backend_cuda_interface = {
     /* .get_default_buffer_type = */ ggml_backend_cuda_get_default_buffer_type,
     /* .set_tensor_async        = */ ggml_backend_cuda_set_tensor_async,
     /* .get_tensor_async        = */ ggml_backend_cuda_get_tensor_async,
-    /* .cpy_tensor_async        = */ ggml_backend_cuda_cpy_tensor_async,
+    /* .cpy_tensor_async        = */ NULL,//ggml_backend_cuda_cpy_tensor_async,
     /* .synchronize             = */ ggml_backend_cuda_synchronize,
     /* .graph_plan_create       = */ NULL,
     /* .graph_plan_free         = */ NULL,
@@ -11592,10 +11592,10 @@ static ggml_backend_i ggml_backend_cuda_interface = {
     /* .graph_compute           = */ ggml_backend_cuda_graph_compute,
     /* .supports_op             = */ ggml_backend_cuda_supports_op,
     /* .offload_op              = */ ggml_backend_cuda_offload_op,
-    /* .event_new               = */ ggml_backend_cuda_event_new,
-    /* .event_free              = */ ggml_backend_cuda_event_free,
-    /* .event_record            = */ ggml_backend_cuda_event_record,
-    /* .event_wait              = */ ggml_backend_cuda_event_wait,
+    /* .event_new               = */ NULL,//ggml_backend_cuda_event_new,
+    /* .event_free              = */ NULL,//ggml_backend_cuda_event_free,
+    /* .event_record            = */ NULL,//ggml_backend_cuda_event_record,
+    /* .event_wait              = */ NULL,//ggml_backend_cuda_event_wait,
     /* .event_synchronize       = */ ggml_backend_cuda_event_synchronize,
 };
morphles commented 5 months ago

@slaren oh wow! Rebuilt on fresh checkout with your patch, and so far I think it works, just tested with single chat with one character in SillyTavern and it seems to be generating sensible stuff (as much as one can expect from model at this time :) ). Tested on couple models, command-r Q6 and noromaid mixtral Q4_K_M. I'll try some more stuff later today, but I think you have here a winning patch! :+1:

morphles commented 5 months ago

Ok testing some more generations, using mixtral, all seems to be working fine! Huge thanks @slaren !