EnigmaConsultant/huntr-poc-tensorrt-gridanchor-plugin-oob-read

GridAnchor_TRT plugin deserialize() heap out-of-bounds read via attacker-controlled mNumLayers / numAspectRatios (too-late length guard)

Summary

The GridAnchor_TRT / GridAnchorRect_TRT plugin in NVIDIA TensorRT-OSS reconstructs its state from a serialized blob using unchecked read<> calls. Counts embedded in the blob (mNumLayers, per-layer numAspectRatios, mNumPriors) are trusted without any bound check against the actual blob length. The only length validation PLUGIN_VALIDATE(d == a + length) runs after every read loop — too late to prevent the out-of-bounds access. A short, malicious blob that declares an inflated numAspectRatios (or mNumLayers) makes read<float>() walk off the end of the heap allocation holding the plugin blob, producing an out-of-bounds heap read.

• Target repo: https://github.com/NVIDIA/TensorRT (TensorRT-OSS plugin library, Apache-2.0)
• Version / commit: 11.1.0.106, HEAD a892d22267d9cd2dedc1a0893e6892ac901f6d3d
• Vulnerable file: plugin/gridAnchorPlugin/gridAnchorPlugin.cpp, constructor GridAnchorGenerator::GridAnchorGenerator(void const* data, size_t length, char const* name), lines 126–159
• Helper: plugin/common/plugin.h, read<>(), lines 100–108
• Struct: include/NvInferPluginUtils.h, struct GridAnchorParameters, lines 88–97
• Class: CWE-125 (Out-of-bounds Read)

Trigger path in real TensorRT

IRuntime::deserializeCudaEngine() parses an attacker-supplied .plan / .engine file containing an IPluginV2Ext layer of type GridAnchor_TRT. The engine reader hands the plugin's inline serialized blob to:

GridAnchorBasePluginCreator::deserializePlugin(name, serialData, serialLength)
    -> new GridAnchorGenerator(serialData, serialLength, name)   // vulnerable ctor

serialData / serialLength are copied verbatim out of the plan file, so their contents and the declared counts are fully attacker-controlled.

Root cause (verbatim from the target)

plugin/common/plugin.h lines 100–108 — read<> is an unconditional 4-byte memcpy that advances the cursor with no bounds check:

template <typename OutType, typename BufferType>
OutType read(BufferType const*& buffer)
{
    static_assert(sizeof(BufferType) == 1, "BufferType must be a 1 byte type.");
    OutType val{};
    std::memcpy(&val, static_cast<void const*>(buffer), sizeof(OutType));
    buffer += sizeof(OutType);
    return val;
}

plugin/gridAnchorPlugin/gridAnchorPlugin.cpp lines 126–159 — the deserializing ctor (verbatim):

GridAnchorGenerator::GridAnchorGenerator(void const* data, size_t length, char const* name)
    : mPluginName(name)
{
    char const *d = reinterpret_cast<char const*>(data), *a = d;
    mNumLayers = read<int32_t>(d);                                   // attacker-controlled
    PLUGIN_CUASSERT(cudaMallocHost((void**) &mNumPriors, mNumLayers * sizeof(int32_t)));
    PLUGIN_CUASSERT(cudaMallocHost((void**) &mDeviceWidths, mNumLayers * sizeof(Weights)));
    PLUGIN_CUASSERT(cudaMallocHost((void**) &mDeviceHeights, mNumLayers * sizeof(Weights)));
    mParam.resize(mNumLayers);
    for (int32_t id = 0; id < mNumLayers; id++)
    {
        // we have to deserialize GridAnchorParameters by hand
        mParam[id].minSize = read<float>(d);
        mParam[id].maxSize = read<float>(d);
        mParam[id].numAspectRatios = read<int32_t>(d);              // attacker-controlled
        mParam[id].aspectRatios = (float*) malloc(sizeof(float) * mParam[id].numAspectRatios);
        for (int32_t i = 0; i < mParam[id].numAspectRatios; ++i)
        {
            mParam[id].aspectRatios[i] = read<float>(d);           // <-- OOB READ
        }
        mParam[id].H = read<int32_t>(d);
        mParam[id].W = read<int32_t>(d);
        for (int32_t i = 0; i < 4; ++i)
        {
            mParam[id].variance[i] = read<float>(d);
        }
        mNumPriors[id] = read<int32_t>(d);
        mDeviceWidths[id] = deserializeToDevice(d, mNumPriors[id]);
        mDeviceHeights[id] = deserializeToDevice(d, mNumPriors[id]);
    }
    PLUGIN_VALIDATE(d == a + length);                               // bound check -- TOO LATE
}

