Code

Nomic Vulkan backend#4456

Code

Merged

cebtenzzre merged 155 commits intomasterggerganov/llama.cpp:masterfrom ceb/nomic-vulkanggerganov/llama.cpp:ceb/nomic-vulkanCopy head branch name to clipboardJan 29, 2024

+4,271-19

Commits

Commits on Oct 5, 2023

• Nomic vulkan backend licensed under the Software for Open Models License (SOM), version 1.0.

  

  niansa

  authored and

  cebtenzzre

  committed
  ba15dfd

• Remove warning which fails on windows.

  

  manyoso

  authored and

  cebtenzzre

  committed
  48a45ea

• remove dynamic deps from kompute build

  Show description for 8563fa0
  

  apage43

  authored and

  cebtenzzre

  committed
  8563fa0

• Switch to a dynamic dispatch table instead of linking hard against libvulkan.

  

  manyoso

  authored and

  cebtenzzre

  committed
  45c8778

• Completely revamp how we do object management with the vulkan backend and

  Show description for b7e2e69
  

  manyoso

  authored and

  cebtenzzre

  committed
  b7e2e69

• Make kompute actually include external SDK headers when requested

  

  apage43

  authored and

  cebtenzzre

  committed
  beee572

• Throw an exception when allocation fails for vulkan.

  

  manyoso

  authored and

  cebtenzzre

  committed
  68cf1df

• vulkan: disambiguate gpus with the same name

  

  apage43

  authored and

  cebtenzzre

  committed
  8bea719

• Don't try and install kompute artifacts.

  

  manyoso

  authored and

  cebtenzzre

  committed
  bd5f639

• Sync from device back to host at begin of new prompt.

  

  manyoso

  authored and

  cebtenzzre

  committed
  4ed25b2

• Only use vulkan with known quant that work.

  

  manyoso

  authored and

  cebtenzzre

  committed
  68aca6b

• Set the singleton to nullptr here.

  

  manyoso

  authored and

  cebtenzzre

  committed
  addac25

• Don't crash on available devices if we can't even create an instance.

  

  manyoso

  authored and

  cebtenzzre

  committed
  2c24d67

Commits on Nov 3, 2023

• Support for gguf.

  

  manyoso

  authored and

  cebtenzzre

  committed
  1b1416d

• kompute : don't fail build because of -Warray-bounds

  Show description for 6b6c73a
  cebtenzzre

  committed
  6b6c73a

• Upload immediately to device.

  

  manyoso

  authored and

  cebtenzzre

  committed
  9e4f8b4

• Add a common boilerplate code via include and elim copy pasta

  

  manyoso

  authored and

  cebtenzzre

  committed
  77135a3

• Consolidate code for mat x vec kernels and use subgroups more extensively.

  

  manyoso

  authored and

  cebtenzzre

  committed
  93306f1

• Move the subgroups and printf into common.

  

  manyoso

  authored and

  cebtenzzre

  committed
  601905e

• Minor cleanup.

  

  manyoso

  authored and

  cebtenzzre

  committed
  5509f74

• Refactor getrows to use common code and get ready for q6_k.

  

  manyoso

  authored and

  cebtenzzre

  committed
  4b223ec

• Add q6_k getrows and mul*vec kernel.

  

  manyoso

  authored and

  cebtenzzre

  committed
  f1c9bc1

• Fix offset into the qh and now we have working vulkan accelerated for gguff'd llama.

  

  manyoso

  authored and

  cebtenzzre

  committed
  06d4b21

• Fixes for norm.

  

  manyoso

  authored and

  cebtenzzre

  committed
  32289aa

• Fixup the upstream CMakelists.txt so we can build just llama.cpp with our branch.

  

  manyoso

  authored and

  cebtenzzre

  committed
  6ac3975

• Change this back to be in agreement with metal and our previous softmax kernel.

  

  manyoso

  authored and

  cebtenzzre

  committed
  de589ce

• Fixes for subgroup size to bring AMD and NVIDIA inline with eachother for all kernels.

  

  manyoso

  authored and

  cebtenzzre

  committed
  bc4b5ed

• kompute : only try to use Vulkan for LLaMA itself

  cebtenzzre

  committed
  24a4a59

• kompute : remove Q6_K from list of supported quant types

  cebtenzzre

  committed
  3d850db

• f16 mv broadcasting fix (gqa fix)

  

  apage43

  authored and

  cebtenzzre

  committed
  9db90cb

• q8 mat*vec

  

  apage43

  authored and

  cebtenzzre

  committed
  ff4212d

• vulkan: implement neox mode for rope

  

  apage43

  authored and

  cebtenzzre

  committed
  020b174

• falcon h2d + reenable vulkan

  

  apage43

  authored and

  cebtenzzre

  committed
  8564f79

