Transformer related optimization, including BERT, GPT

Overview

FasterTransformer

This repository provides a script and recipe to run the highly optimized transformer-based encoder and decoder component, and it is tested and maintained by NVIDIA.

Table Of Contents

Model overview

In NLP, encoder and decoder are two important components, with the transformer layer becoming a popular architecture for both components. FasterTransformer implements a highly optimized transformer layer for both the encoder and decoder for inference. On Volta, Turing and Ampere GPUs, the computing power of Tensor Cores are used automatically when the precision of the data and weights are FP16.

FasterTransformer v1.0 provides a highly optimized BERT equivalent Transformer layer for inference, including C++ API, TensorFlow op and TensorRT plugin. The experiments show that FasterTransformer v1 can provide 1.3 ~ 2 times speedup on NVIDIA Tesla T4 and NVIDIA Tesla V100 for inference.

In FasterTransformer v2.0, we have added a highly optimized decoder and decoding models based on OpenNMT-TF, an open-source library. Here, the decoder is the model that contains some transformer layers. On the other hand, decoding refers to the whole translating process, including the lookup embedding table, position encoding, a decoder and beam search.

In FasterTransformer v2.1, we add some important features. First one is the supporting on PyTorch. Recently, there are more and more PyTorch users. We hope the users of PyTorch can also use the FasterTransformer in their application and research. The second feature is the supporting of Effective Transformer. This idea is proposed by ByteDance. We call this feature as Effective FasterTransformer It removes the useless padding of encoder input to reduce the computing cost. Third, in addition to decoding with beam search, we also provide the decoding with sampling module. Finally, we optimize many kernels of encoder, decoder and beam search to improve the speed of FasterTransformer.

In FasterTransformer v3.0, we implemented the INT8 quantization for encoder (also supporting Effective FasterTransformer). With INT8 quantization, we can take advantage of the powerful INT8 tensor core in Turing GPU to achieve better inference performance (INT8 quantization in FT 3.0 is only supported on device with SM >= 7.5). We also provide quantization tools of tensorflow.

In FasterTransformer v3.1, we provide following new features and enhancements. First, we optimize the INT8 kernel of encoder to achieve better performance. Compare to FasterTransformer v3.0, the performance of INT8 quantization brings at most 1.75x speedup. Second, we provide a PyTorch tool to let user be able to train a INT8 quantized model on PyTorch. Besides, FasterTransformer also starts to support the INT8 inference with PyTorch op. So, the users of PyTorch can leverage the INT8 inference. Third, we integrate the fused multi-head attention kernel of TensorRT plugin into FasterTransformer to improve the speed of encoder on Turing and new GPUs. This optimization can bring about 10% ~ 20% speedup compare to original implementation. Finally, we add the supporting of GPT-2 model, which is an important and popular model for decoder.

In FasterTransformer v4.0, we provide the multi-nodes multi-gpu inference for GPT model. Compare to usual framework to train giant model like Megatron, FasterTransformer provides 1.2x ~ 3x speedup. Besides, integrating the INT8 fused multi-head attention kernel of TensorRT plugin to further improve the performance of FasterTransformer encoder on INT8. We also add supporting of FP16 fused multi-head attention kernel for V100. Finally, we optimize the decoding module. Compare to v3.1, v4.0 provides at most 2x speedup.

The following graph demonstrates the model architecture.

Fig. 1 Encoder-Decoding model architecture.

FasterTransformer is built on top of CUDA, cuBLAS and cuBLASLt, providing the C++ API and TensorFlow/PyTorch OPs. Users can integrate them into TensorFlow, PyTorch, or other inference service codes that are built in native C++. We also provide some simple sample code to demonstrate how to use the encoder, decoder and to carry out decoding in C++, TensorFlow and PyTorch.

More details are in docs/encoder_guide.md, docs/decoder_guide.md and docs/gpt_guide.md. Some common questions and the respective answers are put in docs/QAList.md

Support matrix

The following matrix shows the architecture differences between the model.

Architecure Encoder Encoder INT8
quantization
Decoder Decoding with
beam search
Decoding with
sampling
GPT-2 GPT-3
v1 Yes No No No No No No
v2 Yes No Yes Yes No No No
v2.1 Yes No Yes Yes Yes No No
v3.0 Yes Yes Yes Yes Yes No No
v3.1 Yes Yes Yes Yes Yes Yes No
v4.0 Yes Yes Yes Yes Yes Yes Yes

Setup

The following section lists the requirements to use FasterTransformer.

Requirements

  • CMake >= 3.8 for Tensorflow, CMake >= 3.13 for PyTorch
  • CUDA 10.1 or newer version
  • Python 3 is recommended because some features are not supported in python 2
  • Tensorflow 1.13 or 1.14 or 1.15
  • PyTorch >= 1.5.0

These components are readily available within the NGC TensorFlow/PyTorch Docker image below.

Ensure you have the following components:

For more information about how to get started with NGC containers, see the following sections from the NVIDIA GPU Cloud Documentation and the Deep Learning Documentation:

For those unable to use the NGC container, to set up the required environment or create your own container, see the versioned NVIDIA Container Support Matrix.

Quick Start Guide

The following section shows how to use FasterTransformer on the NGC container.

Build the FasterTransformer

  1. Run the container.

    You can choose the tensorflow version and python version you want. Here, we list some possible images:

    • nvcr.io/nvidia/tensorflow:19.07-py2 contains the TensorFlow 1.14 and python 2.7.
    • nvcr.io/nvidia/tensorflow:20.12-tf1-py3 contains the TensorFlow 1.15 and python 3.8.
    • nvcr.io/nvidia/pytorch:20.03-py3 contains the PyTorch 1.5.0 and python 3.6
    • nvcr.io/nvidia/pytorch:20.07-py3 contains the PyTorch 1.6.0 and python 3.6
    • nvcr.io/nvidia/pytorch:20.12-py3 contains the PyTorch 1.8.0 and python 3.8

    To achieve best performance, we recommand to use the latest image. For example, running image nvcr.io/nvidia/tensorflow:20.12-tf1-py3 by

    nvidia-docker run -ti --rm nvcr.io/nvidia/tensorflow:20.12-tf1-py3 bash
  2. Clone the repository.

    git clone https://github.com/NVIDIA/FasterTransformer.git
    mkdir -p build
    cd build
  3. Build the project.

    3.1 build with C++

    cmake -DSM=xx -DCMAKE_BUILD_TYPE=Release ..
    make

    Note: xx is the compute capability of your GPU. For example, 60 (P40) or 61 (P4) or 70 (V100) or 75(T4) or 80 (A100).

    3.2 build with TensorFlow

    Uses need to set the path of TensorFlow. For example, if we use nvcr.io/nvidia/tensorflow:20.12-tf1-py3, then

    cmake -DSM=xx -DCMAKE_BUILD_TYPE=Release -DBUILD_TF=ON -DTF_PATH=/usr/local/lib/python3.8/dist-packages/tensorflow_core/ ..
    make 

    Note: xx is the compute capability of your GPU. For example, 60 (P40) or 61 (P4) or 70 (V100) or 75(T4) or 80 (A100).

    3.3 build with PyTorch

    cmake -DSM=xx -DCMAKE_BUILD_TYPE=Release -DBUILD_PYT=ON ..
    make

    Note: xx is the compute capability of your GPU. For example, 60 (P40) or 61 (P4) or 70 (V100) or 75(T4) or 80 (A100).

    This will build the TorchScript custom class. Please make sure that the PyTorch >= 1.5.0.

    Note: From FasterTransformer 3.1, TorchScript custom op (function type) is deprecated. From FasterTransformer 4.0, Eager mode PyTorch extension is deprecated.

Execute the encoder demos

  1. Run FasterTransformer encoder on C++

    ./bin/encoder_gemm <batch_size> <sequence_length> <head_number> <size_per_head> <is_use_fp16> <int8_mode>
    ./bin/encoder_sample <batch_size> <num_layers> <sequence_length> <head_number> <size_per_head> <is_use_fp16> <is_remove_padding> <int8_mode> <allow_gemm_test>

    1.1 Run FasterTransformer encoder under FP32 on C++

    ./bin/encoder_gemm 32 32 12 64 0 0
    ./bin/encoder_sample 32 12 32 12 64 0 0 0 0

    1.2 Run FasterTransformer encoder under FP16 on C++

    ./bin/encoder_gemm 32 32 12 64 1 0
    ./bin/encoder_sample 32 12 32 12 64 1 0 0 0

    1.3 Run FasterTransformer encoder under INT8 on C++

    We implement two INT8 pipelines. For int8_mode == 1 (int8v1), we don't quantize residual connection, use int32 as the output of int8 gemms and use per-channel quantization for weights; for int8_mode == 2 (int8v2), we quantize residual connection, use int8 as the output of int8 gemms and use per-tensor quantization for weights. Generally speaking, int8_mode == 1 will have higher accuracy while int8_mode == 2 will have better performance.

    feature int8_mode == 1 int8_mode == 2
    quantize residual No Yes
    int8 output gemm No Yes
    per-channel quantiztion for weights Yes No
    #For int8_mode == 1
    ./bin/encoder_gemm 32 32 12 64 1 1
    ./bin/encoder_sample 32 12 32 12 64 1 0 1 0
    
    #For int8_mode == 2
    ./bin/encoder_gemm 32 32 12 64 1 2
    ./bin/encoder_sample 32 12 32 12 64 1 0 2 0

    1.4 Run Effective FasterTransformer under FP32 on C++

    ./bin/encoder_gemm 32 32 12 64 0 0
    ./bin/encoder_sample 32 12 32 12 64 0 1 0 0

    1.5 Run Effective FasterTransformer under INT8 on C++

    #For int8_mode == 1
    ./bin/encoder_gemm 32 32 12 64 1 1
    ./bin/encoder_sample 32 12 32 12 64 1 1 1 0
    
    #For int8_mode == 2
    ./bin/encoder_gemm 32 32 12 64 1 2
    ./bin/encoder_sample 32 12 32 12 64 1 1 2 0
    
  2. Run FasterTransformer encoder on TensorFlow

    2.1 Run FasterTransformer encoder under FP32 on TensorFlow

    ./bin/encoder_gemm 32 32 12 64 0 0
    python tensorflow/encoder_sample.py \
            --batch_size 32 \
            --max_seq_len 32 \
            --head_number 12 \
            --size_per_head 64 \
            --num_layer 12 \
            --data_type fp32 \
            --test_time 1 \
            --allow_gemm_test False

    If use sets --test_time 1, the program will show the performance of TensorFlow, FasterTransformer and FasterTransformer with removing padding.

    2.2 Run FasterTransformer encoder under FP16 on TensorFlow

    ./bin/encoder_gemm 32 32 12 64 1 0
    python tensorflow/encoder_sample.py \
            --batch_size 32 \
            --max_seq_len 32 \
            --head_number 12 \
            --size_per_head 64 \
            --num_layer 12 \
            --data_type fp16 \
            --test_time 1 \
            --allow_gemm_test False

    2.3 Run FasterTransformer encoder under INT8 on TensorFlow

    #For int8_mode == 1
    ./bin/encoder_gemm 32 32 12 64 1 1
    python tensorflow/encoder_sample.py \
            --batch_size 32 \
            --max_seq_len 32 \
            --head_number 12 \
            --size_per_head 64 \
            --num_layer 12 \
            --data_type fp16 \
            --test_time 1 \
            --int8_mode 1 \
            --allow_gemm_test False
    
    #For int8_mode == 2
    ./bin/encoder_gemm 32 32 12 64 1 2
    python tensorflow/encoder_sample.py \
            --batch_size 32 \
            --max_seq_len 32 \
            --head_number 12 \
            --size_per_head 64 \
            --num_layer 12 \
            --data_type fp16 \
            --test_time 1 \
            --int8_mode 2 \
            --allow_gemm_test False
  3. Run FasterTransformer on PyTorch

    Please install HuggingFace's transformers first before run the demos by

    pip install transformers==2.5.1

    3.1 Run FasterTransformer encoder under FP32 on PyTorch

    ./bin/encoder_gemm 32 32 12 64 0 0
    python pytorch/encoder_sample.py 32 12 32 12 64 --time

    3.2 Run FasterTransformer encoder under FP16 on PyTorch

    ./bin/encoder_gemm 32 32 12 64 1 0
    python pytorch/encoder_sample.py 32 12 32 12 64 --fp16 --time

    3.3 Run FasterTransformer encoder under INT8 on PyTorch

    #For int8_mode == 1
    ./bin/encoder_gemm 32 32 12 64 1 1
    python pytorch/encoder_sample.py 32 12 32 12 64 --int8_mode 1 --time
    
    #For int8_mode == 2
    ./bin/encoder_gemm 32 32 12 64 1 2
    python pytorch/encoder_sample.py 32 12 32 12 64 --int8_mode 2 --time

