TensorRT LLM is an open-sourced library for optimizing Large Language Model (LLM) inference. It provides state-of-the-art optimizations, including custom attention kernels, inflight batching, paged KV caching, quantization (FP8,[FP4](https://www.nvidia.com/en-us/data-center/technologies/blackwell-architecture/), INT4[AWQ](https://arxiv.org/abs/2306.00978), INT8[SmoothQuant](https://arxiv.org/abs/2211.10438), ...), speculative decoding, and much more, to perform inference efficiently on NVIDIA GPUs.

[Architected on PyTorch](https://github.com/NVIDIA/TensorRT-LLM/blob/main/docs/source/torch/arch_overview.md), TensorRT LLM provides a high-level Python[LLM API](https://nvidia.github.io/TensorRT-LLM/quick-start-guide.html#llm-api) that supports a wide range of inference setups - from single-GPU to multi-GPU or multi-node deployments. It includes built-in support for various parallelism strategies and advanced features. The LLM API integrates seamlessly with the broader inference ecosystem, including NVIDIA[Dynamo](https://github.com/ai-dynamo/dynamo) and the[Triton Inference Server](https://github.com/triton-inference-server/server).

[Architected on PyTorch](https://github.com/NVIDIA/TensorRT-LLM/blob/release/1.1/docs/source/developer-guide/overview.md), TensorRT LLM provides a high-level Python[LLM API](https://nvidia.github.io/TensorRT-LLM/quick-start-guide.html#llm-api) that supports a wide range of inference setups - from single-GPU to multi-GPU or multi-node deployments. It includes built-in support for various parallelism strategies and advanced features. The LLM API integrates seamlessly with the broader inference ecosystem, including NVIDIA[Dynamo](https://github.com/ai-dynamo/dynamo) and the[Triton Inference Server](https://github.com/triton-inference-server/server).

TensorRT LLM is designed to be modular and easy to modify. Its PyTorch-native architecture allows developers to experiment with the runtime or extend functionality. Several popular models are also pre-defined and can be customized using[native PyTorch code](./tensorrt_llm/_torch/models/modeling_deepseekv3.py), making it easy to adapt the system to specific needs.

andkspec.cross_mha==False

andkspec.flash_attention==True

andkspec.input_layout!=InputLayout.SEPARATE_Q_K_V)

or (kspec.sm==100

andkspec.dtypein ['fp16','bf16','fp16_fp32','e4m3','e4m3_fp32']

andkspec.head_size==72

andkspec.head_size_v==0

andkspec.sage_block_sizesisNone

andkspec.version==2

andkspec.cross_mha==False

andkspec.flash_attention==True

andkspec.input_layout!=InputLayout.SEPARATE_Q_K_V)

or (kspec.smin [90,100,120]

andkspec.dtypein ['bf16','e4m3_fp32']

PerCudaCtxPerThreadSingletonCreator(CreatorFunc creator, DeleterFunc deleter)

:mCreator{std::move(creator)}

,mDeleter{std::move(deleter)}

,mObservers{new std::unordered_map<CacheKey, std::weak_ptr<T>, hash<CacheKey>>()}

{

}

~PerCudaCtxPerThreadSingletonCreator()

{

std::lock_guard<std::mutex> lk{mMutex};

deletemObservers;

mObservers =nullptr;

}

std::shared_ptr<T>operator()()

{

std::lock_guard<std::mutex> lk{mMutex};

CUcontext ctx{getCurrentCudaCtx()};

std::thread::id thread =std::this_thread::get_id();

autoconst key =std::make_tuple(ctx, thread);

std::shared_ptr<T> result =mObservers[key].lock();

std::shared_ptr<T> result =(*mObservers)[key].lock();

if (result ==nullptr)

{

TLLM_LOG_TRACE("creating singleton instance for CUDA context %lu and thread %lu", ctx, thread);

}

mDeleter(obj);

if (mObservers ==nullptr)

{

return;

}

std::shared_ptr<T> observedObjHolder;// Delay destroy to avoid dead lock.

if (mObservers.find(key) ==mObservers.end())

auto it =mObservers->find(key);

if (it ==mObservers->end())

{

return;

}

observedObjHolder =mObservers.at(key).lock();

observedObjHolder =it->second.lock();

if (observedObjHolder ==nullptr)

{

mObservers.erase(key);

mObservers->erase(it);

}

}};

mObservers.at(key) = result;

(*mObservers)[key] = result;

}

{

mutable std::mutexmMutex;

using CacheKey = std::tuple<CUcontext, std::thread::id>;

std::unordered_map<CacheKey, std::weak_ptr<T>, hash<CacheKey>>mObservers;

std::unordered_map<CacheKey, std::weak_ptr<T>, hash<CacheKey>>*mObservers;

};

{

size_t free_mb;

size_t total_mb;

float free_percent;

};

MemoryInfogetMemoryInfo()

{

size_t free_mem =0, total_mem =0;

TLLM_CUDA_CHECK(cudaMemGetInfo(&free_mem, &total_mem));

size_tconst free_mb = free_mem / (1024 *1024);

size_tconst total_mb = total_mem / (1024 *1024);

floatconst free_percent = (total_mem >0) ? (static_cast<float>(free_mem) / total_mem *100.0f) :0.0f;

return {free_mb, total_mb, free_percent};

}

voidlogMemoryUsage(charconst* operation, CUcontext ctx)

{

autoconst mem =getMemoryInfo();

TLLM_LOG_DEBUG("%s: Context=%p, Free Memory=%zu MB (%.1f%%), Total=%zu MB", operation, ctx, mem.free_mb,

mem.free_percent, mem.total_mb);

}

voidthrowCublasErrorWithMemInfo(charconst* operation, CUcontext ctx, cublasStatus_t status)

{

autoconst mem =getMemoryInfo();

TLLM_THROW(

"Consider reducing kv_cache_config.free_gpu_memory_fraction.",

operation, status, ctx, mem.free_mb, mem.free_percent, mem.total_mb);

}

}// namespace

std::shared_ptr<cublasHandle_t>getCublasHandle()

{

static PerCudaCtxPerThreadSingletonCreator<cublasHandle_t>creator(

[]() ->auto

{

auto handle = std::unique_ptr<cublasHandle_t>(new cublasHandle_t);

TLLM_CUDA_CHECK(cublasCreate(handle.get()));

CUcontext ctx =getCurrentCudaCtx();

logMemoryUsage("Creating cublas handle", ctx);

auto handle = std::make_unique<cublasHandle_t>();

auto status =cublasCreate(handle.get());

if (status != CUBLAS_STATUS_SUCCESS)

{

throwCublasErrorWithMemInfo("cublas handle", ctx, status);

}

return handle;

},

[](cublasHandle_t* handle)

{

TLLM_CUDA_CHECK(cublasDestroy(*handle));

auto status =cublasDestroy(*handle);

if (status != CUBLAS_STATUS_SUCCESS)

{

TLLM_LOG_WARNING("Failed to destroy cublas handle. Status: %d", status);

}

delete handle;

handle =nullptr;

});

returncreator();

}

static PerCudaCtxPerThreadSingletonCreator<cublasLtHandle_t>creator(

[]() ->auto

{

auto handle = std::unique_ptr<cublasLtHandle_t>(new cublasLtHandle_t);

TLLM_CUDA_CHECK(cublasLtCreate(handle.get()));

CUcontext ctx =getCurrentCudaCtx();

logMemoryUsage("Creating cublasLt handle", ctx);

auto handle = std::make_unique<cublasLtHandle_t>();

auto status =cublasLtCreate(handle.get());

if (status != CUBLAS_STATUS_SUCCESS)

{

throwCublasErrorWithMemInfo("cublasLt handle", ctx, status);

}

return handle;

},

[](cublasLtHandle_t* handle)

{

TLLM_CUDA_CHECK(cublasLtDestroy(*handle));

auto status =cublasLtDestroy(*handle);

if (status != CUBLAS_STATUS_SUCCESS)

{

TLLM_LOG_WARNING("Failed to destroy cublasLt handle. Status: %d", status);

}

delete handle;

handle =nullptr;

});

returncreator();

}