• Delete TODO now that we have q8_0.

  

  manyoso

  authored and

  cebtenzzre

  committed
  09d83f0

• add mat*mat ops

  

  apage43

  authored and

  cebtenzzre

  committed
  f0cd38b

• misc vulkan cleanup

  Show description for 46385ee
  

  apage43

  authored and

  cebtenzzre

  committed
  46385ee

• perf: use bigger threadgroups in mm

  

  apage43

  authored and

  cebtenzzre

  committed
  3327d84

• use op param epsilon for norms

  

  apage43

  authored and

  cebtenzzre

  committed
  d5741c0

• q6k mm works

  

  apage43

  authored and

  cebtenzzre

  committed
  b78a94b

• rm commented dbg print

  

  apage43

  authored and

  cebtenzzre

  committed
  4809890

• q4_1 mat*mat

  

  apage43

  authored and

  cebtenzzre

  committed
  cd0257e

• clean up vulkan/cpu switch

  

  apage43

  authored and

  cebtenzzre

  committed
  8dc79ac

• attempted speedups

  

  apage43

  authored and

  cebtenzzre

  committed
  9bc52eb

• attempted speedups 2

  

  apage43

  authored and

  cebtenzzre

  committed
  c1fd645

• use mat*vec shaders for mat*mat

  Show description for cc05a60
  

  apage43

  authored and

  cebtenzzre

  committed
  cc05a60