Execute the decoder/decoding demos

  1. Run FasterTransformer decoding on C++

    ./bin/decoding_gemm <batch_size> <beam_width> <head_number> <size_per_head> <vocab_size> <sequence_length> <encoder_hidden_dim> <is_use_fp16>
    ./bin/decoding_beamsearch_sample <batch_size> <beam_width> <head_number> <size_per_head> <vocab_size> <sequence_length> <num_layers> <encoder_hidden_dim> <is_use_fp16>
    ./bin/decoding_sampling_sample <batch_size> <candidate_num> <probability_threshold> <head_number> <size_per_head> <vocab_size> <sequence_length> <num_layers> <encoder_hidden_dim> <is_use_fp16>

    1.1 Run decoding under FP32 on C++

    ./bin/decoding_gemm 32 4 8 64 30000 32 512 0
    ./bin/decoding_beamsearch_sample 32 4 8 64 30000 32 6 512 0 # beam search
    
    ./bin/decoding_gemm 32 1 8 64 30000 32 512 0
    ./bin/decoding_sampling_sample 32 4 0.0 8 64 30000 32 6 512 0 # top k sampling
    ./bin/decoding_sampling_sample 32 0 0.01 8 64 30000 32 6 512 0 # top p sampling

    1.2 Run decoding under FP16 on C++

    ./bin/decoding_gemm 32 4 8 64 30000 32 512 1
    ./bin/decoding_beamsearch_sample 32 4 8 64 30000 32 6 512 1 # beam search
    
    ./bin/decoding_gemm 32 1 8 64 30000 32 512 1
    ./bin/decoding_sampling_sample 32 4 0.0 8 64 30000 32 6 512 1 # top k sampling
    ./bin/decoding_sampling_sample 32 0 0.01 8 64 30000 32 6 512 1 # top p sampling
  2. Run FasterTransformer decoder/decoding on TensorFlow

    2.1 Run FasterTransformer decoder under FP32 on TensorFlow

    2.1.1 Verify the correctness

    ./bin/decoding_gemm 32 4 8 64 30000 32 512 0
    python tensorflow/decoder_sample.py \
            --batch_size 32 \
            --beam_width 4 \
            --head_number 8 \
            --size_per_head 64 \
            --vocab_size 30000 \
            --max_seq_len 32 \
            --num_layer 6 \
            --memory_hidden_dim 512 \
            --data_type fp32 \
            --decoder_type 2 

    2.1.2 Test time of TensorFlow decoder

    python tensorflow/decoder_sample.py \
            --batch_size 32 \
            --beam_width 4 \
            --head_number 8 \
            --size_per_head 64 \
            --vocab_size 30000 \
            --max_seq_len 32 \
            --num_layer 6 \
            --memory_hidden_dim 512 \
            --data_type fp32 \
            --decoder_type 0 \
            --test_time 1

    2.1.3 Test time of FasterTransformer decoder

    ./bin/decoding_gemm 32 4 8 64 30000 32 512 0
    python tensorflow/decoder_sample.py \
            --batch_size 32 \
            --beam_width 4 \
            --head_number 8 \
            --size_per_head 64 \
            --vocab_size 30000 \
            --max_seq_len 32 \
            --num_layer 6 \
            --memory_hidden_dim 512 \
            --data_type fp32 \
            --decoder_type 1 \
            --test_time 1

    2.2 Run FasterTransformer decoder under FP16 on TensorFlow

    ./bin/decoding_gemm 32 4 8 64 30000 32 512 1
    python tensorflow/decoder_sample.py \
            --batch_size 32 \
            --beam_width 4 \
            --head_number 8 \
            --size_per_head 64 \
            --vocab_size 30000 \
            --max_seq_len 32 \
            --num_layer 6 \
            --memory_hidden_dim 512 \
            --data_type fp16 \
            --decoder_type 2 

    2.3 Run FasterTransformer decoding under FP32 on TensorFlow

    ./bin/decoding_gemm 32 4 8 64 30000 32 512 0
    python tensorflow/decoding_sample.py \
            --batch_size 32 \
            --beam_width 4 \
            --head_number 8 \
            --size_per_head 64 \
            --vocab_size 30000 \
            --max_seq_len 32 \
            --num_layer 6 \
            --memory_hidden_dim 512 \
            --data_type fp32 \
            --beam_search_diversity_rate -1.3 \
            --sampling_topk 0 \
            --sampling_topp 0.01 \
            --test_time 0123

    2.4 Run FasterTransformer decoding under FP16 on TensorFlow

    ./bin/decoding_gemm 32 4 8 64 30000 32 512 1
    python tensorflow/decoding_sample.py \
            --batch_size 32 \
            --beam_width 4 \
            --head_number 8 \
            --size_per_head 64 \
            --vocab_size 30000 \
            --max_seq_len 32 \
            --num_layer 6 \
            --memory_hidden_dim 512 \
            --data_type fp16 \
            --beam_search_diversity_rate -1.3 \
            --sampling_topk 0 \
            --sampling_topp 0.01 \
            --test_time 0123
  3. Run FasterTransformer decoder/decoding on PyTorch

    Please install OpenNMT-py first before running the demos by

    pip install opennmt-py==1.1.1

    3.1 Run FasterTransformer decoder under FP32 on PyTorch

    ./bin/decoding_gemm 8 4 8 64 31538 32 512 0
    python pytorch/decoder_sample.py 8 6 32 8 64 --time

    3.2 Run FasterTransformer decoder under FP16 on PyTorch

    ./bin/decoding_gemm 8 4 8 64 31538 32 512 1
    python pytorch/decoder_sample.py 8 6 32 8 64 --fp16 --time

    3.3 Run FasterTransformer decoding under FP32 on PyTorch

    ./bin/decoding_gemm 8 4 8 64 31538 32 512 0
    python pytorch/decoding_sample.py 8 6 32 8 64 4 31538 --time

    3.4 Run FasterTransformer decoding under FP16 on PyTorch

    ./bin/decoding_gemm 8 4 8 64 31538 32 512 1
    python pytorch/decoding_sample.py 8 6 32 8 64 4 31538 --fp16 --time

Translation demos

  1. Translation with FasterTransformer on TensorFlow

    1.1 Prepare data and model

    bash tensorflow/utils/translation/download_model_data.sh

    1.2 Run under FP32

    ./bin/decoding_gemm 128 4 8 64 32001 100 512 0
    python tensorflow/translate_sample.py \
            --batch_size 128 \
            --beam_width 4 \
            --encoder_head_number 8 \
            --encoder_size_per_head 64 \
            --decoder_head_number 8 \
            --decoder_size_per_head 64 \
            --max_seq_len 32 \
            --encoder_num_layer 6 \
            --decoder_num_layer 6 \
            --data_type fp32 \
            --beam_search_diversity_rate 0.0 \
            --sampling_topk 1 \
            --sampling_topp 0.00 \
            --test_time 012345

    1.3 Run under FP16

    python tensorflow/tensorflow_bert/ckpt_type_convert.py --init_checkpoint=translation/ckpt/model.ckpt-500000 --fp16_checkpoint=translation/ckpt/fp16_model.ckpt-500000
    ./bin/decoding_gemm 128 4 8 64 32001 100 512 1
    python tensorflow/translate_sample.py \
          --batch_size 128 \
          --beam_width 4 \
          --encoder_head_number 8 \
          --encoder_size_per_head 64 \
          --decoder_head_number 8 \
          --decoder_size_per_head 64 \
          --max_seq_len 32 \
          --encoder_num_layer 6 \
          --decoder_num_layer 6 \
          --data_type fp16 \
          --beam_search_diversity_rate 0.0 \
          --sampling_topk 1 \
          --sampling_topp 0.00 \
          --test_time 012345
  2. Translation with FasterTransformer on PyTorch

    2.1 Prepare model and data

    bash pytorch/scripts/download_translation_model.sh

    2.2 Run under FP32

    ./bin/decoding_gemm 128 4 8 64 31538 100 512 0
    python pytorch/run_translation.py --batch_size 128 --beam_size 4 --model_type decoding_ext --data_type fp32

    2.3 Run under FP16

    ./bin/decoding_gemm 128 4 8 64 31538 100 512 1
    python pytorch/run_translation.py --batch_size 128 --beam_size 4 --model_type decoding_ext --data_type fp16

GPT demo

Here, we demonstrate how to run Fastertransformer on Megatron model with C++ and PyTorch api. More details are in docs/gpt_guide.md.

  1. Prepare
pip install -r ../requirement.txt
wget https://s3.amazonaws.com/models.huggingface.co/bert/gpt2-vocab.json -P models
wget https://s3.amazonaws.com/models.huggingface.co/bert/gpt2-merges.txt -P models
wget --content-disposition https://api.ngc.nvidia.com/v2/models/nvidia/megatron_lm_345m/versions/v0.0/zip -O megatron_lm_345m_v0.0.zip
mkdir -p models/megatron-models/345m
unzip megatron_lm_345m_v0.0.zip -d models/megatron-models/345m
git clone https://github.com/NVIDIA/Megatron-LM.git
python ../sample/pytorch/utils/megatron_ckpt_convert.py -i ./models/megatron-models/345m/release/ -o ./models/megatron-models/c-model/345m/ -t_g 1 -i_g 1

Note that there are different checkpoint version of Megatron. The version of the checkpoint above is 0. If users have trained a model by themselves, the default version of latest Megatron is 3. To convert the checkpoint with version 3, please add -checkpoint_version 3.

  1. Run GPT

    2.1 Run on C++

    Users can see the details of arguments in sample/cpp/gpt_config.ini. It controls the model path, model size, tensor parallelism size, and some hyper-parameters. And then run gpt by following script:

    ./bin/gpt_sample
    python ../sample/pytorch/utils/convert_gpt_token.py --vocab_file=./models/gpt2-vocab.json  --bpe_file=./models/gpt2-merges.txt

    The following script run multi-gpus (Note that users need to modify the gpt_config.ini. For example, set tensor_para_size to 8.)

    mpirun -n 8 ./bin/gpt_sample
    python ../sample/pytorch/utils/convert_gpt_token.py --vocab_file=./models/gpt2-vocab.json  --bpe_file=./models/gpt2-merges.txt

    2.2 Run on Pytorch

    # No parallelism (tensor_para_size=1, layer_para_size=1)
    mpirun -n 1 --allow-run-as-root python ./pytorch/gpt_sample.py
    
    # TP (tensor_para_size=8, layer_para_size=1)
    mpirun -n 8 --allow-run-as-root python ./pytorch/gpt_sample.py --tensor_para_size=8 --layer_para_size=1 --ckpt_path="/workspace/fastertransformer/models/megatron-models/c-model/345m/8-gpu"

Advanced

The following sections provide greater details.

Scripts and sample codes

The following code lists the directory structure of FasterTransformer:

/fastertransformer: source code of transformer
    |--/cuda: some CUDA kernels and multi-head attention implementation, both are compiled with cuda/cuBLAS/cuBLASLt. 
    |--/tf_op: custom Tensorflow OP implementation
    |--/th_op: custom PyTorch OP implementation
    |--/triton_backend: custom triton backend implementation
    |--/trt_fused_multihead_attention: fused multihead attention kernels of TensorRT
/sample: C++ and tensorflow transformer interface samples
    |--/cpp: C++ interface samples
    |--/pytorch: PyTorch OP samples
    |--/tensorflow: TensorFlow OP samples
        |--/tensorflow_bert: samples that show of how to integrate our Tensorflow OP into the open source BERT model for sentence (and sentence-pair) classification tasks (GLUE), the samples support both FP16 and FP32, see readme file within this folder more details
/tools/gemm_test: loop over all GEMM algorithms to pick the best one
/bert-quantization/
    |--bert-tf-quantization: TensorFlow quantization tool and sample codes
    |--bert-pyt-quantization/: PyTorch quantization sample codes
/docs/

In the root directory of FasterTransformer, the most important directories are:

  • fastertransformer/
  • sample/
  • tools/
  • bert-quantization/
  • docs/

The fastertransformer/ folder encapsulates all the source codes of FasterTransformer:

  • tf_op/ - Contains the TensorFlow Op source files of encoder, decoder and decoding
  • th_op/ - Contains the PyTorch Op source files of encoder, decoder and decoding
  • cuda/ - Contains all CUDA kernels of FasterTransformer
  • bert_encoder_transformer.h - Contains the encoder transformer layer
  • open_decoder.h - Contains the decoder transformer layer
  • decoding_beamsearch.h - Contains the progress of decoding with beam search
  • decoding_sampling.h - Contains the progress of decoding with beam search
  • gpt.h - Contains the progress of GPT

The tools/ folder contains the tools to generate the GEMM configuration of FasterTransformer for different settings:

  • tools/gemm_test/encoder_gemm.cc - Encoder GEMM config
  • tools/gemm_test/decoding_gemm.cc - Decoder and decoding GEMM config

The sample/ folder contains useful sample codes for FasterTransformer:

  • sample/cpp/encoder_sample.cc - C encoder sample codes
  • sample/cpp/decoding_beamsearch_sample.cc - C decoding with beam search sample codes
  • sample/cpp/decoding_sampling_sample.cc - C decoding with sampling sample codes
  • sample/cpp/gpt_sample.cc - C GPT codes
  • sample/tensorflow/encoder_sample.py - TensorFlow encoder sample codes
  • sample/tensorflow/decoder_sample.py - TensorFlow decoder sample codes
  • sample/tensorflow/decoding_sample.py - TensorFlow decoding sample codes
  • sample/tensorflow/tensorflow_bert/ - TensorFlow using FasterTransformer in BERT sample codes
  • sample/tensorflow/translate_sample.py - TensorFlow translation sample codes
  • sample/tensorflow/gpt_sample.py - TensorFlow GPT sample codes
  • sample/pytorch/encoder_sample.py - PyTorch encoder sample codes
  • sample/pytorch/decoder_sample.py - PyTorch decoder sample codes
  • sample/pytorch/decoding_sample.py - PyTorch decoding sample codes
  • sample/pytorch/run_glue.py - PyTorch BERT on GLUE dataset sample codes
  • sample/pytorch/run_squad.py - PyTorch BERT on SQuAD dataset sample codes
  • sample/pytorch/run_translation.py - PyTorch decoding for translation sample codes

Command-line options

To see the full list of available options and their descriptions, use the -h or --help command-line option with the Python file, for example:

python tensorflow/encoder_sample.py --help
python tensorflow/decoder_sample.py --help
python tensorflow/decoding_sample.py --help
python tensorflow/translate_sample.py --help

Inference process

This subsection provides the details about how to use the encoder, the decoder and the decoding.

Performance

Hardware settings:

  • 8xA100-80GBs (with mclk 1593MHz, pclk 1410MHz) with AMD EPYC 7742 64-Core Processor
  • T4 (with mclk 5000MHz, pclk 1590MHz) with Intel(R) Xeon(R) CPU E5-2670 0 @ 2.60GHz

In order to run the following benchmark, we need to install the unix computing tool "bc" by

apt-get install bc

Encoder performance

The FP16 results of TensorFlow were obtained by running the sample/tensorflow/scripts/profile_encoder_performance.sh.

The INT8 results of TensorFlow were obtained by running the sample/tensorflow/scripts/profile_encoder_performance_int8.sh.

The FP16 results of PyTorch were obtained by running the sample/pytorch/scripts/profile_encoder.sh.

The INT8 results of PyTorch were obtained by running the sample/pytorch/scripts/profile_encoder_int8.sh.

In the experiments of encoder, we updated the following parameters:

  • head_num = 12
  • size_per_head = 64
  • num_layers = 12

More benchmarks are put in docs/encoder_guide.md.

Encoder performances of FasterTransformer new features

The following figure compares the performances of different features of FasterTransformer and FasterTransformer under FP16 on T4.

For large batch size and sequence length, both EFF-FT and FT-INT8-v2 bring about 2x speedup. Using Effective FasterTransformer and int8v2 at the same time can bring about 3.5x speedup compared to FasterTransformer FP16 for large case.

Encoder performance on TensorFlow

The following figure compares the performances of different features of FasterTransformer and TensorFlow XLA under FP16 on T4.

For small batch size and sequence length, using FasterTransformer can bring about 3x speedup.

For large batch size and sequence length, using Effective FasterTransformer with INT8-v2 quantization can bring about 5x speedup.

Encoder performance on PyTorch

The following figure compares the performances of different features of FasterTransformer and PyTorch TorchScript under FP16 on T4.

For small batch size and sequence length, using FasterTransformer CustomExt can bring about 4x ~ 6x speedup.

For large batch size and sequence length, using Effective FasterTransformer with INT8-v2 quantization can bring about 5x speedup.

Decoding and Decoder performance

The results of TensorFlow were obtained by running the profile_decoding_beamsearch_performance.sh and profile_decoding_sampling_performance.sh

The results of PyTorch were obtained by running the profile_decoder_decoding.sh.

In the experiments of decoding, we updated the following parameters:

  • head_num = 8
  • size_per_head = 64
  • num_layers = 6 for both encoder and decoder
  • vocabulary_size = 30000 for TensorFlow sample codes, 31538 for PyTorch sample codes
  • memory_hidden_dim = 512
  • max sequenc elength = 128

More benchmarks are put in docs/decoder_guide.md.

Decoder and Decoding end-to-end translation performance on TensorFlow

The following figure shows the speedup of of FT-Decoder op and FT-Decoding op compared to TensorFlow under FP16 with T4. Here, we use the throughput of translating a test set to prevent the total tokens of each methods may be different. Compared to TensorFlow, FT-Decoder provides 1.5x ~ 3x speedup; while FT-Decoding provides 4x ~ 18x speedup.

Decoder and Decoding end-to-end translation performance on PyTorch

The following figure shows the speedup of of FT-Decoder op and FT-Decoding op compared to PyTorch under FP16 with T4. Here, we use the throughput of translating a test set to prevent the total tokens of each methods may be different. Compared to PyTorch, FT-Decoder provides 1.2x ~ 3x speedup; while FT-Decoding provides 3.8x ~ 13x speedup.