{

if (sm ==89 || sm >=120)

{

return {CutlassTileConfig::CtaShape16x256x128_WarpShape16x64x128,

CutlassTileConfig::CtaShape32x128x64_WarpShape32x32x64,

return {CutlassTileConfig::CtaShape32x128x64_WarpShape32x32x64,

CutlassTileConfig::CtaShape64x128x64_WarpShape64x32x64,

CutlassTileConfig::CtaShape64x64x128_WarpShape32x64x64,

CutlassTileConfig::CtaShape128x64x64_WarpShape64x32x64,

CutlassTileConfig::CtaShape128x256x64_WarpShape64x64x64,

CutlassTileConfig::CtaShape256x128x64_WarpShape64x64x64};

CutlassTileConfig::CtaShape256x128x64_WarpShape64x64x64,

CutlassTileConfig::CtaShape16x256x128_WarpShape16x64x128};

}

{

,mUseTllmGen(tensorrt_llm::common::isSM100Family() && fixedParams.headSize !=80)

,mUseTllmGen(tensorrt_llm::common::isSM100Family() && fixedParams.headSize !=80 && fixedParams.headSize !=72)

{

if (mUseTllmGen)

{

_{FP8 H100, FP16 A100, SXM 80GB GPUs, TP1, ISL/OSL's provided, TensorRT LLM v0.5.0., TensorRT 9.1}

The full data behind these charts & tables and including larger models with higher TP values can be found in TensorRT LLM's[Performance Documentation](https://nvidia.github.io/TensorRT-LLM/latest/performance/perf-overview.html)

The full data behind these charts & tables and including larger models with higher TP values can be found in TensorRT LLM's[Performance Documentation](https://nvidia.github.io/TensorRT-LLM/0.21.0/performance/perf-overview.html)

Stay tuned for a highlight on Llama coming soon!

^{*(1) Largest batch supported on given TP configuration by power of 2.*} ^{*(2) TP = Tensor Parallelism*}

Additional Performance data is available on the[NVIDIA Data Center Deep Learning Product Performance](https://developer.nvidia.com/deep-learning-performance-training-inference/ai-inference) page, & soon in[TensorRT LLM's Performance Documentation](https://nvidia.github.io/TensorRT-LLM/latest/performance/perf-overview.html).

Additional Performance data is available on the[NVIDIA Data Center Deep Learning Product Performance](https://developer.nvidia.com/deep-learning-performance-training-inference/ai-inference) page, & soon in[TensorRT LLM's Performance Documentation](https://nvidia.github.io/TensorRT-LLM/0.21.0/performance/perf-overview.html).

Dynamo also includes built-in support for Kubernetes deployment, monitoring, and metrics collection. The development team is actively working on enabling dynamic instance scaling, further enhancing its suitability for production environments.

For more information on how to use Dynamo with TensorRT-LLM, please refer to[this documentation](https://docs.nvidia.com/dynamo/latest/examples/trtllm.html).

For more information on how to use Dynamo with TensorRT-LLM, please refer to[this documentation](https://docs.nvidia.com/dynamo/latest/backends/trtllm/README.html).

..toctree::