feat: add InternVLA-N1-DualVLN vision-language navigation support

CHANGELOG.md

@@ -1,5 +1,8 @@
 # Release Notes
 
+## Unreleased
+- Added InternVLA-N1-DualVLN vision-language navigation support: direct export of the Qwen2.5-VL planner with the z_latents bridge folded into the graph, and an experimental System-1 runtime for the flow-matching trajectory head
+
 ## 0.10.1
 - Added experimental Dual NVIDIA DGX Spark support with TP=2 tensor-parallel inference over MPI and NCCL
 - Redesigned the experimental OpenAI-compatible server for faster cold launches and lower memory usage

cpp/runtime/config/llmEngineConfig.cpp

@@ -511,6 +511,7 @@ void parseCoreFields(Json const& configJson, LLMEngineConfig& cfg)
     cfg.numKVHeads = getRequired<int32_t>(configJson, "num_key_value_heads");
     cfg.headDim = getRequired<int32_t>(configJson, "head_dim");
     cfg.hiddenSize = getRequired<int32_t>(configJson, "hidden_size");
+    cfg.outputHiddenSize = configJson.value("output_hidden_size", cfg.hiddenSize);
 
     // Top-level: kv_cache_dtype. Required — all current exports write this.
     parseRequiredStateDtype(configJson, "kv_cache_dtype", cfg.kvCacheDtype);
@@ -531,6 +532,7 @@ void parseCoreFields(Json const& configJson, LLMEngineConfig& cfg)
     requirePositive(cfg.numKVHeads, "num_key_value_heads");
     requirePositive(cfg.headDim, "head_dim");
     requirePositive(cfg.hiddenSize, "hidden_size");
+    requirePositive(cfg.outputHiddenSize, "output_hidden_size");
     requirePositive(cfg.maxSupportedBatchSize, "max_batch_size");
     requirePositive(cfg.maxSupportedInputLength, "max_input_len");
     requirePositive(cfg.maxKVCacheCapacity, "max_kv_cache_capacity");

cpp/runtime/config/llmEngineConfig.h

@@ -67,7 +67,16 @@ struct Gemma4MTPKVSharingEntry
 struct LLMEngineConfig
 {
     // --- Core model dimensions ---
-    int32_t hiddenSize{};                //!< Model hidden dimension
+    int32_t hiddenSize{}; //!< Model hidden dimension
+
+    //! Width of the engine's `hidden_states` output, when it differs from
+    //! `hiddenSize`. A model may project the hidden states before emitting them --
+    //! InternVLA-N1 folds its norm + cond_projector into the graph and emits a
+    //! 768-wide bridge tensor -- and the runtime must then allocate and copy that
+    //! width rather than the model width. Parsed from the optional
+    //! `output_hidden_size` key; defaults to `hiddenSize`, so every model that
+    //! emits hidden states at model width is unaffected.
+    int32_t outputHiddenSize{};
     int32_t outputVocabSize{};           //!< Actual output vocab (reduced if vocab reduction active)
     int32_t numAttentionLayers{};        //!< Number of attention layers needing KV cache
     int32_t numKVHeads{};                //!< Number of key-value heads

cpp/runtime/llmRankRuntime.cpp

@@ -1276,8 +1276,8 @@ bool LLMRankRuntime::handleRequest(LLMGenerationRequest const& request, LLMGener
         int32_t const prefillSequenceLength
             = *std::max_element(context.effectivePrefillLengths.begin(), context.effectivePrefillLengths.end());
         mPipelineIO->streamingPrefill.populateFromPrefill(mPipelineIO->inputsEmbeds, mPipelineIO->outputHiddenStates,
-            activeBatchSize, prefillSequenceLength, mDeployment.base.hiddenSize, mMaxRuntimeBatchSize,
-            mDeployment.base.maxSupportedInputLength, stream);
+            activeBatchSize, prefillSequenceLength, mDeployment.base.hiddenSize, mDeployment.base.outputHiddenSize,
+            mMaxRuntimeBatchSize, mDeployment.base.maxSupportedInputLength, stream);
         mLastPrefillLength = prefillSequenceLength;
         mLastInputTokenIds = context.rawBatchedInputIds;
         mHiddenStatesRegistry[0] = &mPipelineIO->streamingPrefill.inputEmbeds;

cpp/runtime/state/pipelineIO.cpp

@@ -130,24 +130,26 @@ void bindBackboneTensorMap(
 } // namespace
 
 void StreamingPrefillBuffers::populateFromPrefill(Tensor const& liveInputEmbeds, Tensor const& liveEngineHiddenStates,
-    int32_t batch, int32_t prefillLen, int32_t hiddenSize, int32_t maxBatch, int32_t maxSeq, cudaStream_t stream)
+    int32_t batch, int32_t prefillLen, int32_t hiddenSize, int32_t outputHiddenSize, int32_t maxBatch, int32_t maxSeq,
+    cudaStream_t stream)
 {
     auto const dtype = nvinfer1::DataType::kHALF;
     if (inputEmbeds.isEmpty())
     {
         inputEmbeds = Tensor(
             {maxBatch, maxSeq, hiddenSize}, DeviceType::kGPU, dtype, "PipelineIO::streamingPrefill.inputEmbeds");
-        engineHiddenStates = Tensor(
-            {maxBatch, maxSeq, hiddenSize}, DeviceType::kGPU, dtype, "PipelineIO::streamingPrefill.engineHiddenStates");
+        engineHiddenStates = Tensor({maxBatch, maxSeq, outputHiddenSize}, DeviceType::kGPU, dtype,
+            "PipelineIO::streamingPrefill.engineHiddenStates");
     }
     check::check(inputEmbeds.reshape({batch, prefillLen, hiddenSize}), "Tensor reshape failed");
-    check::check(engineHiddenStates.reshape({batch, prefillLen, hiddenSize}), "Tensor reshape failed");
+    check::check(engineHiddenStates.reshape({batch, prefillLen, outputHiddenSize}), "Tensor reshape failed");
 
-    size_t const bytes = static_cast<size_t>(batch) * prefillLen * hiddenSize * sizeof(__half);
+    size_t const embedBytes = static_cast<size_t>(batch) * prefillLen * hiddenSize * sizeof(__half);
+    size_t const hiddenBytes = static_cast<size_t>(batch) * prefillLen * outputHiddenSize * sizeof(__half);
     CUDA_CHECK(cudaMemcpyAsync(
-        inputEmbeds.rawPointer(), liveInputEmbeds.rawPointer(), bytes, cudaMemcpyDeviceToDevice, stream));
-    CUDA_CHECK(cudaMemcpyAsync(
-        engineHiddenStates.rawPointer(), liveEngineHiddenStates.rawPointer(), bytes, cudaMemcpyDeviceToDevice, stream));
+        inputEmbeds.rawPointer(), liveInputEmbeds.rawPointer(), embedBytes, cudaMemcpyDeviceToDevice, stream));
+    CUDA_CHECK(cudaMemcpyAsync(engineHiddenStates.rawPointer(), liveEngineHiddenStates.rawPointer(), hiddenBytes,
+        cudaMemcpyDeviceToDevice, stream));
 }
 
 void bindRopeTensors(TensorMap& map, PipelineIO& io, SharedResources& res, LLMEngineConfig const& cfg)