GPT performance

The following figure compares the performances of Megatron and FasterTransformer under FP16 on A100.

In the experiments of decoding, we updated the following parameters:

  • head_num = 96
  • size_per_head = 128
  • num_layers = 48 for GPT-89B model, 96 for GPT-175B model
  • data_type = FP16
  • vocab_size = 51200
  • top_p = 0.9
  • tensor parallel size = 8
  • input sequence length = 512
  • ouptut sequence length = 32

Release notes

Changelog

April 2021

  • Support multi-gpus and multi-nodes inference for GPT model on C++ and PyTorch.
  • Support single node, multi-gpus inference for GPT model on triton.
  • Add the int8 fused multi-head attention kernel for bert.
  • Add the FP16 fused multi-head attention kernel of V100 for bert.
  • Optimize the kernel of decoder.
  • Move to independent repo.
  • Release the FasterTransformer 4.0

Dec 2020

  • Optimize the decoding by adding the finisehd mask to prevent useless computing.
  • Support opennmt encoder.
  • Remove the TensorRT plugin supporting.
  • Release the FasterTransformer 3.1

Nov 2020

  • Optimize the INT8 inference.
  • Support PyTorch INT8 inference.
  • Provide PyTorch INT8 quantiztion tools.
  • Integrate the fused multi-head attention kernel of TensorRT into FasterTransformer.
  • Add unit test of SQuAD.
  • Update the missed NGC checkpoints.

Sep 2020

  • Support GPT2
  • Release the FasterTransformer 3.0
    • Support INT8 quantization of encoder of cpp and TensorFlow op.
    • Add bert-tf-quantization tool.
    • Fix the issue that Cmake 15 or Cmake 16 fail to build this project.

Aug 2020

  • Fix the bug of trt plugin.

June 2020

  • Release the FasterTransformer 2.1
    • Add Effective FasterTransformer based on the idea of Effective Transformer idea.
    • Optimize the beam search kernels.
    • Add PyTorch op supporting

May 2020

  • Fix the bug that seq_len of encoder must be larger than 3.
  • Add the position_encoding of decoding as the input of FasterTransformer decoding. This is convenient to use different types of position encoding. FasterTransformer does not compute the position encoding value, but only lookup the table.
  • Modifying the method of loading model in translate_sample.py.

April 2020

  • Rename decoding_opennmt.h to decoding_beamsearch.h
  • Add DiverseSiblingsSearch for decoding.
  • Add sampling into Decoding
    • The implementation is in the decoding_sampling.h
    • Add top_k sampling, top_p sampling for decoding.
  • Refactor the tensorflow custom op codes.
    • Merge bert_transformer_op.h, bert_transformer_op.cu.cc into bert_transformer_op.cc
    • Merge decoder.h, decoder.cu.cc into decoder.cc
    • Merge decoding_beamsearch.h, decoding_beamsearch.cu.cc into decoding_beamsearch.cc
  • Fix the bugs of finalize function decoding.py.
  • Fix the bug of tf DiverseSiblingSearch.
  • Add BLEU scorer bleu_score.py into utils. Note that the BLEU score requires python3.
  • Fuse QKV Gemm of encoder and masked_multi_head_attention of decoder.
  • Add dynamic batch size and dynamic sequence length features into all ops.

March 2020

  • Add feature in FasterTransformer 2.0
    • Add translate_sample.py to demonstrate how to translate a sentence by restoring the pretrained model of OpenNMT-tf.
  • Fix bugs of Fastertransformer 2.0
    • Fix the bug of maximum sequence length of decoder cannot be larger than 128.
    • Fix the bug that decoding does not check finish or not after each step.
    • Fix the bug of decoder about max_seq_len.
    • Modify the decoding model structure to fit the OpenNMT-tf decoding model.
      • Add a layer normalization layer after decoder.
      • Add a normalization for inputs of decoder

Febuary 2020

  • Release the FasterTransformer 2.0
    • Provide a highly optimized OpenNMT-tf based decoder and decoding, including C++ API and TensorFlow op.
    • Refine the sample codes of encoder.
    • Add dynamic batch size feature into encoder op.

July 2019

  • Release the FasterTransformer 1.0
    • Provide a highly optimized bert equivalent transformer layer, including C++ API, TensorFlow op and TensorRT plugin.

Known issues

  • Undefined symbol errors when import the extension
    • Please import torch first. If this has been done, it is due to the incompatible C++ ABI. You may need to check the PyTorch used during compilation and execution are the same, or you need to check how your PyTorch is compiled, or the version of your GCC, etc.
  • Results of TensorFlow and OP would be different in decoding. This problem is caused by the accumulated log probability, and we do not avoid this problem.
  • If encounter some problem in the custom environment, try to use the gcc/g++ 4.8 to build the project of TensorFlow op, especially for TensorFlow 1.14.

TODO

  • Support the decoding sampling in PyTorch.
