From 0ad78d20a57d87a12cec9dcad2f6ff8dea1895c2 Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Mon, 5 Aug 2024 10:12:46 +0000 Subject: [PATCH 01/12] style --- Dockerfile_amd | 2 +- server/Makefile-flash-att-v2 | 2 +- server/Makefile-vllm | 4 ++-- .../layers/attention/rocm.py | 16 +++++++++++----- .../custom_modeling/flash_llama_modeling.py | 4 ++++ 5 files changed, 19 insertions(+), 9 deletions(-) diff --git a/Dockerfile_amd b/Dockerfile_amd index 51231638..514891a8 100644 --- a/Dockerfile_amd +++ b/Dockerfile_amd @@ -92,7 +92,7 @@ RUN chmod +x ~/mambaforge.sh && \ # Install flash-attention, torch dependencies RUN pip install numpy einops ninja --no-cache-dir -RUN conda install intel::mkl-static intel::mkl-include +RUN conda install mkl-static mkl-include RUN pip uninstall -y triton && \ git clone --depth 1 --single-branch https://github.com/ROCm/triton.git && \ cd triton/python && \ diff --git a/server/Makefile-flash-att-v2 b/server/Makefile-flash-att-v2 index dbddd0f4..03527329 100644 --- a/server/Makefile-flash-att-v2 +++ b/server/Makefile-flash-att-v2 @@ -1,5 +1,5 @@ flash_att_v2_commit_cuda := v2.6.1 -flash_att_v2_commit_rocm := 2554f490101742ccdc56620a938f847f61754be6 +flash_att_v2_commit_rocm := d83c4129a92e4258081f92dfafd34345b3b06130 build-flash-attention-v2-cuda: pip install -U packaging wheel diff --git a/server/Makefile-vllm b/server/Makefile-vllm index f1f80529..bf4a1498 100644 --- a/server/Makefile-vllm +++ b/server/Makefile-vllm @@ -1,5 +1,5 @@ commit_cuda := d243e9dc7e2c9c2e36a4150ec8e64809cb55c01b -commit_rocm := c6ee53b1be97e3bbc791b95f22827501297f8921 +commit_rocm := c06ccbf90a213688a2c6a85d2e7af3da7bc4b41b build-vllm-cuda: if [ ! -d 'vllm' ]; then \ pip install -U ninja packaging --no-cache-dir && \ @@ -13,7 +13,7 @@ install-vllm-cuda: build-vllm-cuda build-vllm-rocm: if [ ! -d 'vllm' ]; then \ pip install -U ninja packaging --no-cache-dir && \ - git clone https://github.com/fxmarty/rocm-vllm.git vllm; \ + git clone https://github.com/mht-sharma/vllm.git vllm; \ fi cd vllm && git fetch && git checkout $(commit_rocm) && \ PYTORCH_ROCM_ARCH="gfx90a;gfx942" python setup.py build diff --git a/server/text_generation_server/layers/attention/rocm.py b/server/text_generation_server/layers/attention/rocm.py index 69e64162..77ba4c92 100644 --- a/server/text_generation_server/layers/attention/rocm.py +++ b/server/text_generation_server/layers/attention/rocm.py @@ -14,7 +14,7 @@ use_triton = os.getenv("ROCM_USE_FLASH_ATTN_V2_TRITON", "").lower() in {"true", ENGINE = "triton" if use_triton else "ck" try: - from vllm._C import cache_ops + import vllm._custom_ops as ops except Exception as e: raise ImportError( f"Could not import vllm paged attention. Make sure your installation is correct. Complete error: {e}" @@ -33,9 +33,7 @@ def reshape_and_cache( key_cache.view(-1, shape[-2], shape[-1])[slots] = key value_cache.view(-1, shape[-2], shape[-1])[slots] = value else: - cache_ops.reshape_and_cache( - key, value, key_cache, value_cache, slots, "auto", 1.0 - ) + ops.reshape_and_cache(key, value, key_cache, value_cache, slots, "auto", 1.0) def paged_attention( @@ -78,7 +76,7 @@ def paged_attention( # V1 to avoid the overhead of reduction. Also, if the number of # sequences or heads is large, we use V1 since there is enough work # to parallelize. - from vllm._C import ops + import vllm._custom_ops as ops use_v1 = max_s <= 8192 and (max_num_partitions == 1 or num_seqs * num_heads > 512) if use_v1: @@ -180,6 +178,7 @@ if ENGINE == "ck": softmax_scale, window_size_left=-1, causal=True, + softcap=0.0, ): if window_size_left <= 0 and window_size_left != -1: raise ValueError("`window_size_left` must be > 0 or -1") @@ -194,12 +193,19 @@ if ENGINE == "ck": out, cu_seqlens, cu_seqlens, + None, + None, + None, + None, max_s, max_s, 0.0, softmax_scale, False, causal, + window_size_left, + 0, + softcap, False, None, ) diff --git a/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py b/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py index 9ea19a87..56d88956 100644 --- a/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py @@ -313,11 +313,15 @@ class LlamaMLP(nn.Module): # TODO: This is a hotfix to be removed & properly refactored. self.quantize = config.quantize + self.hidden_size = config.hidden_size + def forward(self, hidden_states, adapter_data): if ( SYSTEM == "rocm" and self.hidden_act == "silu" and hidden_states.shape[0] == 1 + and self.hidden_size + != 16384 # TODO: Temporary workaround for `LLMM_Silu` kernel not working with LLama3.1 405B; needs refactoring once fixed. and not self.quantize ): out = torch.empty( From 55e6059eb10b13eb1c02849b69d73237fa41dcca Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Mon, 5 Aug 2024 12:47:21 +0000 Subject: [PATCH 02/12] update torch --- Dockerfile_amd | 34 ++++++++++--------- .../models/flash_causal_lm.py | 3 +- 2 files changed, 19 insertions(+), 18 deletions(-) diff --git a/Dockerfile_amd b/Dockerfile_amd index 514891a8..399bc869 100644 --- a/Dockerfile_amd +++ b/Dockerfile_amd @@ -98,24 +98,26 @@ RUN pip uninstall -y triton && \ cd triton/python && \ pip install . -RUN git clone --depth 1 --recursive --single-branch --branch 2.3-patched https://github.com/fxmarty/pytorch.git pytorch && cd pytorch && pip install -r requirements.txt --no-cache-dir +# RUN git clone --depth 1 --recursive --single-branch --branch 2.3-patched https://github.com/fxmarty/pytorch.git pytorch && cd pytorch && pip install -r requirements.txt --no-cache-dir -ARG _GLIBCXX_USE_CXX11_ABI="1" -ARG CMAKE_PREFIX_PATH="/opt/conda" -ARG PYTORCH_ROCM_ARCH="gfx90a;gfx942" -ARG BUILD_CAFFE2="0" \ - BUILD_CAFFE2_OPS="0" \ - USE_CUDA="0" \ - USE_ROCM="1" \ - BUILD_TEST="0" \ - USE_FBGEMM="0" \ - USE_NNPACK="0" \ - USE_QNNPACK="0" \ - USE_XNNPACK="0" \ - USE_FLASH_ATTENTION="1" \ - USE_MEM_EFF_ATTENTION="0" +# ARG _GLIBCXX_USE_CXX11_ABI="1" +# ARG CMAKE_PREFIX_PATH="/opt/conda" +# ARG PYTORCH_ROCM_ARCH="gfx90a;gfx942" +# ARG BUILD_CAFFE2="0" \ +# BUILD_CAFFE2_OPS="0" \ +# USE_CUDA="0" \ +# USE_ROCM="1" \ +# BUILD_TEST="0" \ +# USE_FBGEMM="0" \ +# USE_NNPACK="0" \ +# USE_QNNPACK="0" \ +# USE_XNNPACK="0" \ +# USE_FLASH_ATTENTION="1" \ +# USE_MEM_EFF_ATTENTION="0" -RUN cd pytorch && python tools/amd_build/build_amd.py && python setup.py install +# RUN cd pytorch && python tools/amd_build/build_amd.py && python setup.py install + +RUN pip install --pre torch torchvision torchaudio --index-url https://download.pytorch.org/whl/nightly/rocm6.1 # Set AS recommended: https://github.com/ROCm/triton/wiki/A-script-to-set-program-execution-environment-in-ROCm ENV HIP_FORCE_DEV_KERNARG=1 diff --git a/server/text_generation_server/models/flash_causal_lm.py b/server/text_generation_server/models/flash_causal_lm.py index 36bb2662..9b870447 100644 --- a/server/text_generation_server/models/flash_causal_lm.py +++ b/server/text_generation_server/models/flash_causal_lm.py @@ -1150,8 +1150,7 @@ class FlashCausalLM(Model): elif CUDA_GRAPHS is not None: tuning_sequences = CUDA_GRAPHS else: - # For seqlen = 1, we dispatch to LLMM1 kernel. - tuning_sequences = [2, 3, 4, 5, 6, 7] + tuning_sequences = [1, 2, 3, 4, 5, 6, 7] tunableop_filepath = os.path.join( HUGGINGFACE_HUB_CACHE, From 5788c942a53e98bfb3c63d04d9d111c1c53465ed Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Tue, 6 Aug 2024 10:29:46 +0000 Subject: [PATCH 03/12] ix issues --- Dockerfile_amd | 35 +++++++++---------- .../layers/attention/rocm.py | 2 +- .../models/flash_causal_lm.py | 2 +- 3 files changed, 19 insertions(+), 20 deletions(-) diff --git a/Dockerfile_amd b/Dockerfile_amd index 399bc869..efc80234 100644 --- a/Dockerfile_amd +++ b/Dockerfile_amd @@ -98,26 +98,25 @@ RUN pip uninstall -y triton && \ cd triton/python && \ pip install . -# RUN git clone --depth 1 --recursive --single-branch --branch 2.3-patched https://github.com/fxmarty/pytorch.git pytorch && cd pytorch && pip install -r requirements.txt --no-cache-dir +RUN git clone --depth 1 --recursive --single-branch --branch main https://github.com/pytorch/pytorch.git pytorch && cd pytorch && pip install -r requirements.txt --no-cache-dir && \ + git checkout da320214e66b5af0f7db8fd18a64dbb519d17b27 -# ARG _GLIBCXX_USE_CXX11_ABI="1" -# ARG CMAKE_PREFIX_PATH="/opt/conda" -# ARG PYTORCH_ROCM_ARCH="gfx90a;gfx942" -# ARG BUILD_CAFFE2="0" \ -# BUILD_CAFFE2_OPS="0" \ -# USE_CUDA="0" \ -# USE_ROCM="1" \ -# BUILD_TEST="0" \ -# USE_FBGEMM="0" \ -# USE_NNPACK="0" \ -# USE_QNNPACK="0" \ -# USE_XNNPACK="0" \ -# USE_FLASH_ATTENTION="1" \ -# USE_MEM_EFF_ATTENTION="0" +ARG _GLIBCXX_USE_CXX11_ABI="1" +ARG CMAKE_PREFIX_PATH="/opt/conda" +ARG PYTORCH_ROCM_ARCH="gfx90a;gfx942" +ARG BUILD_CAFFE2="0" \ + BUILD_CAFFE2_OPS="0" \ + USE_CUDA="0" \ + USE_ROCM="1" \ + BUILD_TEST="0" \ + USE_FBGEMM="0" \ + USE_NNPACK="0" \ + USE_QNNPACK="0" \ + USE_XNNPACK="0" \ + USE_FLASH_ATTENTION="1" \ + USE_MEM_EFF_ATTENTION="0" -# RUN cd pytorch && python tools/amd_build/build_amd.py && python setup.py install - -RUN pip install --pre torch torchvision torchaudio --index-url https://download.pytorch.org/whl/nightly/rocm6.1 +RUN cd pytorch && python tools/amd_build/build_amd.py && python setup.py install # Set AS recommended: https://github.com/ROCm/triton/wiki/A-script-to-set-program-execution-environment-in-ROCm ENV HIP_FORCE_DEV_KERNARG=1 diff --git a/server/text_generation_server/layers/attention/rocm.py b/server/text_generation_server/layers/attention/rocm.py index 77ba4c92..da8a4bcd 100644 --- a/server/text_generation_server/layers/attention/rocm.py +++ b/server/text_generation_server/layers/attention/rocm.py @@ -208,7 +208,7 @@ if ENGINE == "ck": softcap, False, None, - ) + )[0] elif ENGINE == "triton": from .flash_attn_triton import triton_attention diff --git a/server/text_generation_server/models/flash_causal_lm.py b/server/text_generation_server/models/flash_causal_lm.py index 9b870447..174bba65 100644 --- a/server/text_generation_server/models/flash_causal_lm.py +++ b/server/text_generation_server/models/flash_causal_lm.py @@ -1159,7 +1159,7 @@ class FlashCausalLM(Model): log_master( logger.info, - f"PyTorch TunableOp (https://github.com/fxmarty/pytorch/tree/2.3-patched/aten/src/ATen/cuda/tunable) is enabled. The warmup may take several minutes, picking the ROCm optimal matrix multiplication kernel for the target lengths {', '.join([str(seqlen) for seqlen in tuning_sequences])}, with typical 5-8% latency improvement for small sequence lengths. The picked GEMMs are saved in the file {tunableop_filepath}. To disable TunableOp, please launch TGI with `PYTORCH_TUNABLEOP_ENABLED=0`.", + f"PyTorch TunableOp is enabled. The warmup may take several minutes, picking the ROCm optimal matrix multiplication kernel for the target lengths {', '.join([str(seqlen) for seqlen in tuning_sequences])}, with typical 5-8% latency improvement for small sequence lengths. The picked GEMMs are saved in the file {tunableop_filepath}. To disable TunableOp, please launch TGI with `PYTORCH_TUNABLEOP_ENABLED=0`.", ) if os.path.isfile(tunableop_filepath): From d61f7e63fa1e940387ae52f6265d803a71f98d8d Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Tue, 6 Aug 2024 12:39:49 +0000 Subject: [PATCH 04/12] fix clone --- Dockerfile_amd | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/Dockerfile_amd b/Dockerfile_amd index efc80234..129a58fc 100644 --- a/Dockerfile_amd +++ b/Dockerfile_amd @@ -92,14 +92,17 @@ RUN chmod +x ~/mambaforge.sh && \ # Install flash-attention, torch dependencies RUN pip install numpy einops ninja --no-cache-dir -RUN conda install mkl-static mkl-include +RUN conda install intel::mkl-static intel::mkl-include RUN pip uninstall -y triton && \ git clone --depth 1 --single-branch https://github.com/ROCm/triton.git && \ cd triton/python && \ pip install . -RUN git clone --depth 1 --recursive --single-branch --branch main https://github.com/pytorch/pytorch.git pytorch && cd pytorch && pip install -r requirements.txt --no-cache-dir && \ - git checkout da320214e66b5af0f7db8fd18a64dbb519d17b27 +RUN git clone --depth 1 --recursive --single-branch --branch main https://github.com/pytorch/pytorch.git pytorch && \ + cd pytorch && git fetch --depth 1 origin da320214e66b5af0f7db8fd18a64dbb519d17b27 && \ + git checkout da320214e66b5af0f7db8fd18a64dbb519d17b27 && \ + pip install -r requirements.txt --no-cache-dir + ARG _GLIBCXX_USE_CXX11_ABI="1" ARG CMAKE_PREFIX_PATH="/opt/conda" From e5578555582e3cc44d871ca631d1b9e8e7d7eeb2 Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Tue, 6 Aug 2024 12:46:35 +0000 Subject: [PATCH 05/12] revert mkl --- Dockerfile_amd | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Dockerfile_amd b/Dockerfile_amd index 129a58fc..d8bb8c47 100644 --- a/Dockerfile_amd +++ b/Dockerfile_amd @@ -92,7 +92,7 @@ RUN chmod +x ~/mambaforge.sh && \ # Install flash-attention, torch dependencies RUN pip install numpy einops ninja --no-cache-dir -RUN conda install intel::mkl-static intel::mkl-include +RUN conda install mkl-static mkl-include RUN pip uninstall -y triton && \ git clone --depth 1 --single-branch https://github.com/ROCm/triton.git && \ cd triton/python && \ From ff0505e7f967b063de3babd70d94a278e7dfed49 Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Wed, 4 Sep 2024 05:46:28 +0000 Subject: [PATCH 06/12] added custom PA --- Dockerfile_amd | 46 +++++++--- server/Makefile-flash-att-v2 | 2 +- server/Makefile-vllm | 2 +- .../layers/attention/cuda.py | 1 + .../layers/attention/ipex.py | 1 + .../layers/attention/rocm.py | 87 ++++++++++++++----- .../custom_modeling/flash_cohere_modeling.py | 1 + .../custom_modeling/flash_dbrx_modeling.py | 1 + .../flash_deepseek_v2_modeling.py | 2 + .../custom_modeling/flash_gemma2_modeling.py | 1 + .../custom_modeling/flash_gemma_modeling.py | 1 + .../custom_modeling/flash_gpt2_modeling.py | 1 + .../custom_modeling/flash_llama_modeling.py | 3 + .../custom_modeling/flash_mistral_modeling.py | 2 + .../custom_modeling/flash_mixtral_modeling.py | 1 + .../custom_modeling/flash_neox_modeling.py | 1 + .../custom_modeling/flash_phi_modeling.py | 1 + .../custom_modeling/flash_qwen2_modeling.py | 1 + .../custom_modeling/flash_rw_modeling.py | 2 + .../flash_santacoder_modeling.py | 1 + .../flash_starcoder2_modeling.py | 1 + .../models/flash_causal_lm.py | 4 +- 22 files changed, 128 insertions(+), 35 deletions(-) diff --git a/Dockerfile_amd b/Dockerfile_amd index d8bb8c47..fd612af5 100644 --- a/Dockerfile_amd +++ b/Dockerfile_amd @@ -39,7 +39,7 @@ COPY launcher launcher RUN cargo build --profile release-opt # Text Generation Inference base image for RoCm -FROM rocm/dev-ubuntu-22.04:6.1.1_hip_update AS base +FROM rocm/dev-ubuntu-22.04:6.2 AS base RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends \ build-essential \ @@ -48,23 +48,25 @@ RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --no-ins curl \ git \ make \ + libmsgpack-dev \ libssl-dev \ + llvm-dev \ g++ \ # Needed to build VLLM & flash. rocthrust-dev \ hipsparse-dev \ hipblas-dev \ - hipblaslt-dev \ + hipcub-dev \ rocblas-dev \ hiprand-dev \ + hipfft-dev \ rocrand-dev \ miopen-hip-dev \ - hipfft-dev \ - hipcub-dev \ hipsolver-dev \ rccl-dev \ cmake \ - python3-dev && \ + python3-dev \ + python3-venv && \ rm -rf /var/lib/apt/lists/* # Keep in sync with `server/pyproject.toml @@ -74,7 +76,30 @@ ARG ROCM_VERSION='6.0.2' ARG PYTHON_VERSION='3.10.10' # Automatically set by buildx ARG TARGETPLATFORM -ENV PATH /opt/conda/bin:$PATH +ENV PATH=/opt/conda/bin:$PATH + +ARG PYTORCH_ROCM_ARCH="gfx90a;gfx942" + +RUN curl -fsSL -v -o cmake-3.30.2-linux-x86_64.sh https://github.com/Kitware/CMake/releases/download/v3.30.2/cmake-3.30.2-linux-x86_64.sh \ + && chmod +x cmake-3.30.2-linux-x86_64.sh \ + && ./cmake-3.30.2-linux-x86_64.sh --skip-license --prefix=/usr/local \ + && rm cmake-3.30.2-linux-x86_64.sh + +RUN pip install joblib msgpack + +# Install HIPBLASLt +ARG HIPBLASLT_BRANCH="6f65c6e" +RUN git clone https://github.com/ROCm/hipBLASLt \ + && cd hipBLASLt \ + && git checkout ${HIPBLASLT_BRANCH} \ + && SCCACHE_IDLE_TIMEOUT=1800 ./install.sh --architecture ${PYTORCH_ROCM_ARCH} \ + && cd build/release \ + && make package +RUN dpkg -i hipBLASLt/build/release/*.deb \ + && sed -i 's/, hipblaslt-dev \(.*\), hipcub-dev/, hipcub-dev/g' /var/lib/dpkg/status \ + && sed -i 's/, hipblaslt \(.*\), hipfft/, hipfft/g' /var/lib/dpkg/status; + # && cd .. \ + # && rm -rf hipBLASLt # TGI seem to require libssl.so.1.1 instead of libssl.so.3 so we can't use ubuntu 22.04. Ubuntu 20.04 has python==3.8, and TGI requires python>=3.9, hence the need for miniconda. # Install mamba @@ -98,15 +123,15 @@ RUN pip uninstall -y triton && \ cd triton/python && \ pip install . +ARG PYTORCH_COMMIT="da320214e66b5af0f7db8fd18a64dbb519d17b27" RUN git clone --depth 1 --recursive --single-branch --branch main https://github.com/pytorch/pytorch.git pytorch && \ - cd pytorch && git fetch --depth 1 origin da320214e66b5af0f7db8fd18a64dbb519d17b27 && \ - git checkout da320214e66b5af0f7db8fd18a64dbb519d17b27 && \ + cd pytorch && git fetch --depth 1 origin ${PYTORCH_COMMIT} && \ + git checkout ${PYTORCH_COMMIT} && \ + git submodule update --init --recursive && \ pip install -r requirements.txt --no-cache-dir - ARG _GLIBCXX_USE_CXX11_ABI="1" ARG CMAKE_PREFIX_PATH="/opt/conda" -ARG PYTORCH_ROCM_ARCH="gfx90a;gfx942" ARG BUILD_CAFFE2="0" \ BUILD_CAFFE2_OPS="0" \ USE_CUDA="0" \ @@ -221,4 +246,3 @@ COPY ./tgi-entrypoint.sh /tgi-entrypoint.sh RUN chmod +x /tgi-entrypoint.sh ENTRYPOINT ["/tgi-entrypoint.sh"] -CMD ["--json-output"] diff --git a/server/Makefile-flash-att-v2 b/server/Makefile-flash-att-v2 index 03527329..74293d9c 100644 --- a/server/Makefile-flash-att-v2 +++ b/server/Makefile-flash-att-v2 @@ -1,5 +1,5 @@ flash_att_v2_commit_cuda := v2.6.1 -flash_att_v2_commit_rocm := d83c4129a92e4258081f92dfafd34345b3b06130 +flash_att_v2_commit_rocm := 3cea2fb6ee54fb7e1aad9db6ac6c9331184b8647 # (Aug28) build-flash-attention-v2-cuda: pip install -U packaging wheel diff --git a/server/Makefile-vllm b/server/Makefile-vllm index bf4a1498..18dcc4a0 100644 --- a/server/Makefile-vllm +++ b/server/Makefile-vllm @@ -1,5 +1,5 @@ commit_cuda := d243e9dc7e2c9c2e36a4150ec8e64809cb55c01b -commit_rocm := c06ccbf90a213688a2c6a85d2e7af3da7bc4b41b +commit_rocm := 4e0929e6e4fa0a3d09d358715c288020ea9dc247 build-vllm-cuda: if [ ! -d 'vllm' ]; then \ pip install -U ninja packaging --no-cache-dir && \ diff --git a/server/text_generation_server/layers/attention/cuda.py b/server/text_generation_server/layers/attention/cuda.py index dff742dc..c623b7f9 100644 --- a/server/text_generation_server/layers/attention/cuda.py +++ b/server/text_generation_server/layers/attention/cuda.py @@ -42,6 +42,7 @@ def paged_attention( block_tables: torch.Tensor, seqlen: Seqlen, max_s: int, + num_kv_heads: int, softcap: Optional[float] = None, ): # Adapted from: https://github.com/vllm-project/vllm/blob/f8a1e39fae05ca610be8d5a78be9d40f5274e5fc/vllm/model_executor/layers/attention.py diff --git a/server/text_generation_server/layers/attention/ipex.py b/server/text_generation_server/layers/attention/ipex.py index e0956b26..f7aada34 100644 --- a/server/text_generation_server/layers/attention/ipex.py +++ b/server/text_generation_server/layers/attention/ipex.py @@ -58,6 +58,7 @@ def paged_attention( block_tables: torch.Tensor, seqlen: Seqlen, max_s: int, + num_kv_heads: int, ): out = torch.empty_like(query) ipex.llm.modules.PagedAttention.single_query_cached_kv_attention( diff --git a/server/text_generation_server/layers/attention/rocm.py b/server/text_generation_server/layers/attention/rocm.py index da8a4bcd..bd033017 100644 --- a/server/text_generation_server/layers/attention/rocm.py +++ b/server/text_generation_server/layers/attention/rocm.py @@ -8,11 +8,17 @@ from loguru import logger major, minor = torch.cuda.get_device_capability() is_sm75 = major == 7 and minor == 5 -_PARTITION_SIZE = 512 + +_PARTITION_SIZE_V1V2 = 512 +_PARTITION_SIZE_CUSTOM = 256 use_triton = os.getenv("ROCM_USE_FLASH_ATTN_V2_TRITON", "").lower() in {"true", "1"} ENGINE = "triton" if use_triton else "ck" +custom_attn_available = os.getenv("VLLM_USE_ROCM_CUSTOM_PAGED_ATTN", "1") != "0" +if custom_attn_available: + from vllm._custom_C import paged_attention_custom + try: import vllm._custom_ops as ops except Exception as e: @@ -45,6 +51,7 @@ def paged_attention( block_tables: torch.Tensor, input_lengths: Seqlen, max_s: int, + num_kv_heads: int, ): # Adapted from: https://github.com/vllm-project/vllm/blob/f8a1e39fae05ca610be8d5a78be9d40f5274e5fc/vllm/model_executor/layers/attention.py # Copyright 2023 The vLLM team. All rights @@ -66,6 +73,22 @@ def paged_attention( # value_cache => [num_blocks, num_heads, head_size, block_size] block_size = value_cache.shape[3] num_seqs, num_heads, head_size = query.shape + + gqa_ratio = num_heads // num_kv_heads + use_custom = ( + custom_attn_available + and (query.dtype == torch.half or query.dtype == torch.bfloat16) + and (head_size == 128 or head_size == 64) + and (block_size == 16 or block_size == 32) + and (gqa_ratio >= 1 and gqa_ratio <= 16) + and max_s <= 32768 + ) + + if not use_custom: + _PARTITION_SIZE = _PARTITION_SIZE_V1V2 + else: + _PARTITION_SIZE = _PARTITION_SIZE_CUSTOM + max_num_partitions = (max_s + _PARTITION_SIZE - 1) // _PARTITION_SIZE input_lengths = input_lengths.input_lengths @@ -78,7 +101,11 @@ def paged_attention( # to parallelize. import vllm._custom_ops as ops - use_v1 = max_s <= 8192 and (max_num_partitions == 1 or num_seqs * num_heads > 512) + use_v1 = ( + max_s <= 8192 + and (max_num_partitions == 1 or num_seqs * num_heads > 512) + and not use_custom + ) if use_v1: ops.paged_attention_v1( out, @@ -110,24 +137,44 @@ def paged_attention( ) max_logits = torch.empty_like(exp_sums) - ops.paged_attention_v2( - out, - exp_sums, - max_logits, - tmp_output, - query, - key_cache, - value_cache, - kv_head_mapping, - softmax_scale, - block_tables, - input_lengths, - block_size, - max_s, - None, - "auto", - 1.0, - ) + if not use_custom: + ops.paged_attention_v2( + out, + exp_sums, + max_logits, + tmp_output, + query, + key_cache, + value_cache, + kv_head_mapping, + softmax_scale, + block_tables, + input_lengths, + block_size, + max_s, + None, + "auto", + 1.0, + ) + else: + paged_attention_custom( + out, + exp_sums, + max_logits, + tmp_output, + query, + key_cache, + value_cache, + num_kv_heads, + softmax_scale, + block_tables, + input_lengths, + block_size, + max_s, + None, + "auto", + ) + return out diff --git a/server/text_generation_server/models/custom_modeling/flash_cohere_modeling.py b/server/text_generation_server/models/custom_modeling/flash_cohere_modeling.py index e02a31d9..46022854 100644 --- a/server/text_generation_server/models/custom_modeling/flash_cohere_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_cohere_modeling.py @@ -313,6 +313,7 @@ class FlashCohereAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.o_proj( diff --git a/server/text_generation_server/models/custom_modeling/flash_dbrx_modeling.py b/server/text_generation_server/models/custom_modeling/flash_dbrx_modeling.py index d3d1d1ef..bc9e8f15 100644 --- a/server/text_generation_server/models/custom_modeling/flash_dbrx_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_dbrx_modeling.py @@ -352,6 +352,7 @@ class DbrxAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_deepseek_v2_modeling.py b/server/text_generation_server/models/custom_modeling/flash_deepseek_v2_modeling.py index 0905d3c2..d02d6cd3 100644 --- a/server/text_generation_server/models/custom_modeling/flash_deepseek_v2_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_deepseek_v2_modeling.py @@ -380,6 +380,7 @@ class DeepseekV2Attention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) # Remove padding. @@ -424,6 +425,7 @@ class DeepseekV2MLP(nn.Module): def forward(self, hidden_states: torch.Tensor, reduce: bool = True): if ( SYSTEM == "rocm" + and hidden_states.dtype == torch.float16 and self.hidden_act == "silu" and hidden_states.shape[0] == 1 and not self.quantize diff --git a/server/text_generation_server/models/custom_modeling/flash_gemma2_modeling.py b/server/text_generation_server/models/custom_modeling/flash_gemma2_modeling.py index de86f514..76d8e8ba 100644 --- a/server/text_generation_server/models/custom_modeling/flash_gemma2_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_gemma2_modeling.py @@ -256,6 +256,7 @@ class FlashGemma2Attention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, softcap=self.softcap, ) diff --git a/server/text_generation_server/models/custom_modeling/flash_gemma_modeling.py b/server/text_generation_server/models/custom_modeling/flash_gemma_modeling.py index 178efadb..42ae24f3 100644 --- a/server/text_generation_server/models/custom_modeling/flash_gemma_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_gemma_modeling.py @@ -248,6 +248,7 @@ class FlashGemmaAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_gpt2_modeling.py b/server/text_generation_server/models/custom_modeling/flash_gpt2_modeling.py index a19cff8c..8c9dc6d6 100644 --- a/server/text_generation_server/models/custom_modeling/flash_gpt2_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_gpt2_modeling.py @@ -247,6 +247,7 @@ class FlashGPT2Attention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py b/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py index 56d88956..9c4cb64a 100644 --- a/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py @@ -235,6 +235,7 @@ class FlashLlamaAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.o_proj( @@ -318,6 +319,7 @@ class LlamaMLP(nn.Module): def forward(self, hidden_states, adapter_data): if ( SYSTEM == "rocm" + and hidden_states.dtype == torch.float16 and self.hidden_act == "silu" and hidden_states.shape[0] == 1 and self.hidden_size @@ -557,6 +559,7 @@ class FlashLlamaForCausalLM(torch.nn.Module): adapter_data: Optional[torch.Tensor] = None, ) -> Tuple[torch.Tensor, Optional[torch.Tensor]]: inputs_embeds = self.embed_tokens(input_ids) + hidden_states = self.model( inputs_embeds, position_ids, diff --git a/server/text_generation_server/models/custom_modeling/flash_mistral_modeling.py b/server/text_generation_server/models/custom_modeling/flash_mistral_modeling.py index dda53ff3..e4ba0295 100644 --- a/server/text_generation_server/models/custom_modeling/flash_mistral_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_mistral_modeling.py @@ -235,6 +235,7 @@ class MistralAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.o_proj( @@ -300,6 +301,7 @@ class MistralMLP(nn.Module): def forward(self, hidden_states, adapter_data): if ( SYSTEM == "rocm" + and hidden_states.dtype == torch.float16 and self.hidden_act == "silu" and hidden_states.shape[0] == 1 and not self.quantize diff --git a/server/text_generation_server/models/custom_modeling/flash_mixtral_modeling.py b/server/text_generation_server/models/custom_modeling/flash_mixtral_modeling.py index 85431c6c..c84c99c6 100644 --- a/server/text_generation_server/models/custom_modeling/flash_mixtral_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_mixtral_modeling.py @@ -292,6 +292,7 @@ class MixtralAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_neox_modeling.py b/server/text_generation_server/models/custom_modeling/flash_neox_modeling.py index b1b03ad7..103c84a7 100644 --- a/server/text_generation_server/models/custom_modeling/flash_neox_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_neox_modeling.py @@ -180,6 +180,7 @@ class FlashNeoxAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.dense(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_phi_modeling.py b/server/text_generation_server/models/custom_modeling/flash_phi_modeling.py index a9e18348..e067e1d8 100644 --- a/server/text_generation_server/models/custom_modeling/flash_phi_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_phi_modeling.py @@ -209,6 +209,7 @@ class FlashPhiAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.dense(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_qwen2_modeling.py b/server/text_generation_server/models/custom_modeling/flash_qwen2_modeling.py index 865cc85d..8b3b8322 100644 --- a/server/text_generation_server/models/custom_modeling/flash_qwen2_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_qwen2_modeling.py @@ -153,6 +153,7 @@ class Qwen2Attention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_rw_modeling.py b/server/text_generation_server/models/custom_modeling/flash_rw_modeling.py index 708641e7..db4c7e7e 100644 --- a/server/text_generation_server/models/custom_modeling/flash_rw_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_rw_modeling.py @@ -223,6 +223,7 @@ class FlashRWAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.dense(attn_output.view(-1, self.num_heads * self.head_size)) @@ -341,6 +342,7 @@ class FlashRWLargeAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.dense( diff --git a/server/text_generation_server/models/custom_modeling/flash_santacoder_modeling.py b/server/text_generation_server/models/custom_modeling/flash_santacoder_modeling.py index c2676782..2d92f3ff 100644 --- a/server/text_generation_server/models/custom_modeling/flash_santacoder_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_santacoder_modeling.py @@ -308,6 +308,7 @@ class FlashMQAttention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.c_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_starcoder2_modeling.py b/server/text_generation_server/models/custom_modeling/flash_starcoder2_modeling.py index e562eb89..534a3792 100644 --- a/server/text_generation_server/models/custom_modeling/flash_starcoder2_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_starcoder2_modeling.py @@ -258,6 +258,7 @@ class Starcoder2Attention(torch.nn.Module): block_tables, input_lengths, max_s, + self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/flash_causal_lm.py b/server/text_generation_server/models/flash_causal_lm.py index 174bba65..35388f49 100644 --- a/server/text_generation_server/models/flash_causal_lm.py +++ b/server/text_generation_server/models/flash_causal_lm.py @@ -1007,12 +1007,12 @@ class FlashCausalLM(Model): else: self.kv_cache = [ ( - torch.empty( + torch.zeros( (num_blocks, num_heads, head_size // x, BLOCK_SIZE, x), dtype=dtype, device=device, ), - torch.empty( + torch.zeros( (num_blocks, num_heads, head_size, BLOCK_SIZE), dtype=dtype, device=device, From 88e2997b9c47b725dcb48d68b5c8ea448b45d591 Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Fri, 6 Sep 2024 12:23:18 +0000 Subject: [PATCH 07/12] style --- Dockerfile_amd | 2 ++ docs/source/installation_amd.md | 6 ++++++ server/text_generation_server/layers/attention/rocm.py | 2 +- 3 files changed, 9 insertions(+), 1 deletion(-) diff --git a/Dockerfile_amd b/Dockerfile_amd index fd612af5..90a1ddf5 100644 --- a/Dockerfile_amd +++ b/Dockerfile_amd @@ -152,6 +152,7 @@ ENV HIP_FORCE_DEV_KERNARG=1 # On MI250 and MI300, performances for flash with Triton FA are slightly better than CK. # However, Triton requires a tunning for each prompt length, which is prohibitive. ENV ROCM_USE_FLASH_ATTN_V2_TRITON=0 +ENV ROCM_USE_CUSTOM_PAGED_ATTN=1 FROM base AS kernel-builder @@ -246,3 +247,4 @@ COPY ./tgi-entrypoint.sh /tgi-entrypoint.sh RUN chmod +x /tgi-entrypoint.sh ENTRYPOINT ["/tgi-entrypoint.sh"] +CMD ["--json-output"] diff --git a/docs/source/installation_amd.md b/docs/source/installation_amd.md index 931a9e3a..8bf60830 100644 --- a/docs/source/installation_amd.md +++ b/docs/source/installation_amd.md @@ -31,6 +31,12 @@ Two implementations of Flash Attention are available for ROCm, the first is [ROC By default, the Composable Kernel implementation is used. However, the Triton implementation has slightly lower latency on MI250 and MI300, but requires a warmup which can be prohibitive as it needs to be done again for each new prompt length. If needed, FA Triton impelmentation can be enabled with `--env ROCM_USE_FLASH_ATTN_V2_TRITON="0"` when launching TGI's docker container. +## Custom PagedAttention + +For better performance on ROCm, a custom Paged Attention kernel is available and is enabled by default. To disable it and fall back to the PagedAttention v2 kernel, set the environment variable `ROCM_USE_CUSTOM_PAGED_ATTN=0`. + +The custom kernel supports bf16 and fp16 data types, block size of 16, head size of 128, a maximum context length of 16k, and GQA ratios between 1 and 16. For other configurations, we use the PagedAttention v2 kernel. + ## Unsupported features The following features are currently not supported in the ROCm version of TGI, and the supported may be extended in the future: diff --git a/server/text_generation_server/layers/attention/rocm.py b/server/text_generation_server/layers/attention/rocm.py index bd033017..58165dc7 100644 --- a/server/text_generation_server/layers/attention/rocm.py +++ b/server/text_generation_server/layers/attention/rocm.py @@ -15,7 +15,7 @@ _PARTITION_SIZE_CUSTOM = 256 use_triton = os.getenv("ROCM_USE_FLASH_ATTN_V2_TRITON", "").lower() in {"true", "1"} ENGINE = "triton" if use_triton else "ck" -custom_attn_available = os.getenv("VLLM_USE_ROCM_CUSTOM_PAGED_ATTN", "1") != "0" +custom_attn_available = os.getenv("ROCM_USE_CUSTOM_PAGED_ATTN", "1") != "0" if custom_attn_available: from vllm._custom_C import paged_attention_custom From 3f2dc6150009b74c71d3cfeb6ea9d12bbc2f9f7c Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Mon, 9 Sep 2024 10:13:59 +0000 Subject: [PATCH 08/12] fix style --- Dockerfile_amd | 1 + docs/source/installation_amd.md | 4 ++++ server/text_generation_server/layers/linear.py | 6 +++++- server/text_generation_server/models/flash_causal_lm.py | 3 ++- 4 files changed, 12 insertions(+), 2 deletions(-) diff --git a/Dockerfile_amd b/Dockerfile_amd index 90a1ddf5..f01b160d 100644 --- a/Dockerfile_amd +++ b/Dockerfile_amd @@ -153,6 +153,7 @@ ENV HIP_FORCE_DEV_KERNARG=1 # However, Triton requires a tunning for each prompt length, which is prohibitive. ENV ROCM_USE_FLASH_ATTN_V2_TRITON=0 ENV ROCM_USE_CUSTOM_PAGED_ATTN=1 +ENV PYTORCH_TUNABLEOP_TUNING_AFTER_WARMUP=0 FROM base AS kernel-builder diff --git a/docs/source/installation_amd.md b/docs/source/installation_amd.md index 8bf60830..070e268e 100644 --- a/docs/source/installation_amd.md +++ b/docs/source/installation_amd.md @@ -25,6 +25,10 @@ Experimentally, on MI300X, we noticed a 6-8% latency improvement when using Tuna TunableOp is enabled by default, the warmup may take 1-2 minutes. In case you would like to disable TunableOp, please pass `--env PYTORCH_TUNABLEOP_ENABLED="0"` when launcher TGI's docker container. +TunableOps tuning is disabled by default after the warmup phase. If you wish to keep tuning enabled for the entire run, set the environment variable `PYTORCH_TUNABLEOP_TUNING_AFTER_WARMUP=1`. + +Note: With tuning enabled, every time a new input shape is encountered, tuning will be performed, which can slow down the first inference for that shape. + ## Flash attention implementation Two implementations of Flash Attention are available for ROCm, the first is [ROCm/flash-attention](https://github.com/ROCm/flash-attention) based on a [Composable Kernel](https://github.com/ROCm/composable_kernel) (CK) implementation, and the second is a [Triton implementation](https://github.com/huggingface/text-generation-inference/blob/main/server/text_generation_server/layers/attention/flash_attn_triton.py). diff --git a/server/text_generation_server/layers/linear.py b/server/text_generation_server/layers/linear.py index 12d7f83a..78815d74 100644 --- a/server/text_generation_server/layers/linear.py +++ b/server/text_generation_server/layers/linear.py @@ -61,7 +61,11 @@ class FastLinearROCm(torch.nn.Module): weight = self.weight bias = self.bias - if SYSTEM == "rocm" and inp.numel() // inp.shape[-1] == 1: + if ( + SYSTEM == "rocm" + and inp.numel() // inp.shape[-1] == 1 + and inp.dtype == torch.float16 + ): batched = False inp_shape = inp.shape diff --git a/server/text_generation_server/models/flash_causal_lm.py b/server/text_generation_server/models/flash_causal_lm.py index 35388f49..7a3f57ab 100644 --- a/server/text_generation_server/models/flash_causal_lm.py +++ b/server/text_generation_server/models/flash_causal_lm.py @@ -1175,7 +1175,8 @@ class FlashCausalLM(Model): log_master(logger.info, f"Warming up TunableOp for seqlen={seqlen}") self.tunableop_warmup(seqlen) torch.cuda.tunable.write_file(tunableop_filepath) - torch.cuda.tunable.tuning_enable(False) + if os.environ.get("PYTORCH_TUNABLEOP_TUNING_AFTER_WARMUP") != "1": + torch.cuda.tunable.tuning_enable(False) else: log_master( logger.info, From f3bc03843003f7ad8f1aa8e85bd29b85264e2b2a Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Wed, 11 Sep 2024 06:52:30 +0000 Subject: [PATCH 09/12] style --- server/text_generation_server/models/flash_causal_lm.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/server/text_generation_server/models/flash_causal_lm.py b/server/text_generation_server/models/flash_causal_lm.py index 7a3f57ab..bb856078 100644 --- a/server/text_generation_server/models/flash_causal_lm.py +++ b/server/text_generation_server/models/flash_causal_lm.py @@ -1162,6 +1162,10 @@ class FlashCausalLM(Model): f"PyTorch TunableOp is enabled. The warmup may take several minutes, picking the ROCm optimal matrix multiplication kernel for the target lengths {', '.join([str(seqlen) for seqlen in tuning_sequences])}, with typical 5-8% latency improvement for small sequence lengths. The picked GEMMs are saved in the file {tunableop_filepath}. To disable TunableOp, please launch TGI with `PYTORCH_TUNABLEOP_ENABLED=0`.", ) + torch.cuda.tunable.set_filename( + tunableop_filepath, insert_device_ordinal=False + ) + if os.path.isfile(tunableop_filepath): log_master( logger.info, From e2f48fae3d32c5b50d75edba271ebe0182dee979 Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Wed, 11 Sep 2024 07:00:29 +0000 Subject: [PATCH 10/12] hide env vart --- docs/source/installation_amd.md | 4 ---- 1 file changed, 4 deletions(-) diff --git a/docs/source/installation_amd.md b/docs/source/installation_amd.md index 070e268e..8bf60830 100644 --- a/docs/source/installation_amd.md +++ b/docs/source/installation_amd.md @@ -25,10 +25,6 @@ Experimentally, on MI300X, we noticed a 6-8% latency improvement when using Tuna TunableOp is enabled by default, the warmup may take 1-2 minutes. In case you would like to disable TunableOp, please pass `--env PYTORCH_TUNABLEOP_ENABLED="0"` when launcher TGI's docker container. -TunableOps tuning is disabled by default after the warmup phase. If you wish to keep tuning enabled for the entire run, set the environment variable `PYTORCH_TUNABLEOP_TUNING_AFTER_WARMUP=1`. - -Note: With tuning enabled, every time a new input shape is encountered, tuning will be performed, which can slow down the first inference for that shape. - ## Flash attention implementation Two implementations of Flash Attention are available for ROCm, the first is [ROCm/flash-attention](https://github.com/ROCm/flash-attention) based on a [Composable Kernel](https://github.com/ROCm/composable_kernel) (CK) implementation, and the second is a [Triton implementation](https://github.com/huggingface/text-generation-inference/blob/main/server/text_generation_server/layers/attention/flash_attn_triton.py). From 0345816477b00be417c602105a950d7482f734c1 Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Wed, 11 Sep 2024 10:52:10 +0000 Subject: [PATCH 11/12] fix mixtral model --- Dockerfile_amd | 1 + 1 file changed, 1 insertion(+) diff --git a/Dockerfile_amd b/Dockerfile_amd index f01b160d..d8f16e7e 100644 --- a/Dockerfile_amd +++ b/Dockerfile_amd @@ -154,6 +154,7 @@ ENV HIP_FORCE_DEV_KERNARG=1 ENV ROCM_USE_FLASH_ATTN_V2_TRITON=0 ENV ROCM_USE_CUSTOM_PAGED_ATTN=1 ENV PYTORCH_TUNABLEOP_TUNING_AFTER_WARMUP=0 +ENV VLLM_MOE_PADDING=0 FROM base AS kernel-builder From 59fd0cbdff68c20d954aafdecd419b4152f34e5e Mon Sep 17 00:00:00 2001 From: Mohit Sharma Date: Thu, 12 Sep 2024 13:16:13 +0000 Subject: [PATCH 12/12] add skinny kernel and merge fixes --- Dockerfile_amd | 10 ++-- .../layers/attention/cuda.py | 1 - .../layers/attention/ipex.py | 1 - .../layers/attention/rocm.py | 47 ++++++++++--------- .../text_generation_server/layers/linear.py | 44 ++++++++++++----- .../custom_modeling/flash_cohere_modeling.py | 6 +-- .../custom_modeling/flash_dbrx_modeling.py | 6 +-- .../flash_deepseek_v2_modeling.py | 6 +-- .../custom_modeling/flash_gemma2_modeling.py | 7 ++- .../custom_modeling/flash_gemma_modeling.py | 7 ++- .../custom_modeling/flash_gpt2_modeling.py | 7 ++- .../custom_modeling/flash_gptj_modeling.py | 5 +- .../custom_modeling/flash_llama_modeling.py | 6 +-- .../custom_modeling/flash_mistral_modeling.py | 6 +-- .../custom_modeling/flash_mixtral_modeling.py | 6 +-- .../custom_modeling/flash_neox_modeling.py | 7 ++- .../custom_modeling/flash_phi_modeling.py | 7 ++- .../custom_modeling/flash_qwen2_modeling.py | 7 ++- .../custom_modeling/flash_rw_modeling.py | 12 ++--- .../flash_santacoder_modeling.py | 7 ++- .../flash_starcoder2_modeling.py | 7 ++- .../models/flash_causal_lm.py | 3 +- .../text_generation_server/models/globals.py | 7 +++ 23 files changed, 121 insertions(+), 101 deletions(-) diff --git a/Dockerfile_amd b/Dockerfile_amd index 1940b985..2aa2a6bc 100644 --- a/Dockerfile_amd +++ b/Dockerfile_amd @@ -152,9 +152,6 @@ ENV HIP_FORCE_DEV_KERNARG=1 # On MI250 and MI300, performances for flash with Triton FA are slightly better than CK. # However, Triton requires a tunning for each prompt length, which is prohibitive. ENV ROCM_USE_FLASH_ATTN_V2_TRITON=0 -ENV ROCM_USE_CUSTOM_PAGED_ATTN=1 -ENV PYTORCH_TUNABLEOP_TUNING_AFTER_WARMUP=0 -ENV VLLM_MOE_PADDING=0 FROM base AS kernel-builder @@ -245,6 +242,13 @@ ENTRYPOINT ["./entrypoint.sh"] # Final image FROM base-copy +ENV ROCM_USE_CUSTOM_PAGED_ATTN=1 +ENV PYTORCH_TUNABLEOP_TUNING_AFTER_WARMUP=0 +ENV VLLM_MOE_PADDING=0 +ENV ATTENTION=paged +ENV USE_PREFIX_CACHING=0 +ENV ROCM_USE_SKINNY_GEMM=1 + COPY ./tgi-entrypoint.sh /tgi-entrypoint.sh RUN chmod +x /tgi-entrypoint.sh diff --git a/server/text_generation_server/layers/attention/cuda.py b/server/text_generation_server/layers/attention/cuda.py index 592350f4..4b588b5c 100644 --- a/server/text_generation_server/layers/attention/cuda.py +++ b/server/text_generation_server/layers/attention/cuda.py @@ -45,7 +45,6 @@ def paged_attention( block_tables: torch.Tensor, seqlen: Seqlen, max_s: int, - num_kv_heads: int, softcap: Optional[float] = None, ): # Adapted from: https://github.com/vllm-project/vllm/blob/f8a1e39fae05ca610be8d5a78be9d40f5274e5fc/vllm/model_executor/layers/attention.py diff --git a/server/text_generation_server/layers/attention/ipex.py b/server/text_generation_server/layers/attention/ipex.py index 83254598..2d1427ae 100644 --- a/server/text_generation_server/layers/attention/ipex.py +++ b/server/text_generation_server/layers/attention/ipex.py @@ -62,7 +62,6 @@ def paged_attention( block_tables: torch.Tensor, seqlen: Seqlen, max_s: int, - num_kv_heads: int, softcap: Optional[float] = None, ): out = torch.empty_like(query) diff --git a/server/text_generation_server/layers/attention/rocm.py b/server/text_generation_server/layers/attention/rocm.py index 3e003acb..0835cb97 100644 --- a/server/text_generation_server/layers/attention/rocm.py +++ b/server/text_generation_server/layers/attention/rocm.py @@ -50,9 +50,8 @@ def paged_attention( kv_head_mapping: torch.Tensor, softmax_scale: float, block_tables: torch.Tensor, - input_lengths: Seqlen, + seqlen: Seqlen, max_s: int, - num_kv_heads: int, softcap: Optional[float] = None, ): # Adapted from: https://github.com/vllm-project/vllm/blob/f8a1e39fae05ca610be8d5a78be9d40f5274e5fc/vllm/model_executor/layers/attention.py @@ -76,6 +75,7 @@ def paged_attention( block_size = value_cache.shape[3] num_seqs, num_heads, head_size = query.shape + num_kv_heads = key_cache.shape[1] gqa_ratio = num_heads // num_kv_heads use_custom = ( custom_attn_available @@ -92,7 +92,7 @@ def paged_attention( _PARTITION_SIZE = _PARTITION_SIZE_CUSTOM max_num_partitions = (max_s + _PARTITION_SIZE - 1) // _PARTITION_SIZE - input_lengths = input_lengths.input_lengths + input_lengths = seqlen.input_lengths out = torch.empty_like(query) @@ -220,10 +220,10 @@ if ENGINE == "ck": def attention( q, - k, - v, - cu_seqlens, - max_s, + key_cache: torch.Tensor, + value_cache: torch.Tensor, + seqlen: Seqlen, + block_tables: torch.Tensor, softmax_scale, window_size_left=-1, causal=True, @@ -237,17 +237,17 @@ if ENGINE == "ck": # We do not need to check window_size_left (not supported) here, so it is already checked ahead of time at model load. return flash_attn_2_cuda.varlen_fwd( q, - k, - v, + key_cache, + value_cache, out, - cu_seqlens, - cu_seqlens, + seqlen.cu_seqlen_q, + seqlen.cu_seqlen_q, None, None, None, None, - max_s, - max_s, + seqlen.max_q, + seqlen.max_k, 0.0, softmax_scale, False, @@ -264,26 +264,27 @@ elif ENGINE == "triton": def attention( q, - k, - v, - cu_seqlens, - max_s, + key_cache: torch.Tensor, + value_cache: torch.Tensor, + seqlen: Seqlen, + block_tables: torch.Tensor, softmax_scale, window_size_left=-1, causal=True, + softcap=0.0, ): out = torch.empty_like(q) # We do not need to check window_size_left (not supported) here, so it is already checked ahead of time at model load. output, _ = triton_attention( q, - k, - v, + key_cache, + value_cache, out, - cu_seqlens, - cu_seqlens, - max_s, - max_s, + seqlen.cu_seqlen_q, + seqlen.cu_seqlen_q, + seqlen.max_q, + seqlen.max_k, causal, softmax_scale, ) diff --git a/server/text_generation_server/layers/linear.py b/server/text_generation_server/layers/linear.py index 78815d74..69b6294b 100644 --- a/server/text_generation_server/layers/linear.py +++ b/server/text_generation_server/layers/linear.py @@ -1,12 +1,19 @@ import torch from text_generation_server.utils.import_utils import SYSTEM from torch.nn import functional as F +import os if SYSTEM == "rocm": - try: - from vllm import _custom_C - except Exception as e: - raise ImportError(f"Could not load `vllm._custom_C`. Full error: {e}") + ROCM_USE_SKINNY_GEMM = os.getenv("ROCM_USE_SKINNY_GEMM", "True").lower() in ( + "true", + "1", + ) + + if ROCM_USE_SKINNY_GEMM: + try: + from vllm import _custom_C + except Exception as e: + raise ImportError(f"Could not load `vllm._custom_C`. Full error: {e}") class FastLinear(torch.nn.Module): @@ -48,6 +55,14 @@ class FastLinearROCm(torch.nn.Module): else: self.bias = None + self.cu_count = torch.cuda.get_device_properties( + device="cuda" + ).multi_processor_count + self.use_skinny_gemm = ( + ROCM_USE_SKINNY_GEMM + and "gfx1" not in torch.cuda.get_device_properties("cuda").gcnArchName + ) + @classmethod def load(cls, config, prefix: str, weights, bias: bool): weight = weights.get_tensor(f"{prefix}.weight") @@ -62,9 +77,9 @@ class FastLinearROCm(torch.nn.Module): bias = self.bias if ( - SYSTEM == "rocm" - and inp.numel() // inp.shape[-1] == 1 + self.use_skinny_gemm and inp.dtype == torch.float16 + and inp.shape[-1] % 8 == 0 ): batched = False inp_shape = inp.shape @@ -73,13 +88,16 @@ class FastLinearROCm(torch.nn.Module): inp = inp.view(-1, inp_shape[-1]) batched = True - m, k = weight.shape[0], inp_shape[1] - out = torch.empty( - inp_shape[0], weight.shape[0], dtype=inp.dtype, device="cuda" - ) - if (k == 8192 and (m == 1280 or m == 7168)) or (k == 3584 and m == 8192): - _custom_C.LLMM1(weight, inp, out, 8) - elif k <= 8192 and k % 8 == 0 and m % 4 == 0: + m, n, k = weight.shape[0], inp_shape[0], inp_shape[1] + if m > 8 and n <= 4: + out = torch.empty( + inp_shape[0], weight.shape[0], dtype=inp.dtype, device=weight.device + ) + _custom_C.wvSpltK(weight, inp, out, n, self.cu_count) + elif m % 4 == 0 and n == 1 and k <= 8192: + out = torch.empty( + inp_shape[0], weight.shape[0], dtype=inp.dtype, device=weight.device + ) _custom_C.LLMM1(weight, inp, out, 4) else: out = F.linear(inp, weight) diff --git a/server/text_generation_server/models/custom_modeling/flash_cohere_modeling.py b/server/text_generation_server/models/custom_modeling/flash_cohere_modeling.py index 8f6cba35..b0e57d68 100644 --- a/server/text_generation_server/models/custom_modeling/flash_cohere_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_cohere_modeling.py @@ -18,6 +18,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -297,8 +298,8 @@ class FlashCohereAttention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else key, - kv_cache[1] if SYSTEM != "ipex" else value, + kv_cache[0] if PAGED_KV else key, + kv_cache[1] if PAGED_KV else value, seqlen, block_tables, self.softmax_scale, @@ -314,7 +315,6 @@ class FlashCohereAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.o_proj( diff --git a/server/text_generation_server/models/custom_modeling/flash_dbrx_modeling.py b/server/text_generation_server/models/custom_modeling/flash_dbrx_modeling.py index 478b5b16..8bce4e57 100644 --- a/server/text_generation_server/models/custom_modeling/flash_dbrx_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_dbrx_modeling.py @@ -13,6 +13,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -336,8 +337,8 @@ class DbrxAttention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv[:, 0], - kv_cache[1] if SYSTEM != "ipex" else kv[:, 1], + kv_cache[0] if PAGED_KV else kv[:, 0], + kv_cache[1] if PAGED_KV else kv[:, 1], seqlen, block_tables, self.softmax_scale, @@ -353,7 +354,6 @@ class DbrxAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_deepseek_v2_modeling.py b/server/text_generation_server/models/custom_modeling/flash_deepseek_v2_modeling.py index 0aa948e7..56136381 100644 --- a/server/text_generation_server/models/custom_modeling/flash_deepseek_v2_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_deepseek_v2_modeling.py @@ -15,6 +15,7 @@ from typing import Any, Dict, List, Optional, Tuple +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed from text_generation_server.layers import ( @@ -363,8 +364,8 @@ class DeepseekV2Attention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else key, - kv_cache[1] if SYSTEM != "ipex" else value, + kv_cache[0] if PAGED_KV else key, + kv_cache[1] if PAGED_KV else value, seqlen, block_tables, self.softmax_scale, @@ -380,7 +381,6 @@ class DeepseekV2Attention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) # Remove padding. diff --git a/server/text_generation_server/models/custom_modeling/flash_gemma2_modeling.py b/server/text_generation_server/models/custom_modeling/flash_gemma2_modeling.py index 6bd4aac5..1ad88801 100644 --- a/server/text_generation_server/models/custom_modeling/flash_gemma2_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_gemma2_modeling.py @@ -18,6 +18,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -25,7 +26,6 @@ from torch import nn from transformers.activations import ACT2FN from transformers.configuration_utils import PretrainedConfig from typing import Optional, List, Tuple -from text_generation_server.utils.import_utils import SYSTEM from text_generation_server.layers.attention import ( paged_attention, attention, @@ -237,8 +237,8 @@ class FlashGemma2Attention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv[:, 0], - kv_cache[1] if SYSTEM != "ipex" else kv[:, 1], + kv_cache[0] if PAGED_KV else kv[:, 0], + kv_cache[1] if PAGED_KV else kv[:, 1], seqlen, block_tables, self.softmax_scale, @@ -257,7 +257,6 @@ class FlashGemma2Attention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, softcap=self.softcap, ) diff --git a/server/text_generation_server/models/custom_modeling/flash_gemma_modeling.py b/server/text_generation_server/models/custom_modeling/flash_gemma_modeling.py index c253e6ab..a401798a 100644 --- a/server/text_generation_server/models/custom_modeling/flash_gemma_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_gemma_modeling.py @@ -18,6 +18,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -25,7 +26,6 @@ from torch import nn from transformers.activations import ACT2FN from transformers.configuration_utils import PretrainedConfig from typing import Optional, List, Tuple -from text_generation_server.utils.import_utils import SYSTEM from text_generation_server.layers.attention import ( paged_attention, attention, @@ -231,8 +231,8 @@ class FlashGemmaAttention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv[:, 0], - kv_cache[1] if SYSTEM != "ipex" else kv[:, 1], + kv_cache[0] if PAGED_KV else kv[:, 0], + kv_cache[1] if PAGED_KV else kv[:, 1], seqlen, block_tables, self.softmax_scale, @@ -249,7 +249,6 @@ class FlashGemmaAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_gpt2_modeling.py b/server/text_generation_server/models/custom_modeling/flash_gpt2_modeling.py index 90382583..33f20b9a 100644 --- a/server/text_generation_server/models/custom_modeling/flash_gpt2_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_gpt2_modeling.py @@ -18,13 +18,13 @@ # See the License for the specific language governing permissions and # limitations under the License. +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed from torch import nn from transformers.activations import ACT2FN from typing import Optional, List, Tuple -from text_generation_server.utils.import_utils import SYSTEM from text_generation_server.layers.attention import ( paged_attention, attention, @@ -231,8 +231,8 @@ class FlashGPT2Attention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else key, - kv_cache[1] if SYSTEM != "ipex" else value, + kv_cache[0] if PAGED_KV else key, + kv_cache[1] if PAGED_KV else value, seqlen, block_tables, self.softmax_scale, @@ -248,7 +248,6 @@ class FlashGPT2Attention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_gptj_modeling.py b/server/text_generation_server/models/custom_modeling/flash_gptj_modeling.py index ef071d46..f2197069 100644 --- a/server/text_generation_server/models/custom_modeling/flash_gptj_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_gptj_modeling.py @@ -18,6 +18,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -192,8 +193,8 @@ class FlashGPTJAttention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else key, - kv_cache[1] if SYSTEM != "ipex" else value, + kv_cache[0] if PAGED_KV else key, + kv_cache[1] if PAGED_KV else value, seqlen, block_tables, self.softmax_scale, diff --git a/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py b/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py index 39218531..6be89297 100644 --- a/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_llama_modeling.py @@ -28,6 +28,7 @@ from torch import nn from transformers.activations import ACT2FN from text_generation_server.utils.import_utils import SYSTEM +from text_generation_server.models.globals import PAGED_KV from text_generation_server.layers.attention import ( paged_attention, attention, @@ -220,8 +221,8 @@ class FlashLlamaAttention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv[:, 0], - kv_cache[1] if SYSTEM != "ipex" else kv[:, 1], + kv_cache[0] if PAGED_KV else kv[:, 0], + kv_cache[1] if PAGED_KV else kv[:, 1], seqlen, block_tables, self.softmax_scale, @@ -237,7 +238,6 @@ class FlashLlamaAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.o_proj( diff --git a/server/text_generation_server/models/custom_modeling/flash_mistral_modeling.py b/server/text_generation_server/models/custom_modeling/flash_mistral_modeling.py index dacda101..3b56bbab 100644 --- a/server/text_generation_server/models/custom_modeling/flash_mistral_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_mistral_modeling.py @@ -18,6 +18,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -218,8 +219,8 @@ class MistralAttention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv_to_cache[:, 0], - kv_cache[1] if SYSTEM != "ipex" else kv_to_cache[:, 1], + kv_cache[0] if PAGED_KV else kv_to_cache[:, 0], + kv_cache[1] if PAGED_KV else kv_to_cache[:, 1], seqlen, block_tables, self.softmax_scale, @@ -236,7 +237,6 @@ class MistralAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.o_proj( diff --git a/server/text_generation_server/models/custom_modeling/flash_mixtral_modeling.py b/server/text_generation_server/models/custom_modeling/flash_mixtral_modeling.py index b35688ec..3451158b 100644 --- a/server/text_generation_server/models/custom_modeling/flash_mixtral_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_mixtral_modeling.py @@ -18,6 +18,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -275,8 +276,8 @@ class MixtralAttention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv_to_cache[:, 0], - kv_cache[1] if SYSTEM != "ipex" else kv_to_cache[:, 1], + kv_cache[0] if PAGED_KV else kv_to_cache[:, 0], + kv_cache[1] if PAGED_KV else kv_to_cache[:, 1], seqlen, block_tables, self.softmax_scale, @@ -293,7 +294,6 @@ class MixtralAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_neox_modeling.py b/server/text_generation_server/models/custom_modeling/flash_neox_modeling.py index 698f0343..2d3be430 100644 --- a/server/text_generation_server/models/custom_modeling/flash_neox_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_neox_modeling.py @@ -18,6 +18,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -26,7 +27,6 @@ from transformers.activations import ACT2FN from transformers.modeling_utils import PreTrainedModel from transformers.models.gpt_neox import GPTNeoXConfig as TransformersGPTNeoXConfig from typing import Optional, List, Tuple -from text_generation_server.utils.import_utils import SYSTEM from text_generation_server.layers.attention import ( paged_attention, attention, @@ -172,8 +172,8 @@ class FlashNeoxAttention(torch.nn.Module): # flash attention attn_output = attention( qkv[:, 0], - kv_cache[0] if SYSTEM != "ipex" else qkv[:, 1], - kv_cache[1] if SYSTEM != "ipex" else qkv[:, 2], + kv_cache[0] if PAGED_KV else qkv[:, 1], + kv_cache[1] if PAGED_KV else qkv[:, 2], seqlen, block_tables, self.softmax_scale, @@ -189,7 +189,6 @@ class FlashNeoxAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.dense(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_phi_modeling.py b/server/text_generation_server/models/custom_modeling/flash_phi_modeling.py index 38e8c884..76e406a7 100644 --- a/server/text_generation_server/models/custom_modeling/flash_phi_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_phi_modeling.py @@ -1,3 +1,4 @@ +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -25,7 +26,6 @@ from text_generation_server.layers.layernorm import ( from text_generation_server.layers.rotary import ( PositionRotaryEmbedding, ) -from text_generation_server.utils.import_utils import SYSTEM class PhiConfig(PretrainedConfig): @@ -194,8 +194,8 @@ class FlashPhiAttention(torch.nn.Module): if cu_seqlen_prefill is not None: attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv[:, 0], - kv_cache[1] if SYSTEM != "ipex" else kv[:, 1], + kv_cache[0] if PAGED_KV else kv[:, 0], + kv_cache[1] if PAGED_KV else kv[:, 1], seqlen, block_tables, self.softmax_scale, @@ -211,7 +211,6 @@ class FlashPhiAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.dense(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_qwen2_modeling.py b/server/text_generation_server/models/custom_modeling/flash_qwen2_modeling.py index d43401ba..0f0dbf5e 100644 --- a/server/text_generation_server/models/custom_modeling/flash_qwen2_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_qwen2_modeling.py @@ -1,3 +1,4 @@ +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -21,7 +22,6 @@ from text_generation_server.layers.rotary import PositionRotaryEmbedding from text_generation_server.layers.layernorm import ( FastRMSNorm, ) -from text_generation_server.utils.import_utils import SYSTEM def load_attention(config, prefix, weights): @@ -137,8 +137,8 @@ class Qwen2Attention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv_to_cache[:, 0], - kv_cache[1] if SYSTEM != "ipex" else kv_to_cache[:, 1], + kv_cache[0] if PAGED_KV else kv_to_cache[:, 0], + kv_cache[1] if PAGED_KV else kv_to_cache[:, 1], seqlen, block_tables, self.softmax_scale, @@ -155,7 +155,6 @@ class Qwen2Attention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_rw_modeling.py b/server/text_generation_server/models/custom_modeling/flash_rw_modeling.py index 765cf39e..ba516881 100644 --- a/server/text_generation_server/models/custom_modeling/flash_rw_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_rw_modeling.py @@ -1,11 +1,11 @@ from typing import List, Optional, Tuple +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed from torch import nn from transformers.configuration_utils import PretrainedConfig from transformers.modeling_utils import PreTrainedModel -from text_generation_server.utils.import_utils import SYSTEM from text_generation_server.layers import ( SpeculativeHead, TensorParallelColumnLinear, @@ -207,8 +207,8 @@ class FlashRWAttention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv[:, 0], - kv_cache[1] if SYSTEM != "ipex" else kv[:, 1], + kv_cache[0] if PAGED_KV else kv[:, 0], + kv_cache[1] if PAGED_KV else kv[:, 1], seqlen, block_tables, self.softmax_scale, @@ -224,7 +224,6 @@ class FlashRWAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.dense(attn_output.view(-1, self.num_heads * self.head_size)) @@ -326,8 +325,8 @@ class FlashRWLargeAttention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv[:, :, 0].contiguous(), - kv_cache[1] if SYSTEM != "ipex" else kv[:, :, 1].contiguous(), + kv_cache[0] if PAGED_KV else kv[:, :, 0].contiguous(), + kv_cache[1] if PAGED_KV else kv[:, :, 1].contiguous(), seqlen, block_tables, self.softmax_scale, @@ -343,7 +342,6 @@ class FlashRWLargeAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.dense( diff --git a/server/text_generation_server/models/custom_modeling/flash_santacoder_modeling.py b/server/text_generation_server/models/custom_modeling/flash_santacoder_modeling.py index 80c280c8..fa074606 100644 --- a/server/text_generation_server/models/custom_modeling/flash_santacoder_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_santacoder_modeling.py @@ -1,3 +1,4 @@ +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -22,7 +23,6 @@ from text_generation_server.layers.gptq import GPTQWeightsLoader from text_generation_server.layers.layernorm import ( FastLayerNorm, ) -from text_generation_server.utils.import_utils import SYSTEM def load_multi_mqa( @@ -293,8 +293,8 @@ class FlashMQAttention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else key_value[:, 0], - kv_cache[1] if SYSTEM != "ipex" else key_value[:, 1], + kv_cache[0] if PAGED_KV else key_value[:, 0], + kv_cache[1] if PAGED_KV else key_value[:, 1], seqlen, block_tables, self.softmax_scale, @@ -310,7 +310,6 @@ class FlashMQAttention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.c_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/custom_modeling/flash_starcoder2_modeling.py b/server/text_generation_server/models/custom_modeling/flash_starcoder2_modeling.py index 0c4ce05a..30d35632 100644 --- a/server/text_generation_server/models/custom_modeling/flash_starcoder2_modeling.py +++ b/server/text_generation_server/models/custom_modeling/flash_starcoder2_modeling.py @@ -18,6 +18,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +from text_generation_server.models.globals import PAGED_KV import torch import torch.distributed @@ -47,7 +48,6 @@ from text_generation_server.layers.rotary import ( PositionRotaryEmbedding, ) from text_generation_server.utils.weights import UnquantizedWeight -from text_generation_server.utils.import_utils import SYSTEM class Starcoder2Config(PretrainedConfig): @@ -242,8 +242,8 @@ class Starcoder2Attention(torch.nn.Module): # flash attention attn_output = attention( query, - kv_cache[0] if SYSTEM != "ipex" else kv_to_cache[:, 0], - kv_cache[1] if SYSTEM != "ipex" else kv_to_cache[:, 1], + kv_cache[0] if PAGED_KV else kv_to_cache[:, 0], + kv_cache[1] if PAGED_KV else kv_to_cache[:, 1], seqlen, block_tables, self.softmax_scale, @@ -260,7 +260,6 @@ class Starcoder2Attention(torch.nn.Module): block_tables, seqlen, max_s, - self.num_key_value_heads, ) return self.o_proj(attn_output.view(-1, self.num_heads * self.head_size)) diff --git a/server/text_generation_server/models/flash_causal_lm.py b/server/text_generation_server/models/flash_causal_lm.py index 6f21e846..c4bf8a57 100644 --- a/server/text_generation_server/models/flash_causal_lm.py +++ b/server/text_generation_server/models/flash_causal_lm.py @@ -1379,6 +1379,7 @@ class FlashCausalLM(Model): cu_seqlen_prefill = torch.tensor( [0, seqlen], device=self.device, dtype=torch.int32 ) + max_s = seqlen seqlen = Seqlen( input_lengths=input_lengths, prefix_lengths=prefix_lens_tensor, @@ -1396,7 +1397,7 @@ class FlashCausalLM(Model): block_tables=None, seqlen=seqlen, slots=slots, - max_s=seqlen, + max_s=max_s, lm_head_indices=None, prefill_cache_indices=None, ) diff --git a/server/text_generation_server/models/globals.py b/server/text_generation_server/models/globals.py index 6c518c2c..f04c6df5 100644 --- a/server/text_generation_server/models/globals.py +++ b/server/text_generation_server/models/globals.py @@ -4,6 +4,7 @@ from loguru import logger from typing import Dict, Optional from text_generation_server.utils.log import log_master +from text_generation_server.utils.import_utils import SYSTEM PREFIX_CACHING = os.getenv("USE_PREFIX_CACHING").lower() in {"1", "true"} log_master(logger.info, f"Using prefix caching = {PREFIX_CACHING}") @@ -52,6 +53,12 @@ CUDA_GRAPHS = cuda_graphs # index in all cases. ADAPTER_TO_INDEX: Optional[Dict[str, int]] = None +PAGED_KV: bool +if SYSTEM in {"rocm", "ipex"}: + PAGED_KV = False +else: + PAGED_KV = True + def set_adapter_to_index(adapter_to_index: Dict[str, int]): global ADAPTER_TO_INDEX