• kompute : enable kp_logger and make it static (#8)

  cebtenzzre

  committed
  21841d3

• kompute : make scripts executable

  cebtenzzre

  committed
  cbc0d1a

• Don't try an allocation on a heap that is smaller than the size we require.

  

  manyoso

  authored and

  cebtenzzre

  committed
  8400015

• Remove unused push constant that was giving validation errors.

  

  manyoso

  authored and

  cebtenzzre

  committed
  752f7eb

• Lower the workgroup count for some shaders by providing a loop that processes

  Show description for 8d9efbf
  

  manyoso

  authored and

  cebtenzzre

  committed
  8d9efbf

• Fix synchronization problem for AMD Radeon with amdvlk driver or windows

  Show description for 74ddf0f
  

  manyoso

  authored and

  cebtenzzre

  committed
  74ddf0f

• vulkan : fix missing break in matmul selection (#9)

  cebtenzzre

  committed
  1c17010

• llama : decide to disable Vulkan before loading tensors (#7)

  cebtenzzre

  committed
  89b7127

• Scale the workgroup count down to allow correct generation for falcon with

  Show description for e006d37
  

  manyoso

  authored and

  cebtenzzre

  committed
  e006d37

• Revert the prompt processing on gpu for now.

  Show description for a5eb001
  

  manyoso

  authored and

  cebtenzzre

  committed
  a5eb001

• Remove this debug code.

  

  manyoso

  authored and

  cebtenzzre

  committed
  ffd0624

• llama : fix Vulkan whitelist (#11)

  cebtenzzre

  committed
  f88b198

Commits on Nov 8, 2023

• kompute : fix issues with debug layers

  cebtenzzre

  committed
  a8cac53

• fix build with external fmtlib (v10)

  Show description for c438c16
  

  cebtenzzre

  and

  ToKiNoBug

  committed
  c438c16

• Merge commit 'ec893798b7a2a803466cc8f063051499ec3d96f7' into HEAD

  cebtenzzre

  committed
  af00cca

Commits on Nov 23, 2023

• vulkan : replace ggml_diag_mask_inf with ggml_add (custom -inf mask)

  cebtenzzre

  committed
  71565eb

• vulkan : rope n_past is now KQ_pos, f16 rope kernel

  cebtenzzre

  committed
  84f7fc4

• vulkan : optimize workgroup sizes

  cebtenzzre

  committed
  39abedd

• Merge commit 'fcca0a700487999d52a525c96d6661e9f6a8703a' into nomic-vulkan

  cebtenzzre

  committed
  f194e1b

• vulkan : assert various kernel requirements

  cebtenzzre

  committed
  a934b2c

• Merge commit '469c9addef75893e6be12edda852d12e840bf064' into nomic-vulkan

  cebtenzzre

  committed
  2a41ba7

• vulkan : handle ggml_scale for n%8 != 0

  Show description for 6474fc8
  cebtenzzre

  committed
  6474fc8

• Merge commit 'e16b9fa4baa8a09c6619b116159830e898050942' into nomic-vulkan

  cebtenzzre

  committed
  fe26e6a

• mention skipped change

  cebtenzzre

  committed
  9c4dfd0

• merge fixup (e16b9fa)

  cebtenzzre

  committed
  02c3309

• Merge commit '4760e7cc0b68570d58f55e8dda469805d1759d0d~' into nomic-vulkan

  cebtenzzre

  committed
  1829f1d

• vulkan : implement YaRN RoPE scaling (#2268)

  Show description for 208cd52
  cebtenzzre

  committed
  208cd52

• Merge commit '4760e7cc0b68570d58f55e8dda469805d1759d0d' into nomic-vulkan

  cebtenzzre

  committed
  23f6d51

• vulkan : sync with "migrate to dynamic graphs"

  cebtenzzre

  committed
  a4bb9c5

• Merge remote-tracking branch 'upstream/master' into nomic-vulkan-redo

  cebtenzzre

  committed
  9ae88ba

Commits on Dec 13, 2023

• relicense Vulkan backend as MIT

  cebtenzzre

  committed
  56430c3

• rename ggml-vulkan -> ggml-kompute

  cebtenzzre

  committed
  3e09e12

• separate shaders from kompute itself

  cebtenzzre

  committed
  27631db

• Merge commit '81bc9214a389362010f7a57f4cbc30e5f83a2d28' into nomic-vulkan

  cebtenzzre

  committed
  747e1ea

• kompute : fix compile warnings

  cebtenzzre

  committed
  b906e12

• move kompute to a submodule

  cebtenzzre

  committed
  9af7f58

• remove script with unclear purpose

  cebtenzzre

  committed
  f7cb0a6

Commits on Dec 14, 2023

• ggml : restore 'static' specifiers

  cebtenzzre

  committed
  c8fd4ba

Commits on Dec 15, 2023

• refactor llama.cpp modifications

  cebtenzzre

  committed
  f58f581

• vulkan : fix free of stack addr in llama_buffer

  cebtenzzre

  committed
  2d2c76a

• kompute : always destroy Manager via the destructor

  cebtenzzre

  committed
  8072706

Commits on Jan 8, 2024

• kompute : fix -Wunused-private-field warnings from clang

  Show description for 44b1a97
  cebtenzzre

  committed
  44b1a97

• Merge commit 'bcc0eb4591bec5ec02fad3f2bdcb1b265052ea56' into ceb/nomic-vulkan

  cebtenzzre

  committed
  8b65f4c

• Merge commit '31f27758faf4a4bd08101a57c7ec3a473f771f86' into ceb/nomic-vulkan

  cebtenzzre

  committed
  3959283

Commits on Jan 10, 2024

• sync xxd commands with GPT4All llama.cpp.cmake

  cebtenzzre

  committed
  904c563

• Merge commit 'd232aca5a73b290e218a2e48b91023d5e994203f' into ceb/nomic-vulkan

  cebtenzzre

  committed
  ae6d682

• Merge branch 'master' of https://github.com/ggerganov/llama.cpp into ceb/nomic-vulkan

  cebtenzzre

  committed
  3773e1a

Commits on Jan 24, 2024

• Merge commit 'e7e4df031b9e29d4b55a4e0b0295187f6b213db1' into HEAD

  cebtenzzre

  committed
  7c527eb

• kompute : initial attempt at ggml-backend v2 support

  cebtenzzre

  committed
  298d6ee

• fix assertion failure

  cebtenzzre

  committed
  5f660da

• attempt to get test-backend-ops working

  cebtenzzre

  committed
  070919d

• add sanity check and fix kompute teardown order

  cebtenzzre

  committed
  cad72e1

• kompute : ignore exceptions in ggml_vk_available_devices (#12)

  Show description for 76474a7
  cebtenzzre

  committed
  76474a7

• kompute : fix rope_f32 and scale ops (#5008)

  

  ggerganov

  authored and

  cebtenzzre

  committed
  d6bd471

• clean up old backend code

  cebtenzzre

  committed
  9431026

• actually fix this assertion

  cebtenzzre

  committed
  e9d5223

• sync op_rope_f16 with recent op_rope_f32 changes

  cebtenzzre

  committed
  729e1a4

• never try to evaluate an empty command buffer

  Show description for 0753073
  cebtenzzre

  committed
  0753073

• fix supported ops for kompute backend

  cebtenzzre

  committed
  2f6a279

• kompute : fix ggml_add kernel (#5027)

  

  ggerganov

  authored and

  cebtenzzre

  committed
  33e8d6a

• minor cleanup

  cebtenzzre

  committed
  cb9ceff

• kompute : fix get_rows dispatch -> 4 less failures

  cebtenzzre

  committed
  0899adf

• kompute : fix op_mul kernel -> 13 less test failures

  cebtenzzre

  committed
  08e23fd

• kompute : fix more dispatch ambiguity -> 12 less failures

  cebtenzzre

  committed
  2755ae3

• kompute : remove broken mulrow kernel -> 1 less test failure

  cebtenzzre

  committed
  0787b80

• fix q4_0/q4_1 mmv, 65 -> 49 failures

  cebtenzzre

  committed
  1a14099

• fix f16 mmv, 49 -> 41 failures

  cebtenzzre

  committed
  2b0f642

• test-backend-ops : add llama test

  cebtenzzre

  committed
  2852902

• test-backend-ops : test scale parameter of ggml_soft_max_ext

  cebtenzzre

  committed
  1450966

• kompute : support scale parameter of softmax

  cebtenzzre

  committed
  308f279

• test-backend-ops : test mask parameter of ggml_soft_max_ext

  cebtenzzre

  committed
  8bd38fe

• kompute : support mask parameter of softmax

  cebtenzzre

  committed
  df687b1

• test-backend-ops : test llama with different batch sizes

  cebtenzzre

  committed
  ebb5f7e

• test-backend-ops : increase max_nmse_err so Llama passes

  cebtenzzre

  committed
  ec68a96

Commits on Jan 25, 2024

• kompute : fix algorithm names

  cebtenzzre

  committed
  987335e

• kompute : fix q8_0 mmv, 41 -> 28 failures

  cebtenzzre

  committed
  f5ac635

• test-backend-ops : add Falcon test

  cebtenzzre

  committed
  1849b85

• test-backend-ops : test larger GELU range

  cebtenzzre

  committed
  6fc99a6

• kompute : fix op_gelu -> Falcon is working on AMDVLK

  cebtenzzre

  committed
  38d1f0c

• test-backend-ops : restore softmax tests

  cebtenzzre

  committed
  11b3050

• kompute : fix basic f16 get_rows, 28 -> 26 failures

  cebtenzzre

  committed
  de9fba0

• kompute : fix basic Q6_K get_rows, 26 -> 24 failures

  cebtenzzre

  committed
  445a373

• kompute : mark last few failing ops as unsupported

  cebtenzzre

  committed
  3fbf052

• test-backend-ops : make Falcon test faster with a smaller model

  cebtenzzre

  committed
  3915194

• kompute : remove unused immintrin.h #include

  cebtenzzre

  committed
  bc28704

• kompute : fix a -Wstrict-aliasing warning

  cebtenzzre

  committed
  91654ff

Commits on Jan 26, 2024

• kompute : remove unnecessary use_mmap=false

  cebtenzzre

  committed
  61a5cf8

• llama : revert unintended whitespace change

  cebtenzzre

  committed
  e6ce5f2

• Merge branch 'master' of https://github.com/ggerganov/llama.cpp into ceb/nomic-vulkan

  cebtenzzre

  committed
  aea8498

• test-backend-ops : comment out Llama and Falcon tests

  cebtenzzre

  committed
  2512799

• test-backend-ops : check all the ops in the test for support in the backends

  slaren

  committed
  8ca33de

• kompute : init device automatically and remove an unnecessary free

  cebtenzzre

  committed
  6af02b1

• ggml-kompute.h : remove anything that doesn't need to be public

  Show description for 2ff2d16
  cebtenzzre

  committed
  2ff2d16

• kompute : fix #includes

  cebtenzzre

  committed
  cdab404

• op_mul_mat_mat_f32.comp : fix missing final newline

  cebtenzzre

  committed
  454baeb

• editorconfig-checker : exclude .gitmodules

  cebtenzzre

  committed
  297fde5

• ci : initial attempt at testing Kompute backend

  cebtenzzre

  committed
  9132485

• main : remove ggml-kompute.h #include

  cebtenzzre

  committed
  57cecad

• kompute : adapt ggml-kompute API to be compatible with C

  cebtenzzre

  committed
  4b0c96a

Commits on Jan 27, 2024

• ci : attempt to fix Vulkan installer path

  cebtenzzre

  committed
  e6edd44

• ci : do not run tests for Kompute (no GPU)

  cebtenzzre

  committed
  050d450

• kompute : use llama_backend_init/llama_backend_free to manage device

  cebtenzzre

  committed
  5304625

Commits on Jan 29, 2024

• kompute : better device management

  cebtenzzre

  committed
  be7c055

• Merge branch 'master' of https://github.com/ggerganov/llama.cpp into ceb/nomic-vulkan

  cebtenzzre

  committed
  da1dc66

• kompute : fix merge issues

  cebtenzzre

  committed
  dc08e51

• kompute : remove llama_load_model_from_file_internal

  cebtenzzre

  committed
  7e11fe0

• vulkan : correctly fix use-after-free in ggml_vk_current_device

  Show description for b932cd7
  cebtenzzre

  committed
  b932cd7

• minor fixup

  cebtenzzre

  committed
  48db724

• fix trailing whitespace

  cebtenzzre

  committed
  1f98dff

• fix incorrect memcpy

  cebtenzzre

  committed
  2998211