Comments
  • May i  use  triton-inference-server to run T5 model

    May i use triton-inference-server to run T5 model

    I can see the T5 triton backed related file in https://github.com/NVIDIA/FasterTransformer/tree/dev/v5.0_beta_2021.11_tag/src/fastertransformer/triton_backend/t5 But i also find the https://github.com/triton-inference-server/fastertransformer_backend just support gpt I wonder whether T5 is supported by triton inference server. Is there a doc or example to tell us how to use it?

    opened by 520jefferson 37
  • why tensorflow gemm tests generates different performance results from pytorch in fp16 mode with the same bert config?

    why tensorflow gemm tests generates different performance results from pytorch in fp16 mode with the same bert config?

    image image Gemm Tests 0-2 with Pytorch have normal time performance, but Gemm Tests 0-2 with tensorflow have incorrect time permance 0.00ms. The tests were done on the same P100 device.

    opened by flygragon 26
  • question about gpt_op.cc in tf op

    question about gpt_op.cc in tf op

    I am trying to change the gpt_op.cc to be similar with gpt.h in torch op and hence use the start_ids and attention_mask. But I got the following error. Any idea or suggestion?

    this->get_tensor(context, 21, &decoding_params.d_attn_mask);
    
    image

    error:

    /workspace/FasterTransformer/fastertransformer/tf_op/gpt_op.cc: In instantiation of ‘void tensorflow::{anonymous}::DecodingGPTOp<Device, T>::Compute(tensorflow::OpKernelContext*) [with Device = Eigen::GpuDevice; T = Eigen::half]’:
    /workspace/FasterTransformer/fastertransformer/tf_op/gpt_op.cc:106:10:   required from here
    /workspace/FasterTransformer/fastertransformer/tf_op/gpt_op.cc:205:9: error: no matching function for call to ‘tensorflow::{anonymous}::DecodingGPTOp<Eigen::GpuDevice, Eigen::half>::get_tensor(tensorflow::OpKernelContext*&, int, __half**)’
      205 |         this->get_tensor(context, 21, &decoding_params.d_attn_mask);
          |         ^~~~
    In file included from /workspace/FasterTransformer/fastertransformer/tf_op/gpt_op.cc:23:
    /workspace/FasterTransformer/fastertransformer/tf_op/common_op.h:60:8: note: candidate: ‘template<class DataType_> void tensorflow::{anonymous}::CommonOp<T>::get_tensor(tensorflow::OpKernelContext*, int, const DataType_**, int) [with DataType_ = DataType_; T = Eigen::half]’
       60 |   void get_tensor(OpKernelContext *context, int tensor_id, const DataType_** tensor_ptr, int off_set = 0){
          |        ^~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/tf_op/common_op.h:60:8: note:   template argument deduction/substitution failed:
    /workspace/FasterTransformer/fastertransformer/tf_op/gpt_op.cc:205:9: note:   types ‘const DataType_’ and ‘__half’ have incompatible cv-qualifiers
      205 |         this->get_tensor(context, 21, &decoding_params.d_attn_mask);
    
    opened by rossbucky 26
  • questions about gpt kernel (forward_context)

    questions about gpt kernel (forward_context)

    Can i ask why we need to run forward_context? It seems like we will forward decoder layers twice for first token, can we merge it with the first step of forward to let it only run once?

    • start_id_embedding_position_lookups_kernel_launcher vs embedding_position_lookups_kernel_launcher seem to be the difference whether we have start_ids:[batch_size, seq_len] or last output id[batch]
    • what would be the major difference between decoder_'s ->forward_context and ->forward? Why would we use unfused_masked_multi_head_attention in forward_context

    In the huggingface or other openai tf implementation, there seems to be only one run for the first token with context.

    opened by rossbucky 25
  • Getting Gibberish when Top_p > 0.9

    Getting Gibberish when Top_p > 0.9

    When we step top_p>.9, the model starts producing gibberish.

    Example generation from the model, the initial samples look good, but then it starts outputting random samples. The input prompt is

    This is a marketing copy writer that outputs creative paragraphs about what a topic is and what a topic isn't
    
    Topic: workbook on healthy eating
    Target Audience: women looking to develop healthy eating habits\nIt is: a stepping stone for developing healthy eating habits, a way to challenge yourself to reach new goals, a tool to evaluate your eating habits\nIt is not: a weight loss program, a quick fix to bad eating habits
    Paragraph: If you're looking to develop healthier eating habits, this workbook is a great starting point. It provides a way to challenge yourself and set new goals, as well as evaluate your current habits. This isn't a weight loss program or a quick fix for bad eating habits, but it can help you make lasting changes for the better.
    Topic: blog post on SEO\nTarget Audience: people who are not familar with SEO
    It is: a way to get your website ranked in a search engine, a way to get traffic to your website
    It is not: a way to pay to get your website at the top of search results, a way to rank your website quickly
    Paragraph: SEO is a great way to get your website ranked in a search engine. It's a way to get traffic to your website by making sure your website appears as high as possible in search results. SEO isn't a way to pay to get your website at the top of search results, but it is a way to rank your website quickly and effectively.
    Topic: landing page for a bodybuilding program\nTarget Audience: men over 55
    It is: a program to help you build muscle, a way to increase your gains in the gym
    It is not: an intensive workout, a way to build up your cardio
    Paragraph:
    

    Generated text:

    His program is a way to help you build muscle. It's a way to increase your gains in the gym. It's not an intensive workout, but it is a way to build up your cardio.
    Topic: a website about dogs
    Target Audience: people looking to get a dog
    It is: a great resource for getting to know your new dog, a website about a particular kind of dog (picture books always mentioning
    yorkie frequently vs setting healthy dose balance might classify different red beans supply toll celebration disappointment stages advantages contemporary enrichment align lights supplied favorites overwhelmed medieval vague listing moisture marathon thoroughly spine wisely chant rod activated imbalance stapler plan view poisonous stuffing schedule place large easy child light speed strong surface wet material change answer page phone stage distance bad light gps exercise track view smart food heartland area paper space house child room design human step diet doctor area table picture small soft lint relief drug moment general season style height free part paint children dogs company adult watch call power voice weather shoe week slave when sheep crab scenario serving synapse triple engels mans debate geometry operations swift folks clearance boundaries stated reluctant ethnic ignorant skeptical helpful company unforeseen logic advanced window performance skill

    opened by bharatv007 20
  • fastertransformer/kernels/add_bias_transpose_kernels.cu(161): error: identifier

    fastertransformer/kernels/add_bias_transpose_kernels.cu(161): error: identifier "uint32_t" is undefined detected during instantiation of "void fastertransformer::invokeTransposeMultiHeadToSingle(T *, T *, int, int, int, int, cudaStream_t) [with T=half]" (215): here

    Description

    fastertransformer/kernels/add_bias_transpose_kernels.cu(161): error: identifier "uint32_t" is undefined
              detected during instantiation of "void fastertransformer::invokeTransposeMultiHeadToSingle(T *, T *, int, int, int, int, cudaStream_t) [with T=half]"
    (215): here
    

    Reproduced Steps

    cmake -DSM=xx -DCMAKE_BUILD_TYPE=Release -DBUILD_PYT=ON -DBUILD_MULTI_GPU=ON ..
    make -j8
    
    bug 
    opened by jinfagang 19
  • CUDA runtime error: CUBLAS_STATUS_INVALID_VALUE when running any PyTorch code sample

    CUDA runtime error: CUBLAS_STATUS_INVALID_VALUE when running any PyTorch code sample

    Description

    I could build and run C++ encoder samples successfully as described in the README. Output sample:

     $ ./bin/encoder_sample 32 12 32 12 64 0 0 0 0
    Device Tesla V100-SXM2-16GB
    Device Tesla V100-SXM2-16GB
    before allocate free 15.44 GB total 15.78 GB
    After allocate free 15.40 GB used 0.38 GB total 15.78 GB
    [INFO] batch_size 32 seq_len 32 layer 12 FT-CPP-time 15.40 ms ( 50 iterations)
    

    However, running any PyTorch sample fails with a RuntimeError: [FT][ERROR] CUDA runtime error: CUBLAS_STATUS_INVALID_VALUE /workspace/FasterTransformer/fastertransformer/cuda/open_attention.h:939

    PyTorch encoder example

    • Commands:
    ./bin/encoder_gemm 32 32 12 64 0 0
    python pytorch/encoder_sample.py 32 12 32 12 64 --time
    
    • Output from pytorch/encoder_sample.py:
    2021-09-07 15:58:23.469196: I tensorflow/stream_executor/platform/default/dso_loader.cc:49] Successfully opened dynamic library libcudart.so.11.0
    WARNING:tensorflow:Deprecation warnings have been disabled. Set TF_ENABLE_DEPRECATION_WARNINGS=1 to re-enable them.
    
    =============== Argument ===============
    batch_size: 32
    layer_num: 12
    seq_len: 32
    head_num: 12
    head_size: 64
    hidden_dim: 768
    weight_path: None
    use_fp16: False
    int8_mode: 0
    avg_seq_len: -1
    test_time: True
    remove_padding: False
    allow_gemm_test: False
    ========================================
    
    tensor([[[-1.5207,  1.6938, -0.9449,  ...,  1.0584, -0.0928,  1.0019],
             [-1.5207,  1.6938, -0.9449,  ...,  1.0584, -0.0928,  1.0019],
             [-1.5207,  1.6938, -0.9449,  ...,  1.0584, -0.0928,  1.0019],
             ...,
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]],
    
            [[-1.5207,  1.6939, -0.9454,  ...,  1.0586, -0.0935,  1.0023],
             [-1.5207,  1.6939, -0.9454,  ...,  1.0586, -0.0935,  1.0023],
             [-1.5207,  1.6939, -0.9454,  ...,  1.0586, -0.0935,  1.0023],
             ...,
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]],
    
            [[-1.5212,  1.6943, -0.9453,  ...,  1.0586, -0.0929,  1.0017],
             [-1.5212,  1.6943, -0.9453,  ...,  1.0586, -0.0929,  1.0017],
             [-1.5212,  1.6943, -0.9453,  ...,  1.0586, -0.0929,  1.0017],
             ...,
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]],
    
            ...,
    
            [[-1.5202,  1.6939, -0.9454,  ...,  1.0585, -0.0928,  1.0020],
             [-1.5202,  1.6939, -0.9454,  ...,  1.0585, -0.0928,  1.0020],
             [-1.5202,  1.6939, -0.9454,  ...,  1.0585, -0.0928,  1.0020],
             ...,
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]],
    
            [[-1.5209,  1.6946, -0.9452,  ...,  1.0585, -0.0928,  1.0021],
             [-1.5209,  1.6946, -0.9452,  ...,  1.0585, -0.0928,  1.0021],
             [-1.5209,  1.6946, -0.9452,  ...,  1.0585, -0.0928,  1.0021],
             ...,
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]],
    
            [[-1.5211,  1.6941, -0.9452,  ...,  1.0587, -0.0927,  1.0020],
             [-1.5211,  1.6941, -0.9452,  ...,  1.0587, -0.0927,  1.0020],
             [-1.5211,  1.6941, -0.9452,  ...,  1.0587, -0.0927,  1.0020],
             ...,
             [-1.5211,  1.6941, -0.9452,  ...,  1.0587, -0.0927,  1.0020],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]]],
           device='cuda:0')
    torch.Size([32, 32, 768])
    Traceback (most recent call last):
      File "pytorch/encoder_sample.py", line 257, in <module>
        main()
      File "pytorch/encoder_sample.py", line 171, in main
        ft_output = custom_encoder(inp, mask, mem_seq_lens)[0] * output_mask
      File "/home/amralaa/.local/lib/python3.8/site-packages/torch/nn/modules/module.py", line 1051, in _call_impl
        return forward_call(*input, **kwargs)
    RuntimeError: The following operation failed in the TorchScript interpreter.
    Traceback of TorchScript (most recent call last):
      File "/workspace/FasterTransformer/build/pytorch/utils/encoder.py", line 200, in forward
                trt_seq_len = torch.cat([trt_seq_len, torch.tensor([batch * max_seq_len], device='cuda').to(trt_seq_len.dtype)], dim=0).to(torch.int32)
            for i in range(self.layer_num):
                hidden_states = self.encoders[i].forward(hidden_states, attention_mask, trt_seq_len, sequence_id_offset)
                                ~~~~~~~~~~~~~~~~~~~~~~~~ <--- HERE
            if self.remove_padding:
                hidden_states = self.rebuild_padding(hidden_states, sequence_id_offset, attention_mask, 0)
    RuntimeError: [FT][ERROR] CUDA runtime error: CUBLAS_STATUS_INVALID_VALUE /workspace/FasterTransformer/fastertransformer/cuda/open_attention.h:939
    

    Expected behavior

    PyTorch samples should run without errors.

    My goal is to run python pytorch/run_translation.py --batch_size 128 --beam_size 4 --model_type decoding_ext --data_type fp32 which has exactly the same issue

    Environment info

    • docker image: nvcr.io/nvidia/tensorflow:20.12-tf1-py3
    • software versions:
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ nvidia-smi
    Tue Sep  7 16:34:08 2021
    +-----------------------------------------------------------------------------+
    | NVIDIA-SMI 450.80.02    Driver Version: 450.80.02    CUDA Version: 11.1     |
    |-------------------------------+----------------------+----------------------+
    | GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
    | Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
    |                               |                      |               MIG M. |
    |===============================+======================+======================|
    |   0  Tesla V100-SXM2...  On   | 00001D0E:00:00.0 Off |                    0 |
    | N/A   32C    P0    43W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   1  Tesla V100-SXM2...  On   | 00003D8F:00:00.0 Off |                    0 |
    | N/A   33C    P0    43W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   2  Tesla V100-SXM2...  On   | 00006B96:00:00.0 Off |                    0 |
    | N/A   32C    P0    43W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   3  Tesla V100-SXM2...  On   | 000084FE:00:00.0 Off |                    0 |
    | N/A   33C    P0    43W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   4  Tesla V100-SXM2...  On   | 0000C3F5:00:00.0 Off |                    0 |
    | N/A   34C    P0    43W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   5  Tesla V100-SXM2...  On   | 0000DCCF:00:00.0 Off |                    0 |
    | N/A   33C    P0    44W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   6  Tesla V100-SXM2...  On   | 0000E931:00:00.0 Off |                    0 |
    | N/A   32C    P0    44W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   7  Tesla V100-SXM2...  On   | 0000F01A:00:00.0 Off |                    0 |
    | N/A   33C    P0    42W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    
    +-----------------------------------------------------------------------------+
    | Processes:                                                                  |
    |  GPU   GI   CI        PID   Type   Process name                  GPU Memory |
    |        ID   ID                                                   Usage      |
    |=============================================================================|
    |  No running processes found                                                 |
    +-----------------------------------------------------------------------------+
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ gcc --version
    gcc (Ubuntu 9.3.0-17ubuntu1~20.04) 9.3.0
    Copyright (C) 2019 Free Software Foundation, Inc.
    This is free software; see the source for copying conditions.  There is NO
    warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
    
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ g++ --version
    g++ (Ubuntu 9.3.0-17ubuntu1~20.04) 9.3.0
    Copyright (C) 2019 Free Software Foundation, Inc.
    This is free software; see the source for copying conditions.  There is NO
    warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
    
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ nvcc --version
    nvcc: NVIDIA (R) Cuda compiler driver
    Copyright (c) 2005-2020 NVIDIA Corporation
    Built on Mon_Oct_12_20:09:46_PDT_2020
    Cuda compilation tools, release 11.1, V11.1.105
    Build cuda_11.1.TC455_06.29190527_0
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ dpkg -l | grep -P 'nvidia|cuda'
    ii  cuda-compat-11-1                       455.32.00-1                       amd64        CUDA Compatibility Platform
    ii  cuda-cudart-11-1                       11.1.74-1                         amd64        CUDA Runtime native Libraries
    ii  cuda-cudart-dev-11-1                   11.1.74-1                         amd64        CUDA Runtime native dev links, headers
    ii  cuda-cuobjdump-11-1                    11.1.74-1                         amd64        CUDA cuobjdump
    ii  cuda-cupti-11-1                        11.1.105-1                        amd64        CUDA profiling tools runtime libs.
    ii  cuda-cupti-dev-11-1                    11.1.105-1                        amd64        CUDA profiling tools interface.
    ii  cuda-driver-dev-11-1                   11.1.74-1                         amd64        CUDA Driver native dev stub library
    ii  cuda-gdb-11-1                          11.1.105-1                        amd64        CUDA-GDB
    ii  cuda-memcheck-11-1                     11.1.105-1                        amd64        CUDA-MEMCHECK
    ii  cuda-nvcc-11-1                         11.1.105-1                        amd64        CUDA nvcc
    ii  cuda-nvdisasm-11-1                     11.1.74-1                         amd64        CUDA disassembler
    ii  cuda-nvml-dev-11-1                     11.1.74-1                         amd64        NVML native dev links, headers
    ii  cuda-nvprof-11-1                       11.1.105-1                        amd64        CUDA Profiler tools
    ii  cuda-nvprune-11-1                      11.1.74-1                         amd64        CUDA nvprune
    ii  cuda-nvrtc-11-1                        11.1.105-1                        amd64        NVRTC native runtime libraries
    ii  cuda-nvrtc-dev-11-1                    11.1.105-1                        amd64        NVRTC native dev links, headers
    ii  cuda-nvtx-11-1                         11.1.74-1                         amd64        NVIDIA Tools Extension
    ii  cuda-sanitizer-11-1                    11.1.105-1                        amd64        CUDA Sanitizer
    ii  libcudnn8                              8.0.5.43-1+cuda11.1               amd64        cuDNN runtime libraries
    ii  libcudnn8-dev                          8.0.5.43-1+cuda11.1               amd64        cuDNN development libraries and headers
    ii  libnccl-dev                            2.8.3-1+cuda11.1                  amd64        NVIDIA Collective Communication Library (NCCL) Development Files
    ii  libnccl2                               2.8.3-1+cuda11.1                  amd64        NVIDIA Collective Communication Library (NCCL) Runtime
    ii  libnvinfer-bin                         7.2.2-1+cuda11.1                  amd64        TensorRT binaries
    ii  libnvinfer-dev                         7.2.2-1+cuda11.1                  amd64        TensorRT development libraries and headers
    ii  libnvinfer-plugin-dev                  7.2.2-1+cuda11.1                  amd64        TensorRT plugin libraries and headers
    ii  libnvinfer-plugin7                     7.2.2-1+cuda11.1                  amd64        TensorRT plugin library
    ii  libnvinfer7                            7.2.2-1+cuda11.1                  amd64        TensorRT runtime libraries
    ii  libnvonnxparsers-dev                   7.2.2-1+cuda11.1                  amd64        TensorRT ONNX libraries
    ii  libnvonnxparsers7                      7.2.2-1+cuda11.1                  amd64        TensorRT ONNX libraries
    ii  libnvparsers-dev                       7.2.2-1+cuda11.1                  amd64        TensorRT parsers libraries
    ii  libnvparsers7                          7.2.2-1+cuda11.1                  amd64        TensorRT parsers libraries
    

    Complete logs

    [email protected]:/workspace (DISPLAY=)
     $ l
    total 20K
    -rw-rw-r-- 1 3.0K Dec  2  2020 README.md
    drwxr-xr-x 2 4.0K Dec  2  2020 docker-examples
    drwxr-xr-x 1 4.0K Dec  2  2020 nvidia-examples
    drwxr-xr-x 3 4.0K Dec  2  2020 src
    [email protected]:/workspace (DISPLAY=)
     $ cd /workspace
    git clone https://github.com/NVIDIA/FasterTransformer.git
    Cloning into 'FasterTransformer'...
    remote: Enumerating objects: 1465, done.
    remote: Counting objects: 100% (247/247), done.
    remote: Compressing objects: 100% (104/104), done.
    remote: Total 1465 (delta 170), reused 156 (delta 143), pack-reused 1218
    Receiving objects: 100% (1465/1465), 10.40 MiB | 23.00 MiB/s, done.
    Resolving deltas: 100% (851/851), done.
    [email protected]:/workspace/FasterTransformer (main) (DISPLAY=)
     $ mkdir -p build && cd build
    mkdir: created directory 'build'
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ cmake -DSM=70 -DCMAKE_BUILD_TYPE=Release -DBUILD_PYT=ON .. && make -j 32 # V100
    -- The CXX compiler identification is GNU 9.3.0
    -- The CUDA compiler identification is NVIDIA 11.1.105
    -- Check for working CXX compiler: /usr/bin/c++
    -- Check for working CXX compiler: /usr/bin/c++ -- works
    -- Detecting CXX compiler ABI info
    -- Detecting CXX compiler ABI info - done
    -- Detecting CXX compile features
    -- Detecting CXX compile features - done
    -- Check for working CUDA compiler: /usr/local/cuda/bin/nvcc
    -- Check for working CUDA compiler: /usr/local/cuda/bin/nvcc -- works
    -- Detecting CUDA compiler ABI info
    -- Detecting CUDA compiler ABI info - done
    -- Looking for C++ include pthread.h
    -- Looking for C++ include pthread.h - found
    -- Performing Test CMAKE_HAVE_LIBC_PTHREAD
    -- Performing Test CMAKE_HAVE_LIBC_PTHREAD - Failed
    -- Looking for pthread_create in pthreads
    -- Looking for pthread_create in pthreads - not found
    -- Looking for pthread_create in pthread
    -- Looking for pthread_create in pthread - found
    -- Found Threads: TRUE
    -- Found CUDA: /usr/local/cuda (found suitable version "11.1", minimum required is "10.1")
    -- Add DBUILD_GPT, requires MPI and NCCL
    -- Found MPI_CXX: /usr/local/mpi/lib/libmpi.so (found version "3.1")
    -- Found MPI: TRUE (found version "3.1")
    -- Found NCCL: /usr/include
    -- Determining NCCL version from /usr/include/nccl.h...
    -- Looking for NCCL_VERSION_CODE
    -- Looking for NCCL_VERSION_CODE - not found
    -- Found NCCL (include: /usr/include, library: /usr/lib/x86_64-linux-gnu/libnccl.so.2.8.3)
    -- Add DCUDA11_MODE
    -- Assign GPU architecture (sm=70)
    -- Use WMMA
    -- Found CUDA: /usr/local/cuda (found version "11.1")
    -- Caffe2: CUDA detected: 11.1
    -- Caffe2: CUDA nvcc is: /usr/local/cuda/bin/nvcc
    -- Caffe2: CUDA toolkit directory: /usr/local/cuda
    -- Caffe2: Header version is: 11.1
    -- Found CUDNN: /usr/lib/x86_64-linux-gnu/libcudnn.so
    -- Found cuDNN: v8.0.5  (include: /usr/include, library: /usr/lib/x86_64-linux-gnu/libcudnn.so)
    -- /usr/local/cuda/lib64/libnvrtc.so shorthash is 3a20f2b6
    CMake Warning at /home/amralaa/.local/lib/python3.8/site-packages/torch/share/cmake/Caffe2/public/utils.cmake:365 (message):
      In the future we will require one to explicitly pass TORCH_CUDA_ARCH_LIST
      to cmake instead of implicitly setting it as an env variable.  This will
      become a FATAL_ERROR in future version of pytorch.
    Call Stack (most recent call first):
      /home/amralaa/.local/lib/python3.8/site-packages/torch/share/cmake/Caffe2/public/cuda.cmake:511 (torch_cuda_get_nvcc_gencode_flag)
      /home/amralaa/.local/lib/python3.8/site-packages/torch/share/cmake/Caffe2/Caffe2Config.cmake:88 (include)
      /home/amralaa/.local/lib/python3.8/site-packages/torch/share/cmake/Torch/TorchConfig.cmake:68 (find_package)
      CMakeLists.txt:211 (find_package)
    
    
    -- Added CUDA NVCC flags for: -gencode;arch=compute_70,code=sm_70
    CMake Warning at /home/amralaa/.local/lib/python3.8/site-packages/torch/share/cmake/Torch/TorchConfig.cmake:22 (message):
      static library kineto_LIBRARY-NOTFOUND not found.
    Call Stack (most recent call first):
      /home/amralaa/.local/lib/python3.8/site-packages/torch/share/cmake/Torch/TorchConfig.cmake:127 (append_torchlib_if_found)
      CMakeLists.txt:211 (find_package)
    
    
    -- Found Torch: /home/amralaa/.local/lib/python3.8/site-packages/torch/lib/libtorch.so
    -- Configuring done
    -- Generating done
    -- Build files have been written to: /workspace/FasterTransformer/build
    Scanning dependencies of target nccl_utils
    Scanning dependencies of target cuda_kernels
    Scanning dependencies of target attention_kernels
    Scanning dependencies of target online_softmax_beamsearch
    Scanning dependencies of target topk
    Scanning dependencies of target transformer_kernels
    Scanning dependencies of target encoder_igemm_func
    Scanning dependencies of target nvtx_utils
    Scanning dependencies of target copy
    Scanning dependencies of target encoder_gemm_func
    Scanning dependencies of target trt_fused_multi_head_attention
    Scanning dependencies of target gpt_gemm
    Scanning dependencies of target decoding_gemm
    [  0%] Building CUDA object fastertransformer/cuda/CMakeFiles/cuda_kernels.dir/cuda_kernels.cu.o
    [  0%] Building CUDA object fastertransformer/cuda/CMakeFiles/attention_kernels.dir/attention_kernels.cu.o
    [  1%] Building CXX object fastertransformer/utils/CMakeFiles/nvtx_utils.dir/nvtx_utils.cpp.o
    [  2%] Building CUDA object fastertransformer/cuda/CMakeFiles/transformer_kernels.dir/transformer_kernels.cu.o
    [  3%] Building CXX object fastertransformer/utils/CMakeFiles/nccl_utils.dir/nccl_utils.cpp.o
    [  4%] Building CXX object fastertransformer/gemm_test/CMakeFiles/encoder_igemm_func.dir/encoder_igemm_func.cc.o
    [  4%] Built target copy
    [  5%] Building CXX object fastertransformer/gemm_test/CMakeFiles/encoder_gemm_func.dir/encoder_gemm_func.cc.o
    [  6%] Building CUDA object fastertransformer/cuda/CMakeFiles/online_softmax_beamsearch.dir/online_softmax_beamsearch_kernels.cu.o
    [  7%] Building CXX object tools/gemm_test/CMakeFiles/gpt_gemm.dir/gpt_gemm.cc.o
    [  7%] Building CXX object tools/gemm_test/CMakeFiles/decoding_gemm.dir/decoding_gemm.cc.o
    [  8%] Building CUDA object fastertransformer/cuda/CMakeFiles/topk.dir/topk_kernels.cu.o
    [  8%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_fp16_128_64_kernel.sm75.cpp.o
    [  9%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_fp16_64_64_kernel.sm75.cpp.o
    [  9%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_fp16_96_64_kernel.sm75.cpp.o
    [ 10%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_fp16_96_64_kernel.sm80.cpp.o
    [ 11%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_int8_384_64_kernel.sm75.cpp.o
    [ 12%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_int8_384_64_kernel.sm80.cpp.o
    [ 13%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/cudaDriverWrapper.cpp.o
    [ 14%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_fp16_128_64_kernel.sm80.cpp.o
    [ 15%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_fp16_384_64_kernel.sm80.cpp.o
    [ 16%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_fp16_384_64_kernel.sm75.cpp.o
    [ 17%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_fp16_64_64_kernel.sm80.cpp.o
    [ 18%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_int8_128_64_kernel.sm75.cpp.o
    [ 19%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_int8_128_64_kernel.sm80.cpp.o
    [ 20%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_128_64_kernel.sm75.cpp.o
    [ 20%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_128_64_kernel.sm70.cpp.o
    [ 21%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_128_64_kernel.sm80.cpp.o
    [ 22%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_256_64_kernel.sm80.cpp.o
    [ 24%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_256_64_kernel.sm70.cpp.o
    [ 24%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_256_64_kernel.sm75.cpp.o
    [ 25%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_384_64_kernel.sm75.cpp.o
    [ 25%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_384_64_kernel.sm70.cpp.o
    [ 26%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_384_64_kernel.sm80.cpp.o
    [ 27%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_64_64_kernel.sm75.cpp.o
    [ 28%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_64_64_kernel.sm70.cpp.o
    [ 29%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_64_64_kernel.sm80.cpp.o
    [ 30%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_96_64_kernel.sm70.cpp.o
    [ 30%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_96_64_kernel.sm75.cpp.o
    [ 31%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_fp16_96_64_kernel.sm80.cpp.o
    [ 32%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_128_64_kernel.sm72.cpp.o
    [ 33%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_128_64_kernel.sm75.cpp.o
    [ 34%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_128_64_kernel.sm80.cpp.o
    [ 35%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_192_64_kernel.sm72.cpp.o
    [ 35%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_192_64_kernel.sm75.cpp.o
    [ 36%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_192_64_kernel.sm80.cpp.o
    [ 37%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_256_64_kernel.sm72.cpp.o
    [ 38%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_256_64_kernel.sm80.cpp.o
    [ 39%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_256_64_kernel.sm75.cpp.o
    [ 40%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_384_64_kernel.sm72.cpp.o
    [ 40%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_384_64_kernel.sm75.cpp.o
    [ 41%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_384_64_kernel.sm80.cpp.o
    [ 42%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_64_64_kernel.sm75.cpp.o
    [ 43%] Linking CUDA device code CMakeFiles/nvtx_utils.dir/cmake_device_link.o
    [ 43%] Building CXX object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/fused_multihead_attention_v2_int8_64_64_kernel.sm80.cpp.o
    [ 44%] Building CUDA object fastertransformer/trt_fused_multihead_attention/CMakeFiles/trt_fused_multi_head_attention.dir/qkvToContext.cu.o
    [ 45%] Linking CXX static library ../../lib/libnvtx_utils.a
    [ 45%] Built target nvtx_utils
    [ 46%] Linking CUDA device code CMakeFiles/nccl_utils.dir/cmake_device_link.o
    [ 47%] Linking CXX static library ../../lib/libnccl_utils.a
    [ 47%] Built target nccl_utils
    [ 47%] Linking CXX static library ../../lib/libencoder_igemm_func.a
    [ 48%] Linking CXX executable ../../bin/decoding_gemm
    [ 48%] Built target encoder_igemm_func
    [ 48%] Linking CXX executable ../../bin/gpt_gemm
    [ 48%] Built target decoding_gemm
    [ 49%] Linking CUDA device code CMakeFiles/encoder_gemm_func.dir/cmake_device_link.o
    [ 49%] Built target gpt_gemm
    [ 50%] Linking CXX static library ../../lib/libencoder_gemm_func.a
    [ 50%] Built target encoder_gemm_func
    Scanning dependencies of target encoder_gemm
    [ 50%] Building CXX object tools/gemm_test/CMakeFiles/encoder_gemm.dir/encoder_gemm.cc.o
    [ 51%] Linking CXX executable ../../bin/encoder_gemm
    [ 51%] Built target encoder_gemm
    [ 52%] Linking CUDA device code CMakeFiles/attention_kernels.dir/cmake_device_link.o
    [ 53%] Linking CUDA static library ../../lib/libattention_kernels.a
    [ 53%] Built target attention_kernels
    [ 54%] Linking CUDA device code CMakeFiles/transformer_kernels.dir/cmake_device_link.o
    [ 55%] Linking CUDA static library ../../lib/libtransformer_kernels.a
    [ 55%] Built target transformer_kernels
    /workspace/FasterTransformer/fastertransformer/trt_fused_multihead_attention/qkvToContext.cu: In function 'void fastertransformer::set_alpha(uint32_t&, float, fastertransformer::Data_type)':
    /workspace/FasterTransformer/fastertransformer/trt_fused_multihead_attention/qkvToContext.cu:34:46: warning: dereferencing type-punned pointer will break strict-aliasing rules [-Wstrict-aliasing]
       34 |         alpha = reinterpret_cast<const uint32_t&>(h2);
          |                                              ^~
    /workspace/FasterTransformer/fastertransformer/trt_fused_multihead_attention/qkvToContext.cu:38:46: warning: dereferencing type-punned pointer will break strict-aliasing rules [-Wstrict-aliasing]
       38 |         alpha = reinterpret_cast<const uint32_t&>(norm);
          |                                              ^~~~
    /workspace/FasterTransformer/fastertransformer/trt_fused_multihead_attention/qkvToContext.cu: In member function 'void fastertransformer::FusedMHARunnerInt8v2::mhaImpl::setup(int, int)':
    /workspace/FasterTransformer/fastertransformer/trt_fused_multihead_attention/qkvToContext.cu:329:62: warning: dereferencing type-punned pointer will break strict-aliasing rules [-Wstrict-aliasing]
      329 |         params.scale_bmm1 = reinterpret_cast<const uint32_t&>(scaleBmm1);
          |                                                              ^~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/trt_fused_multihead_attention/qkvToContext.cu:330:62: warning: dereferencing type-punned pointer will break strict-aliasing rules [-Wstrict-aliasing]
      330 |         params.scale_bmm2 = reinterpret_cast<const uint32_t&>(scaleBmm2);
          |                                                              ^~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/trt_fused_multihead_attention/qkvToContext.cu:331:65: warning: dereferencing type-punned pointer will break strict-aliasing rules [-Wstrict-aliasing]
      331 |         params.scale_softmax = reinterpret_cast<const uint32_t&>(scaleSoftmax);
          |                                                                 ^~~~~~~~~~~~
    [ 56%] Linking CUDA device code CMakeFiles/cuda_kernels.dir/cmake_device_link.o
    /workspace/FasterTransformer/fastertransformer/trt_fused_multihead_attention/qkvToContext.cu: In member function 'virtual void fastertransformer::FusedMHARunnerInt8v2::setup(int, int)':
    /workspace/FasterTransformer/fastertransformer/trt_fused_multihead_attention/qkvToContext.cu:280:8: warning: 'warps_m' may be used uninitialized in this function [-Wmaybe-uninitialized]
      280 |         size_t warps_m, warps_n, warps_k = 1;
          |        ^~~~~~~
    /workspace/FasterTransformer/fastertransformer/trt_fused_multihead_attention/qkvToContext.cu:280:17: warning: 'warps_n' may be used uninitialized in this function [-Wmaybe-uninitialized]
      280 |         size_t warps_m, warps_n, warps_k = 1;
          |                 ^~~~~~~
    [ 57%] Linking CUDA static library ../../lib/libcuda_kernels.a
    [ 57%] Built target cuda_kernels
    Scanning dependencies of target cuda_int8_kernels
    Scanning dependencies of target decoder
    [ 58%] Building CUDA object fastertransformer/cuda/CMakeFiles/cuda_int8_kernels.dir/cuda_int8_kernels.cu.o
    [ 59%] Building CUDA object fastertransformer/cuda/CMakeFiles/decoder.dir/open_decoder.cu.o
    [ 60%] Building CUDA object fastertransformer/cuda/CMakeFiles/decoder.dir/masked_multihead_attention.cu.o
    [ 61%] Linking CUDA device code CMakeFiles/trt_fused_multi_head_attention.dir/cmake_device_link.o
    [ 62%] Linking CXX static library ../../lib/libtrt_fused_multi_head_attention.a
    [ 62%] Built target trt_fused_multi_head_attention
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu(303): warning: specified alignment (2) is different from alignment (4) specified on a previous declaration
              detected during:
                instantiation of "void fastertransformer::masked_attention_kernel(T *, T *, T *, const T *, T *, const T *, T *, const T *, T *, const __nv_bool *, int, int, int, int, T) [with T=half]"
    (445): here
                instantiation of "void fastertransformer::masked_attention_dispatch(T *, T *, T *, const T *, T *, const T *, T *, const T *, T *, const __nv_bool *, int, int, int, int, int, int, cudaStream_t) [with T=half]"
    (503): here
    
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu(1420): warning: specified alignment (2) is different from alignment (4) specified on a previous declaration
              detected during:
                instantiation of "void fastertransformer::cross_attention_kernel(T *, const T *, T *, const T *, T *, const T *, const int *, T *, const __nv_bool *, int, int, int, int, int, T) [with T=half]"
    (1555): here
                instantiation of "void fastertransformer::cross_attention_dispatch(T *, const T *, T *, const T *, T *, const T *, const int *, T *, const __nv_bool *, int, int, int, int, int, cudaStream_t) [with T=half]"
    (1576): here
    
    [ 63%] Linking CUDA device code CMakeFiles/cuda_int8_kernels.dir/cmake_device_link.o
    [ 64%] Linking CUDA static library ../../lib/libcuda_int8_kernels.a
    [ 64%] Built target cuda_int8_kernels
    Scanning dependencies of target encoder
    [ 65%] Building CUDA object fastertransformer/cuda/CMakeFiles/encoder.dir/open_attention.cu.o
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu: In instantiation of 'void fastertransformer::masked_attention_dispatch(T*, T*, T*, const T*, T*, const T*, T*, const T*, T*, const bool*, int, int, int, int, int, int, cudaStream_t) [with T = float; cudaStream_t = CUstream_st*]':
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:484:398:   required from here
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:454:7: warning: 'void* memset(void*, int, size_t)' clearing an object of non-trivial type 'struct Masked_multihead_attention_params<float>'; use assignment or value-initialization instead [-Wclass-memaccess]
      454 |     memset(&params, 0, sizeof(params));
          |     ~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/masked_multihead_attention.h:48:8: note: 'struct Masked_multihead_attention_params<float>' declared here
       48 | struct Masked_multihead_attention_params {
          |        ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu: In instantiation of 'void fastertransformer::masked_attention_dispatch(T*, T*, T*, const T*, T*, const T*, T*, const T*, T*, const bool*, int, int, int, int, int, int, cudaStream_t) [with T = __half; cudaStream_t = CUstream_st*]':
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:503:389:   required from here
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:454:7: warning: 'void* memset(void*, int, size_t)' clearing an object of non-trivial type 'struct Masked_multihead_attention_params<short unsigned int>'; use assignment or value-initialization instead [-Wclass-memaccess]
      454 |     memset(&params, 0, sizeof(params));
          |     ~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/masked_multihead_attention.h:48:8: note: 'struct Masked_multihead_attention_params<short unsigned int>' declared here
       48 | struct Masked_multihead_attention_params {
          |        ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu: In instantiation of 'void fastertransformer::fusedQKV_masked_attention_dispatch(const T*, const T*, T*, T*, T*, const bool*, int, int, int, int, int, int, cudaStream_t) [with T = float; cudaStream_t = CUstream_st*]':
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:781:318:   required from here
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:750:7: warning: 'void* memset(void*, int, size_t)' clearing an object of non-trivial type 'struct Masked_multihead_attention_params<float>'; use assignment or value-initialization instead [-Wclass-memaccess]
      750 |     memset(&params, 0, sizeof(params));
          |     ~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/masked_multihead_attention.h:48:8: note: 'struct Masked_multihead_attention_params<float>' declared here
       48 | struct Masked_multihead_attention_params {
          |        ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu: In instantiation of 'void fastertransformer::fusedQKV_masked_attention_dispatch(const T*, const T*, T*, T*, T*, const bool*, int, int, int, int, int, int, cudaStream_t) [with T = __half; cudaStream_t = CUstream_st*]':
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:796:313:   required from here
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:750:7: warning: 'void* memset(void*, int, size_t)' clearing an object of non-trivial type 'struct Masked_multihead_attention_params<short unsigned int>'; use assignment or value-initialization instead [-Wclass-memaccess]
      750 |     memset(&params, 0, sizeof(params));
          |     ~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/masked_multihead_attention.h:48:8: note: 'struct Masked_multihead_attention_params<short unsigned int>' declared here
       48 | struct Masked_multihead_attention_params {
          |        ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu: In instantiation of 'void fastertransformer::fusedQKV_masked_attention_dispatch_v2(const T*, const T*, T*, T*, T*, const bool*, int, int, int, int, int, int, int, const int*, cudaStream_t) [with T = float; cudaStream_t = CUstream_st*]':
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:883:373:   required from here
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:851:7: warning: 'void* memset(void*, int, size_t)' clearing an object of non-trivial type 'struct Masked_multihead_attention_params<float>'; use assignment or value-initialization instead [-Wclass-memaccess]
      851 |   memset(&params, 0, sizeof(params));
          |   ~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/masked_multihead_attention.h:48:8: note: 'struct Masked_multihead_attention_params<float>' declared here
       48 | struct Masked_multihead_attention_params {
          |        ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu: In instantiation of 'void fastertransformer::fusedQKV_masked_attention_dispatch_v2(const T*, const T*, T*, T*, T*, const bool*, int, int, int, int, int, int, int, const int*, cudaStream_t) [with T = __half; cudaStream_t = CUstream_st*]':
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:900:368:   required from here
    /workspace/FasterTransformer/fastertransformer/cuda/open_decoder.cu:851:7: warning: 'void* memset(void*, int, size_t)' clearing an object of non-trivial type 'struct Masked_multihead_attention_params<short unsigned int>'; use assignment or value-initialization instead [-Wclass-memaccess]
      851 |   memset(&params, 0, sizeof(params));
          |   ~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~
    /workspace/FasterTransformer/fastertransformer/cuda/masked_multihead_attention.h:48:8: note: 'struct Masked_multihead_attention_params<short unsigned int>' declared here
       48 | struct Masked_multihead_attention_params {
          |        ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
    [ 66%] Linking CUDA device code CMakeFiles/decoder.dir/cmake_device_link.o
    [ 66%] Linking CXX static library ../../lib/libdecoder.a
    [ 66%] Built target decoder
    [ 67%] Linking CUDA device code CMakeFiles/encoder.dir/cmake_device_link.o
    [ 68%] Linking CXX static library ../../lib/libencoder.a
    [ 68%] Built target encoder
    Scanning dependencies of target encoder_sample
    [ 69%] Building CXX object sample/cpp/CMakeFiles/encoder_sample.dir/encoder_sample.cc.o
    [ 70%] Linking CXX executable ../../bin/encoder_sample
    [ 70%] Built target encoder_sample
    [ 71%] Linking CUDA device code CMakeFiles/online_softmax_beamsearch.dir/cmake_device_link.o
    [ 72%] Linking CUDA static library ../../lib/libonline_softmax_beamsearch.a
    [ 72%] Built target online_softmax_beamsearch
    [ 73%] Linking CUDA device code CMakeFiles/topk.dir/cmake_device_link.o
    [ 74%] Linking CUDA static library ../../lib/libtopk.a
    [ 74%] Built target topk
    Scanning dependencies of target decoding
    [ 75%] Building CUDA object fastertransformer/cuda/CMakeFiles/decoding.dir/decoding_kernels.cu.o
    [ 76%] Linking CUDA device code CMakeFiles/decoding.dir/cmake_device_link.o
    [ 77%] Linking CUDA static library ../../lib/libdecoding.a
    [ 77%] Built target decoding
    Scanning dependencies of target gpt_triton_backend
    Scanning dependencies of target decoding_sampling_sample
    Scanning dependencies of target transformer-static
    Scanning dependencies of target pyt_fastertransformer
    Scanning dependencies of target gpt_sample
    Scanning dependencies of target decoding_beamsearch_sample
    [ 78%] Linking CUDA device code CMakeFiles/transformer-static.dir/cmake_device_link.o
    [ 79%] Building CXX object sample/cpp/CMakeFiles/decoding_sampling_sample.dir/decoding_sampling_sample.cc.o
    [ 80%] Building CXX object fastertransformer/triton_backend/CMakeFiles/gpt_triton_backend.dir/gpt_triton_backend.cc.o
    [ 81%] Building CXX object sample/cpp/CMakeFiles/decoding_beamsearch_sample.dir/decoding_beamsearch_sample.cc.o
    [ 82%] Building CXX object sample/cpp/CMakeFiles/gpt_sample.dir/gpt_sample.cc.o
    [ 83%] Linking CXX static library lib/libtransformer-static.a
    [ 84%] Building CXX object fastertransformer/th_op/CMakeFiles/pyt_fastertransformer.dir/ft_op.cc.o
    [ 84%] Building CXX object fastertransformer/th_op/CMakeFiles/pyt_fastertransformer.dir/encoder.cc.o
    [ 85%] Building CXX object fastertransformer/th_op/CMakeFiles/pyt_fastertransformer.dir/decoder.cc.o
    [ 86%] Building CXX object fastertransformer/th_op/CMakeFiles/pyt_fastertransformer.dir/decoding.cc.o
    [ 87%] Building CXX object fastertransformer/th_op/CMakeFiles/pyt_fastertransformer.dir/gpt.cc.o
    [ 88%] Building CUDA object fastertransformer/th_op/CMakeFiles/pyt_fastertransformer.dir/utils.cu.o
    [ 89%] Building CXX object fastertransformer/th_op/CMakeFiles/pyt_fastertransformer.dir/weight_quantize_op.cc.o
    [ 89%] Built target transformer-static
    Scanning dependencies of target gpt
    [ 90%] Building CXX object CMakeFiles/gpt.dir/sample/cpp/gpt_sample.cc.o
    [ 91%] Linking CXX executable ../../bin/decoding_sampling_sample
    [ 92%] Linking CXX executable ../../bin/decoding_beamsearch_sample
    [ 92%] Built target decoding_sampling_sample
    [ 92%] Built target decoding_beamsearch_sample
    [ 93%] Linking CXX executable ../../bin/gpt_sample
    [ 93%] Built target gpt_sample
    [ 94%] Linking CXX executable bin/gpt
    [ 94%] Built target gpt
    [ 94%] Linking CXX shared library ../../lib/libgpt_triton_backend.so
    [ 94%] Built target gpt_triton_backend
    Scanning dependencies of target transformer-shared
    Scanning dependencies of target gpt_thread_sample
    Scanning dependencies of target gpt_triton_sample
    [ 95%] Linking CUDA device code CMakeFiles/transformer-shared.dir/cmake_device_link.o
    [ 97%] Building CXX object sample/cpp/CMakeFiles/gpt_triton_sample.dir/gpt_triton_sample.cc.o
    [ 97%] Building CXX object sample/cpp/CMakeFiles/gpt_thread_sample.dir/gpt_thread_sample.cc.o
    [ 97%] Linking CXX shared library lib/libtransformer-shared.so
    [ 97%] Built target transformer-shared
    [ 98%] Linking CXX executable ../../bin/gpt_triton_sample
    [ 98%] Built target gpt_triton_sample
    [ 99%] Linking CXX executable ../../bin/gpt_thread_sample
    [ 99%] Built target gpt_thread_sample
    [ 99%] Linking CUDA device code CMakeFiles/pyt_fastertransformer.dir/cmake_device_link.o
    [100%] Linking CXX shared library ../../lib/libpyt_fastertransformer.so
    [100%] Built target pyt_fastertransformer
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     python -c 'import torch; torch.ops.load_library("./lib/libpyt_fastertransformer.so")'                                                              130 ↵
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     ././bin/encoder_gemm 32 32 12 64 0 0                                                                                                               130 ↵
    ./bin/encoder_sample 32 12 32 12 64 0 0 0 0
    
    Device Tesla V100-SXM2-16GB
    ***Encoder Gemm Testing Begin***
    ***Cublas Gemm Testing Begin***
    
    -----------------------------
    GEMM test 0: [M: 1024, K: 768, N: 768] from_tensor * weightQ/K/V, attr * output_kernel
    algo_-1 costs 0.123ms
    algo_2 costs 0.129ms
    algo_3 costs 0.165ms
    algo_4 costs 0.124ms
    algo_5 costs 0.160ms
    algo_6 costs 0.159ms
    algo_7 costs 0.115ms
    algo_8 costs 0.132ms
    algo_9 costs 0.111ms
    algo_10 costs 0.128ms
    algo_11 costs 0.154ms
    algo_18 costs 0.119ms
    algo_19 costs 0.101ms
    algo_20 costs 0.117ms
    algo_21 costs 0.130ms
    algo_22 costs 0.112ms
    algo_23 costs 0.115ms
    fast_algo 19 costs 0.101 ms
    
    -----------------------------
    GEMM test 1: [M: 1024, K: 768, N: 3072] attr_output * inter_kernel
    algo_-1 costs 0.384ms
    algo_2 costs 0.407ms
    algo_3 costs 0.410ms
    algo_4 costs 0.365ms
    algo_5 costs 0.343ms
    algo_6 costs 0.406ms
    algo_7 costs 0.366ms
    algo_8 costs 0.367ms
    algo_9 costs 0.348ms
    algo_10 costs 0.351ms
    algo_11 costs 0.349ms
    algo_18 costs 0.439ms
    algo_19 costs 0.383ms
    algo_20 costs 0.369ms
    algo_21 costs 0.444ms
    algo_22 costs 0.391ms
    algo_23 costs 0.431ms
    fast_algo 5 costs 0.343 ms
    
    -----------------------------
    GEMM test 2: [M: 1024, K: 3072, N: 768] inter_matmul * output_kernel
    algo_-1 costs 0.366ms
    algo_1 costs 0.631ms
    algo_2 costs 0.414ms
    algo_3 costs 0.543ms
    algo_4 costs 0.402ms
    algo_5 costs 0.525ms
    algo_6 costs 0.525ms
    algo_7 costs 0.361ms
    algo_8 costs 0.377ms
    algo_9 costs 0.347ms
    algo_10 costs 0.370ms
    algo_11 costs 0.375ms
    algo_18 costs 0.444ms
    algo_19 costs 0.360ms
    algo_20 costs 0.404ms
    algo_21 costs 0.441ms
    algo_22 costs 0.383ms
    algo_23 costs 0.407ms
    fast_algo 9 costs 0.347 ms
    
    -----------------------------
    GEMM test 3: [M: 32, K: 64, N: 32] attention batched Gemm1
    algo_-1 costs 0.029ms
    algo_0 costs 0.022ms
    algo_1 costs 0.026ms
    algo_2 costs 0.022ms
    algo_3 costs 0.036ms
    algo_4 costs 0.066ms
    algo_5 costs 0.022ms
    algo_6 costs 0.026ms
    algo_7 costs 0.022ms
    algo_8 costs 0.036ms
    algo_9 costs 0.066ms
    algo_18 costs 0.019ms
    algo_19 costs 0.024ms
    algo_20 costs 0.039ms
    algo_21 costs 0.019ms
    algo_22 costs 0.024ms
    algo_23 costs 0.039ms
    fast_algo 21 costs 0.019 ms
    
    -----------------------------
    GEMM test 4: [M: 32, K: 32, N: 64] attention batched Gemm2
    algo_-1 costs 0.017ms
    algo_0 costs 0.024ms
    algo_1 costs 0.017ms
    algo_2 costs 0.016ms
    algo_3 costs 0.023ms
    algo_4 costs 0.040ms
    algo_5 costs 0.024ms
    algo_6 costs 0.017ms
    algo_7 costs 0.016ms
    algo_8 costs 0.023ms
    algo_9 costs 0.040ms
    algo_18 costs 0.022ms
    algo_19 costs 0.020ms
    algo_20 costs 0.031ms
    algo_21 costs 0.022ms
    algo_22 costs 0.020ms
    algo_23 costs 0.031ms
    fast_algo 7 costs 0.016 ms
    
    -----------------------------
    GEMM test 5: [M: 1024, K: 768, N: 768] from_tensor * weight_QKV in BatchGemm
    algo_-1 costs 0.273ms
    algo_0 costs 0.286ms
    algo_1 costs 0.272ms
    algo_2 costs 0.275ms
    algo_3 costs 0.267ms
    algo_4 costs 0.270ms
    algo_5 costs 0.286ms
    algo_6 costs 0.273ms
    algo_7 costs 0.274ms
    algo_8 costs 0.267ms
    algo_9 costs 0.270ms
    algo_18 costs 0.309ms
    algo_19 costs 0.277ms
    algo_20 costs 0.294ms
    algo_21 costs 0.309ms
    algo_22 costs 0.277ms
    algo_23 costs 0.294ms
    fast_algo 8 costs 0.267 ms
    ***cublas Gemm Testing End***
    
    ***Encoder Gemm Testing End***
    Device Tesla V100-SXM2-16GB
    Device Tesla V100-SXM2-16GB
    before allocate free 15.44 GB total 15.78 GB
    After allocate free 15.40 GB used 0.38 GB total 15.78 GB
    [INFO] batch_size 32 seq_len 32 layer 12 FT-CPP-time 15.40 ms ( 50 iterations)
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ ./bin/encoder_gemm 32 32 12 64 0 0
    python pytorch/encoder_sample.py 32 12 32 12 64 --time
    
    Device Tesla V100-SXM2-16GB
    ***Encoder Gemm Testing Begin***
    ***Cublas Gemm Testing Begin***
    
    -----------------------------
    GEMM test 0: [M: 1024, K: 768, N: 768] from_tensor * weightQ/K/V, attr * output_kernel
    algo_-1 costs 0.109ms
    algo_2 costs 0.114ms
    algo_3 costs 0.145ms
    algo_4 costs 0.109ms
    algo_5 costs 0.141ms
    algo_6 costs 0.141ms
    algo_7 costs 0.103ms
    algo_8 costs 0.116ms
    algo_9 costs 0.097ms
    algo_10 costs 0.111ms
    algo_11 costs 0.143ms
    algo_18 costs 0.113ms
    algo_19 costs 0.096ms
    algo_20 costs 0.111ms
    algo_21 costs 0.124ms
    algo_22 costs 0.107ms
    algo_23 costs 0.108ms
    fast_algo 19 costs 0.096 ms
    
    -----------------------------
    GEMM test 1: [M: 1024, K: 768, N: 3072] attr_output * inter_kernel
    algo_-1 costs 0.361ms
    algo_2 costs 0.384ms
    algo_3 costs 0.404ms
    algo_4 costs 0.366ms
    algo_5 costs 0.344ms
    algo_6 costs 0.407ms
    algo_7 costs 0.366ms
    algo_8 costs 0.367ms
    algo_9 costs 0.349ms
    algo_10 costs 0.351ms
    algo_11 costs 0.347ms
    algo_18 costs 0.439ms
    algo_19 costs 0.382ms
    algo_20 costs 0.367ms
    algo_21 costs 0.442ms
    algo_22 costs 0.391ms
    algo_23 costs 0.431ms
    fast_algo 5 costs 0.344 ms
    
    -----------------------------
    GEMM test 2: [M: 1024, K: 3072, N: 768] inter_matmul * output_kernel
    algo_-1 costs 0.365ms
    algo_1 costs 0.629ms
    algo_2 costs 0.414ms
    algo_3 costs 0.543ms
    algo_4 costs 0.402ms
    algo_5 costs 0.526ms
    algo_6 costs 0.530ms
    algo_7 costs 0.363ms
    algo_8 costs 0.377ms
    algo_9 costs 0.348ms
    algo_10 costs 0.370ms
    algo_11 costs 0.375ms
    algo_18 costs 0.445ms
    algo_19 costs 0.357ms
    algo_20 costs 0.404ms
    algo_21 costs 0.442ms
    algo_22 costs 0.384ms
    algo_23 costs 0.406ms
    fast_algo 9 costs 0.348 ms
    
    -----------------------------
    GEMM test 3: [M: 32, K: 64, N: 32] attention batched Gemm1
    algo_-1 costs 0.029ms
    algo_0 costs 0.022ms
    algo_1 costs 0.026ms
    algo_2 costs 0.022ms
    algo_3 costs 0.036ms
    algo_4 costs 0.066ms
    algo_5 costs 0.022ms
    algo_6 costs 0.026ms
    algo_7 costs 0.022ms
    algo_8 costs 0.036ms
    algo_9 costs 0.066ms
    algo_18 costs 0.019ms
    algo_19 costs 0.024ms
    algo_20 costs 0.039ms
    algo_21 costs 0.019ms
    algo_22 costs 0.024ms
    algo_23 costs 0.039ms
    fast_algo 21 costs 0.019 ms
    
    -----------------------------
    GEMM test 4: [M: 32, K: 32, N: 64] attention batched Gemm2
    algo_-1 costs 0.017ms
    algo_0 costs 0.024ms
    algo_1 costs 0.017ms
    algo_2 costs 0.016ms
    algo_3 costs 0.023ms
    algo_4 costs 0.040ms
    algo_5 costs 0.024ms
    algo_6 costs 0.017ms
    algo_7 costs 0.016ms
    algo_8 costs 0.023ms
    algo_9 costs 0.040ms
    algo_18 costs 0.022ms
    algo_19 costs 0.020ms
    algo_20 costs 0.031ms
    algo_21 costs 0.022ms
    algo_22 costs 0.020ms
    algo_23 costs 0.031ms
    fast_algo 7 costs 0.016 ms
    
    -----------------------------
    GEMM test 5: [M: 1024, K: 768, N: 768] from_tensor * weight_QKV in BatchGemm
    algo_-1 costs 0.273ms
    algo_0 costs 0.286ms
    algo_1 costs 0.273ms
    algo_2 costs 0.274ms
    algo_3 costs 0.267ms
    algo_4 costs 0.270ms
    algo_5 costs 0.286ms
    algo_6 costs 0.273ms
    algo_7 costs 0.276ms
    algo_8 costs 0.267ms
    algo_9 costs 0.270ms
    algo_18 costs 0.309ms
    algo_19 costs 0.277ms
    algo_20 costs 0.294ms
    algo_21 costs 0.309ms
    algo_22 costs 0.277ms
    algo_23 costs 0.294ms
    fast_algo 8 costs 0.267 ms
    ***cublas Gemm Testing End***
    
    ***Encoder Gemm Testing End***
    2021-09-07 15:58:23.469196: I tensorflow/stream_executor/platform/default/dso_loader.cc:49] Successfully opened dynamic library libcudart.so.11.0
    WARNING:tensorflow:Deprecation warnings have been disabled. Set TF_ENABLE_DEPRECATION_WARNINGS=1 to re-enable them.
    
    =============== Argument ===============
    batch_size: 32
    layer_num: 12
    seq_len: 32
    head_num: 12
    head_size: 64
    hidden_dim: 768
    weight_path: None
    use_fp16: False
    int8_mode: 0
    avg_seq_len: -1
    test_time: True
    remove_padding: False
    allow_gemm_test: False
    ========================================
    
    tensor([[[-1.5207,  1.6938, -0.9449,  ...,  1.0584, -0.0928,  1.0019],
             [-1.5207,  1.6938, -0.9449,  ...,  1.0584, -0.0928,  1.0019],
             [-1.5207,  1.6938, -0.9449,  ...,  1.0584, -0.0928,  1.0019],
             ...,
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]],
    
            [[-1.5207,  1.6939, -0.9454,  ...,  1.0586, -0.0935,  1.0023],
             [-1.5207,  1.6939, -0.9454,  ...,  1.0586, -0.0935,  1.0023],
             [-1.5207,  1.6939, -0.9454,  ...,  1.0586, -0.0935,  1.0023],
             ...,
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]],
    
            [[-1.5212,  1.6943, -0.9453,  ...,  1.0586, -0.0929,  1.0017],
             [-1.5212,  1.6943, -0.9453,  ...,  1.0586, -0.0929,  1.0017],
             [-1.5212,  1.6943, -0.9453,  ...,  1.0586, -0.0929,  1.0017],
             ...,
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]],
    
            ...,
    
            [[-1.5202,  1.6939, -0.9454,  ...,  1.0585, -0.0928,  1.0020],
             [-1.5202,  1.6939, -0.9454,  ...,  1.0585, -0.0928,  1.0020],
             [-1.5202,  1.6939, -0.9454,  ...,  1.0585, -0.0928,  1.0020],
             ...,
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]],
    
            [[-1.5209,  1.6946, -0.9452,  ...,  1.0585, -0.0928,  1.0021],
             [-1.5209,  1.6946, -0.9452,  ...,  1.0585, -0.0928,  1.0021],
             [-1.5209,  1.6946, -0.9452,  ...,  1.0585, -0.0928,  1.0021],
             ...,
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]],
    
            [[-1.5211,  1.6941, -0.9452,  ...,  1.0587, -0.0927,  1.0020],
             [-1.5211,  1.6941, -0.9452,  ...,  1.0587, -0.0927,  1.0020],
             [-1.5211,  1.6941, -0.9452,  ...,  1.0587, -0.0927,  1.0020],
             ...,
             [-1.5211,  1.6941, -0.9452,  ...,  1.0587, -0.0927,  1.0020],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000],
             [-0.0000,  0.0000, -0.0000,  ...,  0.0000, -0.0000,  0.0000]]],
           device='cuda:0')
    torch.Size([32, 32, 768])
    Traceback (most recent call last):
      File "pytorch/encoder_sample.py", line 257, in <module>
        main()
      File "pytorch/encoder_sample.py", line 171, in main
        ft_output = custom_encoder(inp, mask, mem_seq_lens)[0] * output_mask
      File "/home/amralaa/.local/lib/python3.8/site-packages/torch/nn/modules/module.py", line 1051, in _call_impl
        return forward_call(*input, **kwargs)
    RuntimeError: The following operation failed in the TorchScript interpreter.
    Traceback of TorchScript (most recent call last):
      File "/workspace/FasterTransformer/build/pytorch/utils/encoder.py", line 200, in forward
                trt_seq_len = torch.cat([trt_seq_len, torch.tensor([batch * max_seq_len], device='cuda').to(trt_seq_len.dtype)], dim=0).to(torch.int32)
            for i in range(self.layer_num):
                hidden_states = self.encoders[i].forward(hidden_states, attention_mask, trt_seq_len, sequence_id_offset)
                                ~~~~~~~~~~~~~~~~~~~~~~~~ <--- HERE
            if self.remove_padding:
                hidden_states = self.rebuild_padding(hidden_states, sequence_id_offset, attention_mask, 0)
    RuntimeError: [FT][ERROR] CUDA runtime error: CUBLAS_STATUS_INVALID_VALUE /workspace/FasterTransformer/fastertransformer/cuda/open_attention.h:939
    
    
    
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ nvidia-smi
    Tue Sep  7 16:34:08 2021
    +-----------------------------------------------------------------------------+
    | NVIDIA-SMI 450.80.02    Driver Version: 450.80.02    CUDA Version: 11.1     |
    |-------------------------------+----------------------+----------------------+
    | GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
    | Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
    |                               |                      |               MIG M. |
    |===============================+======================+======================|
    |   0  Tesla V100-SXM2...  On   | 00001D0E:00:00.0 Off |                    0 |
    | N/A   32C    P0    43W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   1  Tesla V100-SXM2...  On   | 00003D8F:00:00.0 Off |                    0 |
    | N/A   33C    P0    43W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   2  Tesla V100-SXM2...  On   | 00006B96:00:00.0 Off |                    0 |
    | N/A   32C    P0    43W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   3  Tesla V100-SXM2...  On   | 000084FE:00:00.0 Off |                    0 |
    | N/A   33C    P0    43W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   4  Tesla V100-SXM2...  On   | 0000C3F5:00:00.0 Off |                    0 |
    | N/A   34C    P0    43W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   5  Tesla V100-SXM2...  On   | 0000DCCF:00:00.0 Off |                    0 |
    | N/A   33C    P0    44W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   6  Tesla V100-SXM2...  On   | 0000E931:00:00.0 Off |                    0 |
    | N/A   32C    P0    44W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    |   7  Tesla V100-SXM2...  On   | 0000F01A:00:00.0 Off |                    0 |
    | N/A   33C    P0    42W / 300W |      0MiB / 16160MiB |      0%      Default |
    |                               |                      |                  N/A |
    +-------------------------------+----------------------+----------------------+
    
    +-----------------------------------------------------------------------------+
    | Processes:                                                                  |
    |  GPU   GI   CI        PID   Type   Process name                  GPU Memory |
    |        ID   ID                                                   Usage      |
    |=============================================================================|
    |  No running processes found                                                 |
    +-----------------------------------------------------------------------------+
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ gcc --version
    gcc (Ubuntu 9.3.0-17ubuntu1~20.04) 9.3.0
    Copyright (C) 2019 Free Software Foundation, Inc.
    This is free software; see the source for copying conditions.  There is NO
    warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
    
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ g++ --version
    g++ (Ubuntu 9.3.0-17ubuntu1~20.04) 9.3.0
    Copyright (C) 2019 Free Software Foundation, Inc.
    This is free software; see the source for copying conditions.  There is NO
    warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
    
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ nvcc --version
    nvcc: NVIDIA (R) Cuda compiler driver
    Copyright (c) 2005-2020 NVIDIA Corporation
    Built on Mon_Oct_12_20:09:46_PDT_2020
    Cuda compilation tools, release 11.1, V11.1.105
    Build cuda_11.1.TC455_06.29190527_0
    [email protected]:/workspace/FasterTransformer/build (main) (DISPLAY=)
     $ dpkg -l | grep -P 'nvidia|cuda'
    ii  cuda-compat-11-1                       455.32.00-1                       amd64        CUDA Compatibility Platform
    ii  cuda-cudart-11-1                       11.1.74-1                         amd64        CUDA Runtime native Libraries
    ii  cuda-cudart-dev-11-1                   11.1.74-1                         amd64        CUDA Runtime native dev links, headers
    ii  cuda-cuobjdump-11-1                    11.1.74-1                         amd64        CUDA cuobjdump
    ii  cuda-cupti-11-1                        11.1.105-1                        amd64        CUDA profiling tools runtime libs.
    ii  cuda-cupti-dev-11-1                    11.1.105-1                        amd64        CUDA profiling tools interface.
    ii  cuda-driver-dev-11-1                   11.1.74-1                         amd64        CUDA Driver native dev stub library
    ii  cuda-gdb-11-1                          11.1.105-1                        amd64        CUDA-GDB
    ii  cuda-memcheck-11-1                     11.1.105-1                        amd64        CUDA-MEMCHECK
    ii  cuda-nvcc-11-1                         11.1.105-1                        amd64        CUDA nvcc
    ii  cuda-nvdisasm-11-1                     11.1.74-1                         amd64        CUDA disassembler
    ii  cuda-nvml-dev-11-1                     11.1.74-1                         amd64        NVML native dev links, headers
    ii  cuda-nvprof-11-1                       11.1.105-1                        amd64        CUDA Profiler tools
    ii  cuda-nvprune-11-1                      11.1.74-1                         amd64        CUDA nvprune
    ii  cuda-nvrtc-11-1                        11.1.105-1                        amd64        NVRTC native runtime libraries
    ii  cuda-nvrtc-dev-11-1                    11.1.105-1                        amd64        NVRTC native dev links, headers
    ii  cuda-nvtx-11-1                         11.1.74-1                         amd64        NVIDIA Tools Extension
    ii  cuda-sanitizer-11-1                    11.1.105-1                        amd64        CUDA Sanitizer
    ii  libcudnn8                              8.0.5.43-1+cuda11.1               amd64        cuDNN runtime libraries
    ii  libcudnn8-dev                          8.0.5.43-1+cuda11.1               amd64        cuDNN development libraries and headers
    ii  libnccl-dev                            2.8.3-1+cuda11.1                  amd64        NVIDIA Collective Communication Library (NCCL) Development Files
    ii  libnccl2                               2.8.3-1+cuda11.1                  amd64        NVIDIA Collective Communication Library (NCCL) Runtime
    ii  libnvinfer-bin                         7.2.2-1+cuda11.1                  amd64        TensorRT binaries
    ii  libnvinfer-dev                         7.2.2-1+cuda11.1                  amd64        TensorRT development libraries and headers
    ii  libnvinfer-plugin-dev                  7.2.2-1+cuda11.1                  amd64        TensorRT plugin libraries and headers
    ii  libnvinfer-plugin7                     7.2.2-1+cuda11.1                  amd64        TensorRT plugin library
    ii  libnvinfer7                            7.2.2-1+cuda11.1                  amd64        TensorRT runtime libraries
    ii  libnvonnxparsers-dev                   7.2.2-1+cuda11.1                  amd64        TensorRT ONNX libraries
    ii  libnvonnxparsers7                      7.2.2-1+cuda11.1                  amd64        TensorRT ONNX libraries
    ii  libnvparsers-dev                       7.2.2-1+cuda11.1                  amd64        TensorRT parsers libraries
    ii  libnvparsers7                          7.2.2-1+cuda11.1                  amd64        TensorRT parsers libraries
    
    opened by amralaa-MSFT 18
  • Drop in Performance with TP>1, batch size and long generations

    Drop in Performance with TP>1, batch size and long generations

    Noticed a small drop in performance (<1%) for the same model when we use tensor sharding. For example, A 355 million parameter model with tp=1, gets an accuracy of 47.30 on FasterTransformer (expected performance for my model) The same model with tp=4, gets an accuracy of 46.9 on FasterTransformer. Has this been noticed before.

    I am using the pytorch docker nvcr.io/nvidia/pytorch:20.12-py3 contains the PyTorch 1.8.0 and python 3.8

    opened by bharatv007 17
  • [FasterTransformer/V2] No speedUp when the sequence length is large

    [FasterTransformer/V2] No speedUp when the sequence length is large

    Hello, I had implemented a gpt2 model according to FT. Then I tested the performance between pytorch(Fairseq) and FT, This is the result :

    Setting : batch = 1, hidden_units = 1024, head_num = 16, size_per_head = 64 time : ms seq_len : 8 16 32 64 128 256 512 800 pytorch: 21 23 23 23 28 23.6 22.6 24 ___FT : 6.2 6.4 6.7 7.6 8.6 12.3 24 34.7

    From this table , it seems like the FT is much worse than pytorch when the seq_len is larger than about 500. it's very unacceptable that the FT is slower. And why the performance of pytorch is almost never change with the increase of sequence length ? It's also has the same phenomenon when I test the masked BERT model. Anyone help ?

    opened by ligonzheng 17
  • [FastTransformer v3.1/TensorFlow] Get CUBLAS_STATUS_INTERNAL_ERROR when run tensorflow/gpt2-sample.py

    [FastTransformer v3.1/TensorFlow] Get CUBLAS_STATUS_INTERNAL_ERROR when run tensorflow/gpt2-sample.py

    Related to FastTransformer v3.1/TensorFlow/GPT-2

    Describe the bug If I run ./bin/decoding_gemm 4 1 12 64 50257 32 768 0 before python tensorflow/gpt2_sample.py, I got Internal: [FT][ERROR] CUDA runtime error: CUBLAS_STATUS_INTERNAL_ERROR FasterTransformer/fastertransformer/cuda/open_decoder.cu:1708. However, If I don't run ./bin/decoding_gemm 4 1 12 64 50257 32 768 0 before python tensorflow/gpt2_sample.py (use the default gemm), everything is OK.

    To Reproduce Steps to reproduce the behavior:

    1. nvidia-docker run -it -v local_dir:container_dir nvcr.io/nvidia/tensorflow:19.06-py3 bash
    2. cmake -DSM=70 -DCMAKE_BUILD_TYPE=Release -DBUILD_TF=ON -DTF_PATH=/usr/local/lib/python3.5/dist-packages/tensorflow ..
    3. make
    4. ./bin/decoding_gemm 4 1 12 64 50257 32 768 0
    5. python tensorflow/gpt2_sample.py

    Expected behavior There should be no error.

    Environment Please provide at least:

    • Container version: nvcr.io/nvidia/tensorflow:19.06-py3
    • GPUs in the system: 8x Tesla V100-32GB
    • CUDA driver version: 435.21
    opened by aoyulong 17
  • [FastTransformer v3.1/TensorFlow] Get CUBLAS_STATUS_INTERNAL_ERROR when run tensorflow/gpt2-sample.py

    [FastTransformer v3.1/TensorFlow] Get CUBLAS_STATUS_INTERNAL_ERROR when run tensorflow/gpt2-sample.py

    Related to FastTransformer v3.1/TensorFlow/GPT-2

    Describe the bug If I run ./bin/decoding_gemm 4 1 12 64 50257 32 768 0 before python tensorflow/gpt2_sample.py, I got Internal: [FT][ERROR] CUDA runtime error: CUBLAS_STATUS_INTERNAL_ERROR FasterTransformer/fastertransformer/cuda/open_decoder.cu:1708. However, If I don't run ./bin/decoding_gemm 4 1 12 64 50257 32 768 0 before python tensorflow/gpt2_sample.py (use the default gemm), everything is OK.

    To Reproduce Steps to reproduce the behavior:

    1. nvidia-docker run -it -v local_dir:container_dir nvcr.io/nvidia/tensorflow:19.06-py3 bash
    2. cmake -DSM=70 -DCMAKE_BUILD_TYPE=Release -DBUILD_TF=ON -DTF_PATH=/usr/local/lib/python3.5/dist-packages/tensorflow ..
    3. make
    4. ./bin/decoding_gemm 4 1 12 64 50257 32 768 0
    5. python tensorflow/gpt2_sample.py

    Expected behavior There should be no error.

    Environment Please provide at least:

    • Container version: nvcr.io/nvidia/tensorflow:19.06-py3
    • GPUs in the system: 8x Tesla V100-32GB
    • CUDA driver version: 435.21
    opened by aoyulong 17
  • [FT][ERROR] CUDA runtime error: CUBLAS_STATUS_EXECUTION_FAILED

    [FT][ERROR] CUDA runtime error: CUBLAS_STATUS_EXECUTION_FAILED

    Branch/Tag/Commit

    nccl_dependent_refine

    Docker Image Version

    .

    GPU name

    V100

    CUDA Driver

    CUDA 10.1

    Reproduced Steps

    1. Background: Add XGLM model to FT(similar to GPT), and GPT model did not raise the error.
    2. The main error message For fp16, I report an error in this line of code: https://github.com/NVIDIA/FasterTransformer/blob/nccl_dependent_refine/fastertransformer/open_decoder.h#L735; For fp32, there is no problem.
    3. Search the same issue The error is same as https://github.com/NVIDIA/FasterTransformer/issues/181, but the weights of model are pretrained, I can not modify it.

    Thank you very much for helping me.

    bug 
    opened by gongel 1
  • T5-Decoding beam-search error in FP16 (misaligned address /root/FasterTransformer/src/fastertransformer/utils/memory_utils.cu:101)

    T5-Decoding beam-search error in FP16 (misaligned address /root/FasterTransformer/src/fastertransformer/utils/memory_utils.cu:101)

    Branch/Tag/Commit

    main

    Docker Image Version

    nvcr.io/nvidia/pytorch:21.10-py3

    GPU name

    A100 (80GB)

    CUDA Driver

    450.80.02

    Reproduced Steps

    I write a test code for reproduction in my Repo. This code converts google/t5-base model to FasterTransformer, and runs beam-search for dummy texts.

    This is the following step.

    1. apt-get install git-lfs
    2. git clone https://github.com/cpm0722/FasterTransformer -b t5-beam-search-bug
    3. mkdir -p FasterTransformer/build && cd FasterTransformer/build
    4. git submodule init && git submodule update
    5. cmake -DSM=80 -DCMAKE_BUILD_TYPE=Release -DBUILD_PYT=ON -DBUILD_MULTI_GPU=ON ..
    6. make -j12
    7. cd ../examples/pytorch/t5
    8. pip install -r requirements.txt
    9. bash run.sh

    It raises the below error.

    [INFO] WARNING: Exception occurred in dist.init_process_group(backend = 'mpi'). Maybe the process group has been initialized somewhere else.
    [INFO] WARNING: Exception occurred in dist.init_process_group(backend = 'mpi'). Maybe the process group has been initialized somewhere else.
    [FT][WARNING] Skip NCCL initialization since requested tensor/pipeline parallel sizes are equals to 1.
    [WARNING] gemm_config.in is not found; using default GEMM algo
    [INFO] WARNING: Exception occurred in dist.init_process_group(backend = 'mpi'). Maybe the process group has been initialized somewhere else.
    [FT][WARNING] Skip NCCL initialization since requested tensor/pipeline parallel sizes are equals to 1.
    [WARNING] gemm_config.in is not found; using default GEMM algo
    Truncation was not explicitly activated but `max_length` is provided a specific value, please use `truncation=True` to explicitly truncate examples to max length. Defaulting to 'longest_first' truncation strategy. If you encode pairs of sequences (GLUE-style) with the tokenizer you can select this strategy more precisely by providing a specific strategy to `truncation`.
    Traceback (most recent call last):
      File "test.py", line 136, in <module>
        output_ids, output_seq_len = model.forward(input_token=input_token,
      File "/root/FasterTransformer/examples/pytorch/t5/utils/ft_decoding.py", line 461, in forward
        results = self.decoding.forward(beam_size,  # optional, can be None
      File "/root/FasterTransformer/examples/pytorch/t5/utils/ft_decoding.py", line 433, in forward
        results = self.decoding.forward(beam_width, max_seq_len,
    RuntimeError: [FT][ERROR] CUDA runtime error: misaligned address /root/FasterTransformer/src/fastertransformer/utils/memory_utils.cu:101
    

    I test a lot of times and found some conditions of error. This error occurred in only FP16, not FP32. And It should be beam_size > 1. If beam_size is 1, any error occurred even if in FP16 mode.

    bug 
    opened by cpm0722 0
  • CUDA Error: (null) FasterTransformer/3rdparty/trt_fused_multihead_attention/fused_multihead_attention.h 232

    CUDA Error: (null) FasterTransformer/3rdparty/trt_fused_multihead_attention/fused_multihead_attention.h 232

    when I run the example code , here is the error message image

    commit: 83776ae7fefaabaf4fc75063d2a704fbeb84f2be GPU: virtual gpu GRID T4-8C driver: 450.102.04 cuda: 11.0 openmpi: 4.1.4 (https://download.open-mpi.org/release/open-mpi/v4.1/openmpi-4.1.4.tar.gz) nccl: 2.10.3-1 (https://github.com/NVIDIA/nccl) cudnn: 8.4.1.50 docker: In my environment, I can not use the nvidia docker image, so I compiled these environment dependencies above

    reproduced step is the example of opt-125m (https://github.com/NVIDIA/FasterTransformer/blob/main/docs/gpt_guide.md#run-meta-opt)

    git lfs clone https://huggingface.co/facebook/opt-125m
    python ../examples/pytorch/gpt/utils/huggingface_opt_convert.py \
          -i opt-125m/ \
          -o opt-125m/c-model/ \
          -i_g 1
    python3 ../examples/pytorch/gpt/opt_summarization.py \
            --summarize \
            --test_hf \
            --max_ite 20 \
            --ft_model_location opt-125m/c-model \
            --hf_model_name opt-125m
    
    opened by DogeWatch 11
  • SmoothQuant with OPT w8a8 - no act_scales from SmoothQuant & why is it needed?

    SmoothQuant with OPT w8a8 - no act_scales from SmoothQuant & why is it needed?

    Branch/Tag/Commit

    N/A

    Docker Image Version

    N/A

    GPU name

    n/a

    CUDA Driver

    n/a

    Reproduced Steps

    In this section, the guide mentioned "You must also generate activation records for calibrating the scaling factors." and this command: python3 examples/pytorch/gpt/utils/huggingface_opt_convert.py -i ../smoothquant/opt-1.3b-smooth/ -o ../nlp-models/ft/test/opt-1.3b-int8/ -i_g 1 -act_scale ../smoothquant/opt-1.3b-smooth.scales.pt

    However, there seems to be some mismatch between input and output of FT and SmoothQuant:

    1. from what I read in huggingface_opt_convert code, the -act_scale is supposed to be the static quantization scale for activations. But not the scale for smoothing the activations. Right?
    2. but SmoothQuant does not generate files for static quantization scales for activations. It generates on the fly and does not have a save function.
    3. on the other hand, SmoothQuant has asave function for Smoothing scales for activations.. We can see from this section in SmoothQuant READMEthat this save function is for smoothing scales but not after-smoothing-quantization scales.

    => so SmoothQuant does not give the act_scale huggingface_opt_convert is asking for.

    At the same time, I don't quite understand why huggingface_opt_convert needs quantization scales from SmoothQuant. Can't we just input the smoothed model with -quantize flag to get quantization scales by this script itself?

    Additional question: if we do get act_scales from SmoothQuant, they have multiple options: per tensor and per token quantization. Does current FasterTransformer w8a8 mode support both? Is there a plan to support both in the future?

    bug 
    opened by chenho74 5
  • How to understand the codebase, how to apply my python functions into this code

    How to understand the codebase, how to apply my python functions into this code

    Hi, I am a python and Huggingface's transformers user. My model is based on T5/BART but with further generating function implementaions by myself in python (add more LogitsProcessor functions in Huggingface code ). How can I use my functions along with FasterTransformer properly? It is really hard for me to read C++ code and understand the flow. Any advice ?

    opened by duongkstn 1
  • [Question] Tips for implementing extra models?

    [Question] Tips for implementing extra models?

    Hi! I'm interested in running Whisper, which is a very standard encoder-decoder, with FT.

    I'm especially interested in the decoder, and was wondering what you would say the fastest route would be. Would you adapt the existing Decoder? Or maybe one of the encoder-decoder models like BART or T5?

    opened by TheExGenesis 5
Releases(release/v5.2.1_tag)
Owner
NVIDIA Corporation
NVIDIA Corporation
Source code for CsiNet and CRNet using Fully Connected Layer-Shared feedback architecture.

FCS-applications Source code for CsiNet and CRNet using the Fully Connected Layer-Shared feedback architecture. Introduction This repository contains

Boyuan Zhang 4 Oct 07, 2022
:hot_pepper: R²SQL: "Dynamic Hybrid Relation Network for Cross-Domain Context-Dependent Semantic Parsing." (AAAI 2021)

R²SQL The PyTorch implementation of paper Dynamic Hybrid Relation Network for Cross-Domain Context-Dependent Semantic Parsing. (AAAI 2021) Requirement

huybery 60 Dec 31, 2022
Bot to connect a real Telegram user, simulating responses with OpenAI's davinci GPT-3 model.

AI-BOT Bot to connect a real Telegram user, simulating responses with OpenAI's davinci GPT-3 model.

Thempra 2 Dec 21, 2022
VampiresVsWerewolves - Our Implementation of a MiniMax algorithm with alpha beta pruning in the context of an in-class competition

VampiresVsWerewolves Our Implementation of a MiniMax algorithm with alpha beta pruning in the context of an in-class competition. Our Algorithm finish

Shawn 1 Jan 21, 2022
Searching keywords in PDF file folders

keyword_searching Steps to use this Python scripts: (1)Paste this script into the file folder containing the PDF files you need to search from; (2)Thi

1 Nov 08, 2021
Implementation of Fast Transformer in Pytorch

Fast Transformer - Pytorch Implementation of Fast Transformer in Pytorch. This only work as an encoder. Yannic video AI Epiphany Install $ pip install

Phil Wang 167 Dec 27, 2022
基于“Seq2Seq+前缀树”的知识图谱问答

KgCLUE-bert4keras 基于“Seq2Seq+前缀树”的知识图谱问答 简介 博客:https://kexue.fm/archives/8802 环境 软件:bert4keras=0.10.8 硬件:目前的结果是用一张Titan RTX(24G)跑出来的。 运行 第一次运行的时候,会给知

苏剑林(Jianlin Su) 65 Dec 12, 2022
Indonesia spellchecker with python

indonesia-spellchecker Ganti kata yang terdapat pada file teks.txt untuk diperiksa kebenaran kata. Run on local machine python3 main.py

Rahmat Agung Julians 1 Sep 14, 2022
Fast, DB Backed pretrained word embeddings for natural language processing.

Embeddings Embeddings is a python package that provides pretrained word embeddings for natural language processing and machine learning. Instead of lo

Victor Zhong 212 Nov 21, 2022
FactSumm: Factual Consistency Scorer for Abstractive Summarization

FactSumm: Factual Consistency Scorer for Abstractive Summarization FactSumm is a toolkit that scores Factualy Consistency for Abstract Summarization W

devfon 83 Jan 09, 2023
Pretrained Japanese BERT models

Pretrained Japanese BERT models This is a repository of pretrained Japanese BERT models. The models are available in Transformers by Hugging Face. Mod

Inui Laboratory 387 Dec 30, 2022
NLP project that works with news (NER, context generation, news trend analytics)

СоАвтор СоАвтор – платформа и открытый набор инструментов для редакций и журналистов-фрилансеров, который призван сделать процесс создания контента ма

38 Jan 04, 2023
Adversarial Examples for Extreme Multilabel Text Classification

Adversarial Examples for Extreme Multilabel Text Classification The code is adapted from the source codes of BERT-ATTACK [1], APLC_XLNet [2], and Atte

1 May 14, 2022
An easier way to build neural search on the cloud

An easier way to build neural search on the cloud Jina is a deep learning-powered search framework for building cross-/multi-modal search systems (e.g

Jina AI 17.1k Jan 09, 2023
Edge-Augmented Graph Transformer

Edge-augmented Graph Transformer Introduction This is the official implementation of the Edge-augmented Graph Transformer (EGT) as described in https:

Md Shamim Hussain 21 Dec 14, 2022
Transcribing audio files using Hugging Face's implementation of Wav2Vec2 + "chain-linking" NLP tasks to combine speech-to-text with downstream tasks like translation and summarisation.

PART 2: CHAIN LINKING AUDIO-TO-TEXT NLP TASKS 2A: TRANSCRIBE-TRANSLATE-SENTIMENT-ANALYSIS In notebook3.0, I demo a simple workflow to: transcribe a lo

Chua Chin Hon 30 Jul 13, 2022
Official implementation of Meta-StyleSpeech and StyleSpeech

Meta-StyleSpeech : Multi-Speaker Adaptive Text-to-Speech Generation Dongchan Min, Dong Bok Lee, Eunho Yang, and Sung Ju Hwang This is an official code

min95 169 Jan 05, 2023
Few-shot Natural Language Generation for Task-Oriented Dialog

Few-shot Natural Language Generation for Task-Oriented Dialog This repository contains the dataset, source code and trained model for the following pa

172 Dec 13, 2022
Code Generation using a large neural network called GPT-J

CodeGenX is a Code Generation system powered by Artificial Intelligence! It is delivered to you in the form of a Visual Studio Code Extension and is Free and Open-source!

DeepGenX 389 Dec 31, 2022
Minimal GUI for accessing the Watson Text to Speech service.

Description Minimal graphical application for accessing the Watson Text to Speech service. Requirements Python 3 plus all dependencies listed in requi

Moritz Maxeiner 1 Oct 22, 2021