Because numAspectRatios (and mNumLayers, mNumPriors) come straight from the blob and read<> performs no bound check, a short blob with an inflated numAspectRatios drives the aspect-ratio loop past the end of the buffer. The only guard, PLUGIN_VALIDATE(d == a + length), executes only after all loops — i.e. after the out-of-bounds reads already occurred.

PoC

A host-only AddressSanitizer harness faithfully reproduces the vulnerable ctor. read<> and struct GridAnchorParameters are copied verbatim; the only substitutions are cudaMallocHost -> malloc (both allocate host-accessible memory) and deserializeToDevice -> an equivalent host cursor-advance. Neither substitution touches the vulnerable read<> sequence or the cursor arithmetic that overflows, so no GPU is required — the overflow is a pure host-side read of the serialized buffer.

Malicious blob (16 bytes): mNumLayers=1, minSize=1.0, maxSize=2.0, numAspectRatios=100, then the blob ends. The aspect-ratio read loop needs 400 more bytes; iteration i=0 already reads at d == a + length, one byte past the allocation.

Files (in this repo):

• gridAnchor_harness.cpp — positive PoC (crashes under ASan)
• gridAnchor_negctl.cpp — negative control (well-formed blob, runs clean)

Build:

clang++ -std=c++17 -O0 -g -fsanitize=address -fno-omit-frame-pointer \
    gridAnchor_harness.cpp -o gridAnchor_asan
clang++ -std=c++17 -O0 -g -fsanitize=address -fno-omit-frame-pointer \
    gridAnchor_negctl.cpp -o gridAnchor_negctl

Captured evidence (verbatim)

Positive PoC (./gridAnchor_asan):

[harness] blob length = 16 bytes; numAspectRatios = 100 (needs 400 extra bytes)
[harness] calling GridAnchorGenerator deserialize ctor...
=================================================================
==355729==ERROR: AddressSanitizer: heap-buffer-overflow on address 0x7b5d3fde0080 at pc 0x55ed96b7061f bp 0x7ffc023324e0 sp 0x7ffc02331ca0
READ of size 4 at 0x7b5d3fde0080 thread T0
    #0 0x55ed96b7061e in __asan_memcpy
    #1 0x55ed96bbd1b6 in float nvinfer1::plugin::read<float, char>(char const*&) gridAnchor_harness.cpp:79:5
    #2 0x55ed96bb8f6a in gridAnchorDeserialize(void const*, unsigned long) gridAnchor_harness.cpp:143:42
    #3 0x55ed96bb82ef in main gridAnchor_harness.cpp:195:9

0x7b5d3fde0080 is located 0 bytes after 16-byte region [0x7b5d3fde0070,0x7b5d3fde0080)
allocated by thread T0 here:
    #0 0x55ed96b729c8 in malloc
    #1 0x55ed96bb824c in main gridAnchor_harness.cpp:186:32

SUMMARY: AddressSanitizer: heap-buffer-overflow in __asan_memcpy
==355729==ABORTING

The overflowing read is at gridAnchor_harness.cpp:143 — mParam[id].aspectRatios[i] = read<float>(d) — exactly the target's line 137 (gridAnchorPlugin.cpp).

Negative control (./gridAnchor_negctl) — well-formed 72-byte blob whose declared counts match the payload:

[negctl] well-formed blob length = 72 bytes
[negctl] deserialize returned cleanly, PLUGIN_VALIDATE passed

No ASan report; PLUGIN_VALIDATE(d == a + length) passes. This confirms the crash is specific to the malformed length/count mismatch and not an artifact of the harness.

Impact

Loading a malicious TensorRT engine/plan (a common distribution artifact, e.g. shipped alongside a model) that contains a crafted GridAnchor_TRT plugin blob triggers an out-of-bounds heap read during deserializeCudaEngine(). Consequences range from denial-of-service (crash) to potential information disclosure (deserialized float values are stored into aspectRatios / device buffers derived from out-of-bounds memory).

Suggested fix

Validate remaining buffer length before each read<> (or before entering each count-driven loop): check that d + N*sizeof(T) <= a + length prior to reading, and reject mNumLayers / numAspectRatios / mNumPriors values that would exceed the blob. The existing PLUGIN_VALIDATE(d == a + length) should be complemented by per-read bound checks.

Dedup note

This is a distinct plugin from previously reported TensorRT-OSS deserialize OOB issues. The same anti-pattern (unchecked read<> with a trailing PLUGIN_VALIDATE) recurs across several plugins; each plugin's ctor is a separate code path and a separate finding:

• DecodeBbox3D / PriorBox / Region / PillarScatter / FlattenConcat / RPROI / EmbLayerNorm — reported separately.
• This report covers GridAnchorGenerator (GridAnchor_TRT / GridAnchorRect_TRT), specifically the mNumLayers and per-layer numAspectRatios count fields.

No CVE is currently assigned to this specific GridAnchorGenerator deserialize path at the time of writing.