ChunkLlama

repository·main·Indexed 19 days ago

https://github.com/hkunlp/chunkllama

A training-free method for scaling the long-context capabilities of Large Language Models (LLMs) without retraining. It supports Dual Chunk Attention (DCA) and integrates with vLLM for inference. The library provides replacement functions for Llama2, Mistral, Mixtral, and Qwen models, as well as utilities for Flash Decoding and fine-tuning on long conversations.

Tokens
41.5K
Snippets
135
Records
179
Agent score
66%

What's inside ChunkLlama

  1. Overview of vLLM features and capabilities

    main

    vLLM is a high-throughput library designed for LLM inference and serving. It achieves high performance through several core technologies:

    Performance Optimizations

    • PagedAttention: Efficient management of attention key and value (KV) memory.
    • Continuous Batching: Efficiently handles incoming requests by batching them continuously.
    • Fast Execution: Utilizes CUDA/HIP graphs for accelerated model execution.
    • Quantization Support: Supports GPTQ, AWQ, SqueezeLLM, and FP8 KV Cache.
    • Optimized Kernels: Custom CUDA kernels for speed.

    Flexibility and Integration

    • Model Support: Seamlessly integrates with Hugging Face models, including Transformer-like (e.g., Llama), Mixture-of-Experts (e.g., Mixtral), and Multi-modal (e.g., LLaVA) architectures.
    • Distributed Inference: Supports Tensor parallelism for multi-GPU setups.
    • API Compatibility: Provides an OpenAI-compatible API server.
    • Hardware Support: Works with NVIDIA GPUs, AMD GPUs, Intel CPUs, and Intel GPUs.
    • Advanced Features: Supports streaming outputs, parallel sampling, beam search, and (experimentally) prefix caching and multi-LoRA support.
  2. Understand the vLLM Paged Attention Kernel

    main
    vLLM uses a custom multi-head query attention kernel located at csrc/attention/attention_kernels.cu. This kernel is specifically optimized for vLLM's paged KV cache system, where key and value caches are stored in separate blocks. The kernel achieves high performance through a specialized memory layout and access pattern designed to optimize data movement from global memory to shared memory.
  3. ChunkLlama License and Usage Terms

    main

    ChunkLlama is licensed under the Apache License 2.0, which requires the preservation of copyright and license notices.

    Important Restrictions:

    • Data and weights are under the CC-BY-NC 4.0 License.
    • They are licensed for research use only and are strictly non-commercial.
    • Models trained using the provided dataset must not be used for purposes outside of research.
  4. Requirements for vLLM with AWS Neuron

    main

    To use vLLM (version 0.3.3 onwards) for model inferencing and serving on AWS Trainium/Inferentia with Neuron SDK, ensure your environment meets these requirements:

    • OS: Linux
    • Python: 3.8 -- 3.11
    • Accelerator: NeuronCore_v2 (available in trn1 or inf2 instances)
    • PyTorch: 2.0.1 or 2.1.1
    • AWS Neuron SDK: 2.16 or 2.17 (Verified on Python 3.8)

    Note on Capabilities:

    • Paged Attention: Currently not supported in Neuron SDK.
    • Batching: Naive continuous batching is supported via transformers-neuronx.
    • Data Types: Supported types are FP16 and BF16.
  5. What is Automatic Prefix Caching (APC)

    main
    Automatic Prefix Caching (APC) is a feature in vLLM that caches the KV cache of existing queries. When a new query shares a common prefix with a previously processed query, vLLM can reuse the cached KV cache for that shared part. This allows the engine to skip the computation of the shared prefix, reducing latency and increasing throughput during the prefilling phase.
  6. Reduce and write out Attention results (LV)

    main

    After the Value dot product, the accumulated results in accs must be reduced and written to global memory.

    Workflow:

    1. Intra-Warp Reduction: Threads within a warp reduce their accs values so each thread holds the accumulation for its assigned head positions across all tokens in a block.
    2. Inter-Warp Reduction: A reduction is performed across all warps. This involves using shared memory (out_smem) where upper warps write to shared memory and lower warps update the accs values by reading from it.
    3. Global Memory Write: The final results are written from local registers to the output global memory using a calculated out_ptr which points to the specific sequence, head, and partition assigned to the thread.
    // Intra-warp reduction
    for (int i = 0; i < NUM_ROWS_PER_THREAD; i++) {
       float acc = accs[i];
       for (int mask = NUM_V_VECS_PER_ROW / 2; mask >= 1; mask /= 2) {
          acc += VLLM_SHFL_XOR_SYNC(acc, mask);
       }
       accs[i] = acc;
    }
    
    // Output pointer calculation and write-out
    scalar_t* out_ptr = out + seq_idx * num_heads * max_num_partitions * HEAD_SIZE
                    + head_idx * max_num_partitions * HEAD_SIZE
                    + partition_idx * HEAD_SIZE;
    
    for (int i = 0; i < NUM_ROWS_PER_THREAD; i++) {
       const int row_idx = lane / NUM_V_VECS_PER_ROW + i * NUM_ROWS_PER_ITER;
       if (row_idx < HEAD_SIZE && lane % NUM_V_VECS_PER_ROW == 0) {
           from_float(*(out_ptr + row_idx), accs[i]);
       }
    }
  7. How Query Data is Fetched and Stored

    main

    Query data is stored in global memory and fetched into shared memory (q_vecs).

    1. Pointer Calculation: Each thread calculates its own q_ptr to point to its assigned query token data: const scalar_t* q_ptr = q + seq_idx * q_stride + head_idx * HEAD_SIZE;
    2. Shared Memory Storage: Query data is read into q_vecs, which is stored in shared memory to allow multiple threads/warps to access it multiple times: __shared__ Q_vec q_vecs[THREAD_GROUP_SIZE][NUM_VECS_PER_THREAD];
    3. Memory Coalescing: Data is read such that neighboring threads in a thread group handle different rows of vecs, enabling efficient memory coalescing.
    const scalar_t* q_ptr = q + seq_idx * q_stride + head_idx * HEAD_SIZE;
    
    __shared__ Q_vec q_vecs[THREAD_GROUP_SIZE][NUM_VECS_PER_THREAD];
  8. Use Tool Calling in the Chat Completion API

    main

    vLLM supports named function calling in the chat completion API.

    Limitations & Requirements:

    • tool_choice options auto and required are not yet supported.
    • To use a tool, you must define the function in the tools parameter and explicitly specify the function name in the tool_choice parameter.
    • Caller Responsibility: You must manually prompt the model with the tool information; vLLM does not automatically manipulate the prompt for tools.
    • vLLM uses guided decoding to ensure the model's response matches the JSON schema defined in the tools parameter.
  9. Core Concepts of vLLM Paged Attention

    main

    To understand the kernel implementation, you must understand these abstractions:

    • Sequence: Represents a client request. In this single-query kernel, num_seqs equals the total number of tokens being processed in a batch.
    • Context: The set of generated tokens in a sequence.
    • Vec: A list of elements fetched/calculated together. VEC_SIZE (for Q/K) and V_VEC_SIZE (for V) are sized so thread groups/threads can fetch 16 bytes at a time.
    • Thread group: A group of THREAD_GROUP_SIZE threads that fetches and calculates one query token and one key token at a time.
    • Block: A unit of the KV cache storing a fixed BLOCK_SIZE of tokens for one head.
    • Warp: A group of 32 threads (WARP_SIZE) that executes simultaneously. A warp processes the calculation between one query token and key tokens of one entire block (potentially across multiple blocks/iterations).
    • Thread block: A group of NUM_THREADS that can access the same shared memory. In this kernel, a thread block processes the calculation for one query token against a whole context.
    • Grid: The collection of thread blocks with shape (num_heads, num_seqs, max_num_partitions). Each thread block handles one head, one sequence, and one partition.
  10. Understand the vLLM KV block eviction policy

    main

    When the KV cache is full, vLLM uses a specific eviction policy to decide which blocks to remove. The policy follows this priority order:

    1. Reference Count: Evict blocks with a reference count (number of current requests using the block) of 0.
    2. LRU (Least Recently Used): If multiple blocks have a reference count of 0, evict the one that was accessed least recently.
    3. Longest Prefix: If multiple blocks have the same last access time, evict the block at the end of the longest prefix (the one with the maximum number of blocks preceding it).

    This policy is designed to mimic the behavior of RadixAttention for models with full attention, prioritizing the eviction of unused leaf nodes in a prefix tree.

  11. Understand the QK (Query-Key) dot product mechanism in Paged Attention

    main

    In the Paged Attention kernel, the Query-Key (QK) stage calculates the dot product between a query token and various key tokens.

    Mechanism:

    1. Fetch Query: The query data for one token is fetched and stored in q_vecs.
    2. Iterate Keys: The kernel iterates through different k_ptrs (pointing to different tokens) and prepares k_vecs in an inner loop.
    3. Dot Product: A dot multiplication is performed between q_vecs and k_vecs using Qk_dot<>::dot.

    Key Detail: Although each thread only fetches a portion of the query and key data at a time, the Qk_dot<>::dot function performs a cross-thread group reduction. This ensures the returned qk value is the full dot product result for the entire query and key token data, even if HEAD_SIZE is larger than the data fetched by a single thread (e.g., if HEAD_SIZE is 128 and THREAD_GROUP_SIZE is 2, each thread handles 64 elements, but the result is the full 128-element dot product).

    q_vecs = ...
    for ... {
       k_ptr = ...
       for ... {
          k_vecs[i] = ...
       }
       ...
       float qk = scale * Qk_dot<scalar_t, THREAD_GROUP_SIZE>::dot(q_vecs[thread_group_offset], k_vecs);
    }