@@ -457,7 +459,7 @@ PipelineIO PipelineIO::createForLLM(LLMEngineConfig const& cfg, cudaStream_t str
     // streaming consumers (Qwen3-Omni Talker) read it; if the engine emits
     // hidden_states but no consumer is set, the buffer is harmless write-target;
     // if the engine has no hidden_states output the binding is silently skipped.
-    io.outputHiddenStates = Tensor({cfg.maxSupportedBatchSize, maxSeqLen, cfg.hiddenSize}, DeviceType::kGPU,
+    io.outputHiddenStates = Tensor({cfg.maxSupportedBatchSize, maxSeqLen, cfg.outputHiddenSize}, DeviceType::kGPU,
         nvinfer1::DataType::kHALF, "PipelineIO::outputHiddenStates");
 
     if (cfg.ropeConfig.type == RopeType::kMRope)

cpp/runtime/state/pipelineIO.h

@@ -45,13 +45,20 @@ struct StreamingPrefillBuffers
     Tensor inputEmbeds;        //!< Prefill-time layer-0 input embeddings.
     Tensor engineHiddenStates; //!< Prefill-time engine hidden_states output.
 
-    //! Allocate on first call (sized to the worst case `{maxBatch, maxSeq, hiddenSize}`),
-    //! reshape to the current request's `{batch, prefillLen, hiddenSize}`, and copy
-    //! from the live PipelineIO buffers on `stream`. Subsequent calls reuse the same
-    //! allocation. Must be invoked after prefill and before the first decode step on
-    //! the same stream so the copies precede any overwrite of `outputHiddenStates`.
+    //! Allocate on first call (sized to the worst case `{maxBatch, maxSeq, dim}`),
+    //! reshape to the current request's `{batch, prefillLen, dim}`, and copy from the
+    //! live PipelineIO buffers on `stream`. Subsequent calls reuse the same allocation.
+    //! Must be invoked after prefill and before the first decode step on the same
+    //! stream so the copies precede any overwrite of `outputHiddenStates`.
+    //!
+    //! The two tensors are sized independently: input embeddings are always model
+    //! width, but a model may project its hidden states before emitting them, in
+    //! which case `outputHiddenSize` is the narrower emitted width. Copying the
+    //! model width out of a narrower engine output would read past the valid data
+    //! and hand the consumer a plausible-looking buffer whose tail is garbage.
     void populateFromPrefill(Tensor const& liveInputEmbeds, Tensor const& liveEngineHiddenStates, int32_t batch,
-        int32_t prefillLen, int32_t hiddenSize, int32_t maxBatch, int32_t maxSeq, cudaStream_t stream);
+        int32_t prefillLen, int32_t hiddenSize, int32_t outputHiddenSize, int32_t maxBatch, int32_t maxSeq,
+        cudaStream_t stream);
 };
 
 //! All tensors flowing through the inference pipeline.

docs/source/user_guide/examples/vla/index.md

@@ -8,14 +8,19 @@ runtime:
 |---|---|---|---|
 | [Alpamayo-R1](alpamayo.md) | camera frames, instruction, past trajectory | future acceleration/curvature trajectory | `action_inference` |
 | [Cosmos3-Edge policy](cosmos3.md) | observation image or frame list, instruction | robot action chunk | `cosmos3_policy_inference` |
+| [InternVLA-N1-DualVLN](internvla_n1.md) | navigation frames, instruction | future waypoint trajectory | `internvla_n1_dual_system_inference` |
 
-Both workflows export on CPU, build all required TensorRT engines on the target,
-and invoke one end-to-end runtime executable.
+These workflows export on CPU, build all required TensorRT engines on the target,
+and invoke one end-to-end runtime executable. InternVLA-N1 is the exception in one
+respect: its two systems run at different rates, so System 2 plans on a background
+thread while System 1 keeps sampling from the newest plan, rather than running the
+whole graph in one pass.
 
 ```{toctree}
 :maxdepth: 1
 :hidden:
 
 alpamayo.md
 cosmos3.md
+internvla_n1.md
 ```