https://github.com/ggerganov/llama.cpp/pull/1827 Skip to content Toggle navigation Sign up * Product + Actions Automate any workflow + Packages Host and manage packages + Security Find and fix vulnerabilities + Codespaces Instant dev environments + Copilot Write better code with AI + Code review Manage code changes + Issues Plan and track work + Discussions Collaborate outside of code Explore + All features + Documentation + GitHub Skills + Blog * Solutions For + Enterprise + Teams + Startups + Education By Solution + CI/CD & Automation + DevOps + DevSecOps Case Studies + Customer Stories + Resources * Open Source + GitHub Sponsors Fund open source developers + The ReadME Project GitHub community articles Repositories + Topics + Trending + Collections * Pricing [ ] * # In this repository All GitHub | Jump to | * No suggested jump to results * # In this repository All GitHub | Jump to | * # In this user All GitHub | Jump to | * # In this repository All GitHub | Jump to | Sign in Sign up You signed in with another tab or window. Reload to refresh your session. You signed out in another tab or window. Reload to refresh your session. You switched accounts on another tab or window. Reload to refresh your session. {{ message }} ggerganov / llama.cpp Public * Notifications * Fork 4.4k * Star 31.3k * Code * Issues 296 * Pull requests 70 * Discussions * Actions * Projects 4 * Wiki * Security * Insights More * Code * Issues * Pull requests * Discussions * Actions * Projects * Wiki * Security * Insights New issue Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community. Pick a username [ ] Email Address [ ] Password [ ] [ ] Sign up for GitHub By clicking "Sign up for GitHub", you agree to our terms of service and privacy statement. We'll occasionally send you account related emails. Already on GitHub? Sign in to your account Jump to bottom CUDA full GPU acceleration, KV cache in VRAM #1827 Open JohannesGaessler wants to merge 17 commits into ggerganov:master base: master Choose a base branch [ ] Branches Tags Could not load branches Branch not found: {{ refName }} {{ refName }} default Could not load tags Nothing to show {{ refName }} default Are you sure you want to change the base? Some commits from the old base branch may be removed from the timeline, and old review comments may become outdated. Change base from JohannesGaessler:cuda-full-gpu-2 Open CUDA full GPU acceleration, KV cache in VRAM #1827 JohannesGaessler wants to merge 17 commits into ggerganov:master from JohannesGaessler:cuda-full-gpu-2 +798 -119 Conversation 62 Commits 17 Checks 21 Files changed 11 Conversation This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters. Learn more about bidirectional Unicode characters Show hidden characters JohannesGaessler Copy link Collaborator @JohannesGaessler JohannesGaessler commented Jun 12, 2023 * edited This PR adds GPU acceleration for all remaining ggml tensors that didn't yet have it. Especially for long generations this makes a large difference because the KV cache is still CPU only on master and gets larger as the context fills up. Prompt processing is also significantly faster because the large batch size allows the more effective use of GPUs. For the following performance numbers PP is prompt processing, and TG 128/1024 are the generation of 128/1024 tokens with an empty prompt: GPU Test Model t/s master t/s PR Speedup RTX 3090 PP 7b q4_0 452 990 2.19 RTX 3090 PP 13b q4_0 318 645 2.03 RTX 3090 PP 33b q4_0 155 292 1.88 RTX 3090 TG 128 7b q4_0 45.66 62.66 1.37 RTX 3090 TG 128 13b q4_0 29.94 38.21 1.28 RTX 3090 TG 128 33b q4_0 14.89 17.61 1.18 RTX 3090 TG 1024 7b q4_0 31.08 56.34 1.81 RTX 3090 TG 1024 13b q4_0 20.16 35.20 1.75 RTX 3090 TG 1024 33b q4_0 10.37 16.49 1.59 Note that I was only using a single thread for the PR since multiple threads have no benefit when all computations are on the GPU but they still add overhead. I added CUDA kernels for scale, cpy, diag_mask_inf, and soft_max. I also added two special kernels for doing matrix vector multiplication with permuted or not contiguous inputs; they are used in conjunction with the KV cache. Changes to ggml.c: I added a utility function ggml_is_permuted. Things that are still to do: * [*] Fix VRAM memory leaks. * [*] Fix memory usage prints. * [*] Check performance for lower-end GPUs and add a --low-vram option if necessary. * [*] Check Windows performance and maybe disable features. * [*] General code cleanup. Sorry, something went wrong. 87 mirek190, slaren, Glavin001, Extraltodeus, avischiffmann, oobabooga, sequoiar, titoBouzout, franchb, transitive-bullshit, and 77 more reacted with thumbs up emoji 25 thushan, ggerganov, pabl-o-ce, bxh-io, dakl, theolivenbaum, mdawid, fredsadaghiani, melodysdreamj, DerekChia, and 15 more reacted with hooray emoji [?] 11 vihangd, mudler, Josua996, garanews, ShivBhosale, raphaelrk, AttilaBerczik, LiliumSancta, hugoabonizio, Pospelove, and alexl83 reacted with heart emoji 48 Green-Sky, Extraltodeus, avischiffmann, TheSeamau5, tmostak, megupta, sequoiar, andrewschreiber, ijt, DaseinPhaos, and 38 more reacted with rocket emoji All reactions * 87 reactions * 25 reactions * [?] 11 reactions * 48 reactions @JohannesGaessler JohannesGaessler added performance Speed related topics hardware Hardware related labels Jun 12, 2023 @KerfuffleV2 Copy link Collaborator KerfuffleV2 commented Jun 12, 2023 * edited Check performance for lower-end GPUs and add a --low-vram option if necessary. Unfortunately, I think that will be necessary. I have a 6GB 1060 - I can't even run main with 2048 context and -ngl 0 on a 65B model. Using -ngl 0 -c 1800 uses about 4.5GB VRAM even though the log output claims total VRAM used: 512 MB Since I upgraded my CPU, offloading to the GPU isn't really worthwhile anymore except for prompt ingestion. It does make a big difference there so it would be unfortunate if just using cuBLAS prompt input was impossible for full context on larger models. edit: By the way, let me know if/when further testing with this old hardware would be helpful. I can probably assist. 1 ruchir29 reacted with eyes emoji All reactions * 1 reaction Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 12, 2023 Really excited to try this! Amazing work! 3 oobabooga, lin72h, and avideci reacted with thumbs up emoji All reactions * 3 reactions Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 12, 2023 * edited Definitely seems faster and I can see it using 10% more GPU. Unfortunately when I ask it to return 512 tokens I'm getting an abort before the end of generation: CUDA error 1 at ggml-cuda.cu:1920: invalid argument Testing on H100 80GB. Ubuntu 20.04, gcc 9.4.0. CUDA toolkit 12.0.1. nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2023 NVIDIA Corporation Built on Fri_Jan__6_16:45:21_PST_2023 Cuda compilation tools, release 12.0, V12.0.140 Build cuda_12.0.r12.0/compiler.32267302_0 Compiled with: make clean && LLAMA_CUBLAS=1 make -j Command line arguments: ./main --color -ngl 60 --temp 0.7 --repeat_penalty 1.1 -n 512 --ignore-eos -m /workspace/WizardLM-30B-Uncensored.ggmlv3.q4_0.bin -p "USER: write a story about llamas\nASSISTANT:" Full output: [pytorch2] ubuntu@h100:/workspace/git/cuda_llama git:(cuda-full-gpu-2) $ ./main --color -ngl 60 --temp 0.7 --repeat_penalty 1.1 -n 512 --ignore-eos -m /workspace/WizardLM-30B-Uncensored.ggmlv3.q4_0.bin -p "USER: write a story about llamas\nASSISTANT:" main: build = 678 (0fe5ff2) main: seed = 1686600421 ggml_init_cublas: found 1 CUDA devices: Device 0: NVIDIA H100 PCIe llama.cpp: loading model from /workspace/WizardLM-30B-Uncensored.ggmlv3.q4_0.bin llama_model_load_internal: format = ggjt v3 (latest) llama_model_load_internal: n_vocab = 32001 llama_model_load_internal: n_ctx = 512 llama_model_load_internal: n_embd = 6656 llama_model_load_internal: n_mult = 256 llama_model_load_internal: n_head = 52 llama_model_load_internal: n_layer = 60 llama_model_load_internal: n_rot = 128 llama_model_load_internal: ftype = 2 (mostly Q4_0) llama_model_load_internal: n_ff = 17920 llama_model_load_internal: n_parts = 1 llama_model_load_internal: model size = 30B llama_model_load_internal: ggml ctx size = 0.13 MB llama_model_load_internal: using CUDA for GPU acceleration llama_model_load_internal: mem required = 2532.68 MB (+ 3124.00 MB per state) llama_model_load_internal: allocating batch_size x 1 MB = 512 MB VRAM for the scratch buffer llama_model_load_internal: offloading 60 layers to GPU llama_model_load_internal: total VRAM used: 17736 MB .................................................................................................... llama_init_from_file: kv self size = 780.00 MB system_info: n_threads = 13 / 26 | AVX = 1 | 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 | VSX = 0 | sampling: repeat_last_n = 64, repeat_penalty = 1.100000, presence_penalty = 0.000000, frequency_penalty = 0.000000, top_k = 40, tfs_z = 1.000000, top_p = 0.950000, typical_p = 1.000000, temp = 0.700000, mirostat = 0, mirostat_lr = 0.100000, mirostat_ent = 5.000000 generate: n_ctx = 512, n_batch = 512, n_predict = 512, n_keep = 0 USER: write a story about llamas\nASSISTANT:Once upon a time, in the high mountains of South America, there lived a group of llamas. They were known for their long necks and soft wool coats that kept them warm during the cold winter months. The llamas roamed the rugged terrain, grazing on grasses and plants that grew in the harsh environment. One day, a young llama named Lluvia decided to explore beyond her usual range. She wandered through fields of wildflowers and up steep ravines, until she came to a stream. As she drank from the cool water, she noticed a group of other llamas on the opposite bank. They were dressed in colorful saddles and had bags strapped to their backs. Curious, Lluvia crossed the stream and approached the group. "Hello," she said. "What are you doing here?" The leader of the group, a wise old llama named Tupac, replied, "We are on a journey to deliver supplies to a village at the base of the mountains. The people there need our help." Lluvia was impressed by the llamas' bravery and kindness. She asked if she could join them on their journey, and they welcomed her with open hearts. Together, the group set out down the mountain, carrying their heavy loads with ease. They passed through fields of crops and orchards, and Lluvia learned about the different plants and animals that lived in the region. Finally, they arrived at the village, where they were greeted by grateful locals who thanked them for their help. The llamas unloaded their supplies and spent the night in a cozy barn, surrounded by the warmth of the community. In the morning, Lluvia said goodbye to her new friends and returned to her own herd. But she never forgot the experience of helping others and the sense of purpose it had given her. From that day forward, she made a promise to herself to always be kind and compassionate, just like the wise old llama Tupac. And so, Lluvia went on to live a long and fulfilling life, inspiring others with her courage and generosity, and always remembering the lessons she had learned on her journey with the kind-CUDA error 1 at ggml-cuda.cu:1920: invalid argument [pytorch2] ubuntu@h100:/workspace/git/cuda_llama git:(cuda-full-gpu-2) $ Reducing to -n 400 resulted in successful completion. All reactions Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 12, 2023 @TheBloke I can reproduce the issue. It seems to have to do with -ngl 60; with those parameters some tensors are still on the CPU and the error happens when copying the data to the GPU. Setting -ngl 61 moves the non-repeating layers to the GPU as well and seems to work as intended. 2 roy09 and lin72h reacted with thumbs up emoji All reactions * 2 reactions Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 12, 2023 * edited Ahh OK, thanks. I did -ngl 60 because I thought that'd be the whole model. Yes that fixes it, thank you. Now tested fine up to 2000 tokens. EDIT: these benchmark figures were with --threads 13, which I learned later is very suboptimal. They at not representative of the real speed of this PR My first benchmark results are, testing with 400 tokens returned, results averaged over 4 runs H100 + Intel(R) Xeon(R) Platinum 8480+: * 7B q4_0 : + main branch: 38.23 ms/token + PR: 33.49 ms/token (best was 31.38ms) * 30B q4_0 : + main branch: 100.33ms/token + PR: 85.61 ms/token (best was 79.28ms) All reactions Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 12, 2023 Did you set --threads 1? You didn't do it in the command you posted previously and it should give a performance boost due to lower overhead. All reactions Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 12, 2023 * edited Oh wow! My apologies, I missed that in your OP. H100 + Intel(R) Xeon(R) Platinum 8480+. 500 tokens returned. * 7B q4_K_S : + main branch (--threads 13): 39.2ms/tok (25.51 tokens/s) + PR (--threads 1): 16.4 ms/token (60.97 tokens/s) + = 2.39x * 30B q4_K_S: + main branch (--threads 13): 92.69 ms/token (10.79 tokens/s) + PR (--threads 1): 53.71 ms/token (18.62 tokens/s) + = 1.73x This is really huge. This is the first time I've ever had GGML outperform GPTQ on any system. The max performance I've ever had from GPTQ 7B is 98 tokens/s, but that requires both a fast GPU (4090) and a top single-core CPU (i9-13900K). This system I'm on now has top GPU (H100), but middling single-core performing CPU, and it does around 45 tokens/s on 7B GPTQ 4bit. So on 7B on this server, GGML is now beating GPTQ 4bit by a healthy margin. I notice that llama.cpp is also pegged at 100% of one core, so I'm assuming a faster single-core CPU would likewise scale llama.cpp's figure. So soon I'll test on that 4090 + i9-13900K system and see if llama.cpp can get close to that 100 t/s mark. I just re-tested 30B GPTQ 4bit on this H100 system and got 20 t//s - still ahead of GGML, but only a tiny bit. This is really game-changing. Well done! 13 Glavin001, mattkanwisher, thushan, pabl-o-ce, th5, phantom32-0, nightlyworker, lin72h, jacobfriedman, AmineDjeghri, and 3 more reacted with hooray emoji All reactions * 13 reactions Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 12, 2023 I notice that llama.cpp is also pegged at 100% of one core, so I'm assuming a faster single-core CPU would likewise scale llama.cpp's figure. Currently the llama.cpp main thread always waits for the GPU computation to finish before arranging the next tensor. But it should be possible to change the logic to instead start preparing the next tensor immediately. Then I think CPU performance will be largely irrelevant. 1 lin72h reacted with thumbs up emoji All reactions * 1 reaction Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 12, 2023 I notice that llama.cpp is also pegged at 100% of one core, so I'm assuming a faster single-core CPU would likewise scale llama.cpp's figure. Currently the llama.cpp main thread always waits for the GPU computation to finish before arranging the next tensor. But it should be possible to change the logic to instead start preparing the next tensor immediately. Then I think CPU performance will be largely irrelevant. That would be incredible. It's always been both counter-intuitive and frustrating how GPTQ/pytorch inference is held back by CPU. Especially as it means that many high-end CPUs - the kinds you find on servers - actually do really poorly compared to gaming CPUs like the i9-13900K. I mentioned that with GPTQ 4bit my top 7B is 98 tokens/s with 4090 + i9-13900K. My normal figure on most servers I use is around 30 tokens /s, simply because server CPUs like AMD EPYC have really bad single-core performance. If you could completely de-couple from single core performance then GGML is going to leave pytorch/GPTQ in the dust for most home/ single-prompt users. All reactions Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 12, 2023 I don't think I'll be putting that in this PR though. My top priority is to make sure that the things that are already in it work. Also Python is literally 100x slower than C++ so I'm not sure whether the difference will be as large for llama.cpp. All reactions Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 12, 2023 I just tested on the 4090 + i9-13900K system and it's amazing: * 7B q4_K_S: + New llama.cpp performance: 109.29 tokens/s + AutoGPTQ CUDA 7B GPTQ 4bit: 98 tokens/s * 30B q4_K_S: + New PR llama.cpp performance: 29.11 tokens/s + AutoGPTQ CUDA 30B GPTQ 4bit: 35 tokens/s 2 gotzmann and milyiyo reacted with rocket emoji All reactions * 2 reactions Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 12, 2023 * edited I don't think I'll be putting that in this PR though. My top priority is to make sure that the things that are already in it work. Also Python is literally 100x slower than C++ so I'm not sure whether the difference will be as large for llama.cpp. Quite understand. Got to have something left to improve next week :) On the 4090 + i9-13900K it's still pegged at 100% CPU and GPU utilisation is at around 69%. So looks like there's still a decent margin there for further optimisations, if you can decouple it from CPU. Could it reach 150 tokens/s on 7B 4bit? 1 lin72h reacted with eyes emoji All reactions * 1 reaction Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 12, 2023 Could it reach 150 tokens/s on 7B 4bit? You could possibly find that out by compiling with LLAMA_GPROF and digging through the profiling data but I'll tell you right now that the largest remaining use of CPU on my system with an RTX 3090 was the sampling of the repetition penalty which has to be done sequentially and would not be affected at all by what I have in mind. So maybe I was wrong and single core performance will still make a difference. All reactions Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 12, 2023 * edited Fair enough. Well it's fantastic either way. All reactions Sorry, something went wrong. @LevanKvirkvelia Copy link LevanKvirkvelia commented Jun 13, 2023 * edited Hi, great PR, does this PR allows you to keep the cache after the generation is done? I think statefulness should speedup chats and libraries like Microsoft Guidance and Jsonformer All reactions Sorry, something went wrong. @cmp-nct Copy link cmp-nct commented Jun 13, 2023 Stunning work, can't wait testing it tomorrow. Given the benchmark results this is a huge step forward.. All reactions Sorry, something went wrong. @thushan Copy link thushan commented Jun 13, 2023 * edited This is huge news, it really whips the llama.cpp's ass on CUDA Early attempt this morning we're getting ~2.5-2.8x perf increase on 4090s and about 1.8-2x on 3090Ti. 8 kyushuadamu, geerlingguy, lin72h, jacobfriedman, rlapray, bermanboris, hugows, and houmii reacted with laugh emoji All reactions * 8 reactions Sorry, something went wrong. @ggerganov Copy link Owner ggerganov commented Jun 13, 2023 Congrats @JohannesGaessler! I'll yet be taking a detailed look in the PR, but I can already tell you have done a really good job based on the feedback. Well deserved for all your hard work I notice that llama.cpp is also pegged at 100% of one core, so I'm assuming a faster single-core CPU would likewise scale llama.cpp's figure. Currently the llama.cpp main thread always waits for the GPU computation to finish before arranging the next tensor. But it should be possible to change the logic to instead start preparing the next tensor immediately. Then I think CPU performance will be largely irrelevant. We will refactor and improve the CPU threading and synchronization logic of ggml soon 1 lin72h reacted with thumbs up emoji [?] 17 thushan, pabl-o-ce, deadprogram, Rafaelblsilva, nightlyworker, junaid33, Volrath50, PiotrDabkowski, Priestru, ericvolp12, and 7 more reacted with heart emoji All reactions * 1 reaction * [?] 17 reactions Sorry, something went wrong. @ggerganov ggerganov added the high priority Very important issue label Jun 13, 2023 @deadprogram Copy link Collaborator deadprogram commented Jun 13, 2023 Fantastic work @JohannesGaessler All reactions Sorry, something went wrong. @github-actions github-actions bot mentioned this pull request Jun 13, 2023 Hacker News Daily Top 30 @2023-06-13 meixger/hackernews-daily#268 Open @Priestru Copy link Priestru commented Jun 13, 2023 Would it support multi GPU? 3 lin72h, avideci, and vid reacted with thumbs up emoji 3 lin72h, foolsh, and curtisgray reacted with eyes emoji All reactions * 3 reactions * 3 reactions Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 13, 2023 Hi, great PR, does this PR allows you to keep the cache after the generation is done? I think statefulness should speedup chats and libraries like Microsoft Guidance and Jsonformer @LevanKvirkvelia I am not changing anything other than the location where the KC cache is stored. We will refactor and improve the CPU threading and synchronization logic of ggml soon @ggerganov What I'm talking about is something different. Currently cudaDeviceSynchronize() is called every time after a calculation is started. But it would be more efficient to instead synchronize at a later point when the results are actually needed which would allow the parallel rather than sequential use of CPU and GPU. Unless the threading PR implements this too. Would it support multi GPU? @Priestru It already does for the computationally most expensive operations. There is no multi GPU support for the KV cache specifically; before I add that I'll need to see if that would even be worthwhile in the first place. All reactions Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 13, 2023 @TheBloke Are you setting LLAMA_CUDA_DMMV_X (default 32) and LLAMA_CUDA_DMMV_Y (default 1) at compile time? These values determine how much data the GPU processes at once for the computationally most expensive operations and setting higher values is beneficial on fast GPUs (but make sure they are powers of 2). On my RTX 3090 setting LLAMA_CUDA_DMMV_X=64 LLAMA_CUDA_DMMV_Y=2 increases performance by 20%. All reactions Sorry, something went wrong. @KerfuffleV2 Copy link Collaborator KerfuffleV2 commented Jun 13, 2023 Just want to make sure you're aware this currently prevents using cuBLAS with large models on low VRAM GPUs like my 6GB 1060, even with no GPU offloading. I can't use full context (-c 2048) with -ngl 0 - it still runs out of VRAM. All reactions Sorry, something went wrong. @Dampfinchen Copy link Dampfinchen commented Jun 13, 2023 * edited Sadly with this PR, I get ggml-cuda.cu:1248: out of memory immediately after loading the 13B model. In the regular build, this was my speed (15 GPU layers for my RTX 2060, 6 threads, prompt of 1800 Tokens) Cold run: llama_print_timings: load time = 44221.63 ms llama_print_timings: sample time = 144.82 ms / 180 runs ( 0.80 ms per token) llama_print_timings: prompt eval time = 59045.43 ms / 1849 tokens ( 31.93 ms per token) llama_print_timings: eval time = 77643.51 ms / 179 runs ( 433.76 ms per token) llama_print_timings: total time = 142706.14 ms (I've used 13 layers on the new master build and this PR which uses around 3,4 GB VRAM) So yeah, while this is good work and a big step forward for GGML in general, the biggest strength of GGML compared to GPTQ is that you can run large models on GPUs with insufficient VRAM at decent speed is suffering drastically with this approach. So while it is great that a 33B model can now run entirely on a RTX 4090, it's not so great for us that want to run larger models exceeding our GPU's VRAM. 1 ruchir29 reacted with confused emoji All reactions * 1 reaction Sorry, something went wrong. @xueyuanl xueyuanl mentioned this pull request Jun 13, 2023 Daily Hacker News 13-06-2023 xueyuanl/daily-hackernews#1009 Open @EwoutH Copy link EwoutH commented Jun 13, 2023 Awesome work! Here's the link to the Reddit thread on r/LocalLLaMA for reference, a lot of people share their benchmarks and experiences on this PR. All reactions Sorry, something went wrong. 37 hidden items Load more... JohannesGaessler added 4 commits June 13, 2023 17:02 @JohannesGaessler Removed obsolete code, fixed multi GPU 8e3057b @JohannesGaessler Free KV cache CUDA buffers upon deletion ed65874 @JohannesGaessler Free CUDA scratch buffer upon llama_model deletion 20e76a0 @JohannesGaessler Fixed incorrect index when going out of context e8528d4 This was referenced Jun 13, 2023 Your Issue Title jacky1234/blogPages#28 Open Your Issue Title jacky1234/blogPages#29 Open Your Issue Title jacky1234/blogPages#30 Open @TheBloke Copy link Contributor TheBloke commented Jun 13, 2023 * edited Here's my results offloading 65B q4_K_M on an A10 + Intel(R) Xeon(R) Platinum 8358 CPU @ 2.60GHz branch: master -m guanaco-65B.ggmlv3.q4_K_M.bin -ngl 44 -t 14 -n 200 llama_print_timings: eval time = 61957.43 ms / 199 runs ( 311.34 ms per token) llama_print_timings: eval time = 63633.71 ms / 199 runs ( 319.77 ms per token) llama_print_timings: eval time = 63416.40 ms / 199 runs ( 318.68 ms per token) branch: cuda-full-gpu-2 -m guanaco-65B.ggmlv3.q4_K_M.bin -ngl 41 -t 14 -n 200 llama_print_timings: eval time = 63243.73 ms / 199 runs ( 317.81 ms per token) llama_print_timings: eval time = 64970.85 ms / 199 runs ( 326.49 ms per token) llama_print_timings: eval time = 63186.68 ms / 199 runs ( 317.52 ms per token) branch: cuda-full-gpu-2-no-k -m guanaco-65B.ggmlv3.q4_K_M.bin -ngl 42 -t 14 -n 200 llama_print_timings: eval time = 63346.43 ms / 199 runs ( 318.32 ms per token) llama_print_timings: eval time = 62468.10 ms / 199 runs ( 313.91 ms per token) llama_print_timings: eval time = 62256.63 ms / 199 runs ( 312.85 ms per token) branch: cuda-full-gpu-2-no-v -m guanaco-65B.ggmlv3.q4_K_M.bin -ngl 42 -t 14 -n 200 llama_print_timings: eval time = 61643.87 ms / 199 runs ( 309.77 ms per token) llama_print_timings: eval time = 62827.90 ms / 199 runs ( 315.72 ms per token) llama_print_timings: eval time = 62729.78 ms / 199 runs ( 315.23 ms per token) branch: cuda-full-gpu-2-no-kv -m guanaco-65B.ggmlv3.q4_K_M.bin -ngl 44 -t 14 -n 200 llama_print_timings: eval time = 62337.81 ms / 199 runs ( 313.26 ms per token) llama_print_timings: eval time = 62527.46 ms / 199 runs ( 314.21 ms per token) llama_print_timings: eval time = 62564.66 ms / 199 runs ( 314.40 ms per token) (I preceded this by benchmarking --threads 8 through --threads 16 on master, and established that --threads 14 was quickest) [?] 1 pabl-o-ce reacted with heart emoji All reactions * [?] 1 reaction Sorry, something went wrong. This was referenced Jun 13, 2023 [HackerNews Top 10]-06-14 00-02 jacky1234/blogPages#31 Open HackerNews Top 10 @23yyy-06-14 jacky1234/blogPages#32 Open HackerNews Top 10 @2023-06-14 jacky1234/blogPages#33 Open slaren slaren reviewed Jun 13, 2023 View reviewed changes ggml-cuda.cu Outdated if (nb0 == ts && nb1 == ts*ne0/bs) { return cudaMemcpyAsync(dst_char, x, i1_diff*nb1, cudaMemcpyHostToDevice, stream); return cudaMemcpyAsync(dst_ptr, x, i1_diff*nb1, cudaMemcpyHostToDevice, stream); Copy link Collaborator @slaren slaren Jun 13, 2023 There was a problem hiding this comment. Choose a reason for hiding this comment The reason will be displayed to describe this comment to others. Learn more. [Choose a reason] Hide comment I think this path is never hit currently, but I believe the parameter should be kind rather than cudaMemcpyHostToDevice here. Sorry, something went wrong. All reactions Copy link Collaborator @slaren slaren Jun 13, 2023 There was a problem hiding this comment. Choose a reason for hiding this comment The reason will be displayed to describe this comment to others. Learn more. [Choose a reason] Hide comment On minor note, I was really confused by the name of this function at first (sounds like "tensor 2D" rather than "tensor to device"). Sorry, something went wrong. All reactions Copy link Collaborator Author @JohannesGaessler JohannesGaessler Jun 13, 2023 * edited There was a problem hiding this comment. Choose a reason for hiding this comment The reason will be displayed to describe this comment to others. Learn more. [Choose a reason] Hide comment I agree that the name is currently bad but I didn't yet have a good idea for another name. Sorry, something went wrong. All reactions slaren slaren reviewed Jun 13, 2023 View reviewed changes ggml-cuda.cu Show resolved Hide resolved @JohannesGaessler VRAM KV cache based on -ngl, fixed info prints cc60183 @JohannesGaessler JohannesGaessler force-pushed the cuda-full-gpu-2 branch from 6967a44 to cc60183 Compare June 13, 2023 17:11 @gotzmann Copy link gotzmann commented Jun 13, 2023 * edited Wow, so now performance is on par with exllama? Got those numbers for 30B model from there turboderp/exllama#16 gpu architecture perf power -------------------------------------------------------------- H100 PCIe Hopper 34.1 t/s 300W 4090 Ada Lovelace 32.5 t/s 450W A6000 Ada Ada Lovelace 31.6 t/s 300W L40 Ada Lovelace 27.7 t/s 300W A100 SXM4 Ampere 25.8 t/s 400W 3090 Ampere 22.7 t/s 350W A6000 Ampere 21.4 t/s 300W A40 Ampere 21.4 t/s 300W A5000 Ampere 17.8 t/s 230W All reactions Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 13, 2023 * edited Hey Johannes, new commit seems to work really well. On the A10 I can now run 65B with the same number of layers (44) as I can in master, with the same performance. Did you see my earlier messages about the assert crash I get when I try to compile with LLAMA_CUDA_DMMV_Y to > 1? All reactions Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 13, 2023 I did see that the parameter is causing you trouble but I don't understand why it's happening. I find it highly unlikely that the problem was only introduced with this PR though. I can't really fix the bug unless I have a way to reproduce it; if you know how to use GNU Debugger that could give me some useful information but I suspect you don't. All reactions Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 13, 2023 You are correct on both counts! Yes, master is crashing in the same way. I've noticed that it's not crashing on the A10 system. So in total I have tested: * 4090 + CUDA 11.8 = crash with assert * H100 + CUDA 12.0.1 = crash with assert * A10 + CUDA 11.8 = no crash. The 4090 and H100 are both a later compute version than the A10; 8.9 and 9.0 respectively. Don't know if that could be a factor? I'm happy to run through gdb commands if you can describe them, but yes I have no experience of using it myself. All reactions Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 13, 2023 Can you tell me the line again for the assertion error using the latest commit that I pushed? It's no longer the same line due to changes in the code. All reactions Sorry, something went wrong. @TheBloke Copy link Contributor TheBloke commented Jun 13, 2023 Sure. It's: .GGML_ASSERT: ggml-cuda.cu:2221: nrows % GGML_CUDA_DMMV_Y == 0 In master it's: .GGML_ASSERT: ggml-cuda.cu:1740: nrows % GGML_CUDA_DMMV_Y == 0 All reactions Sorry, something went wrong. @JohannesGaessler Added a --low-vram option dba1452 @lhl Copy link lhl commented Jun 13, 2023 * edited Wow, so now performance is on par with exllama? Got those numbers for 30B model from there turboderp/exllama#16 Those numbers are a bit out of date. I think that the numbers linked from the exllama README might be a better guideline on performance, but it's going to depend a lot on your individual system where the bottlenecks are. As a point of reference, I just tested the latest pulls from each codebase. My system is a Ryzen 5950X w/ DDR4-3800 RAM, (dedicated no output) RTX 4090 on Arch Linux: * On llama-33b q4_0, the llama.cpp branch I get 17.62t/s on full context (run with -ngl 99 -n 2048 --ignore-eos)- the 4090 hits 290W and sustains about 50% utilization, one CPU core is pegged at 100% * With exllama, on a GPTQ 4-bit 128g, on 2048t I get 40.40t/s - my 4090 hits 440W and maintains 100% GPU utilization, CPU core hits (but doesn't stay at) 100% during the run. All reactions Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 13, 2023 @TheBloke The commands should be: make clean && make LLAMA_CUBLAS=1 LLAMA_CUDA_DMMV_X=64 LLAMA_CUDA_DMMV_Y=2 LLAMA_DEBUG=1 gdb -args b exit r up up up p nrows p *tensor To get out of GDB, press CTRL+D. All reactions Sorry, something went wrong. @JohannesGaessler Fixed Windows performance 51830ee @TheBloke Copy link Contributor TheBloke commented Jun 13, 2023 Thanks very much for the info. In case this is helpful: I learned that that the issue seems to occur only when the non-repeating layers are offloaded. Eg on 7B it happens at -ngl 33, with 13B at 41, and 33B at 61. Offloading fewer layers runs OK. Here's the gdb log: [pytorch2] ubuntu@h100:/workspace/git/llama_cuda_test/cuda-full-gpu-2 git:(cuda-full-gpu-2) $ make clean && make LLAMA_CUBLAS=1 LLAMA_CUDA_DMMV_X=64 LLAMA_CUDA_DMMV_Y=2 LLAMA_DEBUG=1 ; gdb -args ./main --threads 1 -ngl 33 -m /workspace/WizardLM-7B-uncensored.ggmlv3.q4_K_S.bin -p "once" I llama.cpp build info: I UNAME_S: Linux I UNAME_P: x86_64 I UNAME_M: x86_64 I CFLAGS: -I. -O3 -std=c11 -fPIC -DNDEBUG -Wall -Wextra -Wpedantic -Wcast-qual -Wdouble-promotion -Wshadow -Wstrict-prototypes -Wpointer-arith -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS I CXXFLAGS: -I. -I./examples -O3 -std=c++11 -fPIC -DNDEBUG -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS I LDFLAGS: I CC: cc (Ubuntu 9.4.0-1ubuntu1~20.04.1) 9.4.0 I CXX: g++ (Ubuntu 9.4.0-1ubuntu1~20.04.1) 9.4.0 rm -vf *.o main quantize quantize-stats perplexity embedding benchmark-matmult save-load-state server vdot build-info.h removed 'common.o' removed 'ggml-cuda.o' removed 'ggml.o' removed 'k_quants.o' removed 'llama.o' removed 'main' removed 'quantize' removed 'quantize-stats' removed 'perplexity' removed 'embedding' removed 'vdot' removed 'build-info.h' I llama.cpp build info: I UNAME_S: Linux I UNAME_P: x86_64 I UNAME_M: x86_64 I CFLAGS: -I. -O3 -std=c11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wdouble-promotion -Wshadow -Wstrict-prototypes -Wpointer-arith -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include I CXXFLAGS: -I. -I./examples -O3 -std=c++11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include I LDFLAGS: -g -lcublas -lculibos -lcudart -lcublasLt -lpthread -ldl -lrt -L/usr/local/cuda/lib64 -L/opt/cuda/lib64 -L/targets/x86_64-linux/lib I CC: cc (Ubuntu 9.4.0-1ubuntu1~20.04.1) 9.4.0 I CXX: g++ (Ubuntu 9.4.0-1ubuntu1~20.04.1) 9.4.0 cc -I. -O3 -std=c11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wdouble-promotion -Wshadow -Wstrict-prototypes -Wpointer-arith -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include -c ggml.c -o ggml.o g++ -I. -I./examples -O3 -std=c++11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include -c llama.cpp -o llama.o g++ -I. -I./examples -O3 -std=c++11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include -c examples/common.cpp -o common.o cc -I. -O3 -std=c11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wdouble-promotion -Wshadow -Wstrict-prototypes -Wpointer-arith -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include -c -o k_quants.o k_quants.c nvcc --forward-unknown-to-host-compiler -arch=native -DGGML_CUDA_DMMV_X=64 -DGGML_CUDA_DMMV_Y=2 -I. -I./examples -O3 -std=c++11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include -Wno-pedantic -c ggml-cuda.cu -o ggml-cuda.o nvcc warning : incompatible redefinition for option 'optimize', the last value of this option was used g++ -I. -I./examples -O3 -std=c++11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include examples/main/main.cpp ggml.o llama.o common.o k_quants.o ggml-cuda.o -o main -g -lcublas -lculibos -lcudart -lcublasLt -lpthread -ldl -lrt -L/usr/local/cuda/lib64 -L/opt/cuda/lib64 -L/targets/x86_64-linux/lib ==== Run ./main -h for help. ==== g++ -I. -I./examples -O3 -std=c++11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include examples/quantize/quantize.cpp ggml.o llama.o k_quants.o ggml-cuda.o -o quantize -g -lcublas -lculibos -lcudart -lcublasLt -lpthread -ldl -lrt -L/usr/local/cuda/lib64 -L/opt/cuda/lib64 -L/targets/x86_64-linux/lib g++ -I. -I./examples -O3 -std=c++11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include examples/quantize-stats/quantize-stats.cpp ggml.o llama.o k_quants.o ggml-cuda.o -o quantize-stats -g -lcublas -lculibos -lcudart -lcublasLt -lpthread -ldl -lrt -L/usr/local/cuda/lib64 -L/opt/cuda/lib64 -L/targets/x86_64-linux/lib g++ -I. -I./examples -O3 -std=c++11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include examples/perplexity/perplexity.cpp ggml.o llama.o common.o k_quants.o ggml-cuda.o -o perplexity -g -lcublas -lculibos -lcudart -lcublasLt -lpthread -ldl -lrt -L/usr/local/cuda/lib64 -L/opt/cuda/lib64 -L/targets/x86_64-linux/lib g++ -I. -I./examples -O3 -std=c++11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include examples/embedding/embedding.cpp ggml.o llama.o common.o k_quants.o ggml-cuda.o -o embedding -g -lcublas -lculibos -lcudart -lcublasLt -lpthread -ldl -lrt -L/usr/local/cuda/lib64 -L/opt/cuda/lib64 -L/targets/x86_64-linux/lib g++ -I. -I./examples -O3 -std=c++11 -fPIC -O0 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wno-multichar -pthread -march=native -mtune=native -DGGML_USE_K_QUANTS -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include -I/targets/x86_64-linux/include pocs/vdot/vdot.cpp ggml.o k_quants.o ggml-cuda.o -o vdot -g -lcublas -lculibos -lcudart -lcublasLt -lpthread -ldl -lrt -L/usr/local/cuda/lib64 -L/opt/cuda/lib64 -L/targets/x86_64-linux/lib GNU gdb (Ubuntu 9.2-0ubuntu1~20.04.1) 9.2 Copyright (C) 2020 Free Software Foundation, Inc. License GPLv3+: GNU GPL version 3 or later This is free software: you are free to change and redistribute it. There is NO WARRANTY, to the extent permitted by law. Type "show copying" and "show warranty" for details. This GDB was configured as "x86_64-linux-gnu". Type "show configuration" for configuration details. For bug reporting instructions, please see: . Find the GDB manual and other documentation resources online at: . For help, type "help". Type "apropos word" to search for commands related to "word"... Reading symbols from ./main... (gdb) b exit Breakpoint 1 at 0xa7a0 (gdb) r Starting program: /workspace/git/llama_cuda_test/cuda-full-gpu-2/main --threads 1 -ngl 33 -m /workspace/WizardLM-7B-uncensored.ggmlv3.q4_K_S.bin -p once [Thread debugging using libthread_db enabled] Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1". main: build = 680 (cc60183) main: seed = 1686688523 [New Thread 0x7ffdc6430000 (LWP 1002419)] ggml_init_cublas: found 1 CUDA devices: Device 0: NVIDIA H100 PCIe [New Thread 0x7ffdc5c2f000 (LWP 1002420)] llama.cpp: loading model from /workspace/WizardLM-7B-uncensored.ggmlv3.q4_K_S.bin llama_model_load_internal: format = ggjt v3 (latest) llama_model_load_internal: n_vocab = 32001 llama_model_load_internal: n_ctx = 512 llama_model_load_internal: n_embd = 4096 llama_model_load_internal: n_mult = 256 llama_model_load_internal: n_head = 32 llama_model_load_internal: n_layer = 32 llama_model_load_internal: n_rot = 128 llama_model_load_internal: ftype = 14 (mostly Q4_K - Small) llama_model_load_internal: n_ff = 11008 llama_model_load_internal: n_parts = 1 llama_model_load_internal: model size = 7B llama_model_load_internal: ggml ctx size = 0.07 MB llama_model_load_internal: using CUDA for GPU acceleration llama_model_load_internal: mem required = 1862.39 MB (+ 1026.00 MB per state) llama_model_load_internal: allocating batch_size x 1 MB = 512 MB VRAM for the scratch buffer llama_model_load_internal: offloading 32 repeating layers to GPU llama_model_load_internal: offloading non-repeating layers to GPU llama_model_load_internal: offloaded 33/35 layers to GPU llama_model_load_internal: total VRAM used: 4058 MB .GGML_ASSERT: ggml-cuda.cu:2221: nrows % GGML_CUDA_DMMV_Y == 0 Thread 1 "main" received signal SIGABRT, Aborted. __GI_raise (sig=sig@entry=6) at ../sysdeps/unix/sysv/linux/raise.c:50 50 ../sysdeps/unix/sysv/linux/raise.c: No such file or directory. (gdb) up #1 0x00007ffff0f53859 in __GI_abort () at abort.c:79 79 abort.c: No such file or directory. (gdb) up #2 0x000055555561f0f2 in ggml_cuda_transform_tensor (data=0x7ffc426aa380, tensor=0x7ffd31400220) at ggml-cuda.cu:2221 2221 GGML_ASSERT(nrows % GGML_CUDA_DMMV_Y == 0); (gdb) up #3 0x00005555555b22e9 in llama_model_loader::load_all_data (this=0x5555a3b9c220, progress_callback=0x5555555a529a <::_FUN(float, void *)>, progress_callback_user_data=0x7fffffff9670, lmlock=0x0) at llama.cpp:788 788 ggml_cuda_transform_tensor(lt.data, lt.ggml_tensor); (gdb) p nrows No symbol "nrows" in current context. (gdb) p *tensor No symbol "tensor" in current context. (gdb) PS. If debugging by proxy is too painful, you're welcome to SSH in directly if you send me an SSH public key. All reactions Sorry, something went wrong. @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 13, 2023 Access via SSH would be useful. I've sent you a friend request on Discord. All reactions Sorry, something went wrong. @JohannesGaessler Cleaned up code, added comments a47072b @JohannesGaessler JohannesGaessler marked this pull request as ready for review June 13, 2023 22:01 @JohannesGaessler Copy link Collaborator Author JohannesGaessler commented Jun 13, 2023 * edited Alright, I think this PR is now feature complete. I've added a --low-vram option that disables the CUDA scratch buffer. Performance on Windows is absolutely terrible but with the changes that I did it's now slightly less terrible given enough VRAM and it's at least not more terrible than it was before I added the VRAM scratch buffer. All reactions Sorry, something went wrong. Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment Reviewers @slaren slaren slaren left review comments At least 1 approving review is required to merge this pull request. Assignees No one assigned Labels hardware Hardware related high priority Very important issue performance Speed related topics Projects None yet Milestone No milestone Development Successfully merging this pull request may close these issues. None yet 17 participants @JohannesGaessler @KerfuffleV2 @TheBloke @LevanKvirkvelia @cmp-nct @thushan @ggerganov @deadprogram @Priestru @Dampfinchen @EwoutH @foolsh @lhl @shouyiwang @LoganDark @gotzmann @slaren Add this suggestion to a batch that can be applied as a single commit. This suggestion is invalid because no changes were made to the code. Suggestions cannot be applied while the pull request is closed. Suggestions cannot be applied while viewing a subset of changes. Only one suggestion per line can be applied in a batch. Add this suggestion to a batch that can be applied as a single commit. Applying suggestions on deleted lines is not supported. You must change the existing code in this line in order to create a valid suggestion. Outdated suggestions cannot be applied. This suggestion has been applied or marked resolved. Suggestions cannot be applied from pending reviews. Suggestions cannot be applied on multi-line comments. Suggestions cannot be applied while the pull request is queued to merge. Footer (c) 2023 GitHub, Inc. Footer navigation * Terms * Privacy * Security * Status * Docs * Contact GitHub * Pricing * API * Training * Blog * About You can't perform that action at this time.