[Web] Compatibility with PagedKVCache in WebGPU #16554
Merged
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
MasterJH5574
approved these changes
Feb 12, 2024
There was a problem hiding this comment.
LGTM. Thank you @CharlieFRuan and @rickzx for carrying through!
CharlieFRuan
added a commit
to mlc-ai/web-llm
that referenced
this pull request
Feb 13, 2024
PagedKVCache is introduced in MLC-LLM a while back to unite the interface for KVCache. This PR makes WebLLM compatible with the new PagedKVCache interface, encapsulating it with the goal that WebLLM users will not notice any difference. This PR is equivalent to the changes to `llm_chat.cc` in mlc-ai/mlc-llm#1651, and should address issues like mlc-ai/mlc-llm#1628. There are still existing model compilation issues regarding `workgroup_size` (since WebGPU, unlike most other backends, can only support 256 number of threads). We will address this issue more elegantly soon; for now, compiling llama-based models require manually changing kernel sizes as shown in [this branch](https://github.com/CharlieFRuan/mlc-llm/tree/local-workgroupSize-webLLM-kvCache). This PR is also largely dependent on apache/tvm#16554.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
Suggestion cannot be applied right now. Please check back later.
This PR introduces various WebGPU changes to accommodate the new
PagedKVCacheinterface. All changes below are essential for making models that use PagedKVCache runnable under WebGPU:storage_rewrite.ccAttentionKVCachetoAttentionKVCacheLegacyfor the old KVcache interface inlm_support.cc; includepaged_kv_cache.ccwhen makingwasm_runtimesubsequentlywhile (true)rather thanwhile (1)maxStorageBuffersPerShaderStagerather than the default 8 from the WebGPU device when initializing runtime; this is required for new kernels introduced in PagedKVCachedeviceCopyToCPU(), when raw bytes to write are not multiples of 4, we pad them, as required by WebGPU'swriteBuffer().Co-authored-by: Rick Zhou [email protected]