[RFC]: Changes in vLLM Model Development

[WoosukKwon]

Motivation.

vLLM has long pursued the goal of building a unified, fast inference engine across diverse models and hardware. We remain committed to this goal, but we have come to believe that vLLM's current model implementation principles need to be revisited in order to improve both performance and development velocity. Two lessons stand out:

• We were overly reluctant to modify model code directly for op fusion and kernel dispatching. Instead, we accumulated abstractions and compiler passes that ultimately made the code opaque and hard to optimize.
• We pursued the unrealistic ideal of a single model definition that works well on every hardware. In practice, each hardware backend benefits from its own model implementation, since different hardware favors different fusion strategies.
• AI coding agents have made manual op fusion and kernel writing much easier, and the agents work best with raw model code.

To address these issues, we are making three fundamental changes to how vLLM implements and optimizes models:

1. Removing full-graph torch.compile requirement
2. Separation of model code across hardware backends
3. A clear model interface

1. Removing full-graph torch.compile requirement

vLLM enables full-graph torch.compile by default and has relied on it for model-specific optimizations. Two problems have emerged:

• Performance. Achieving speed-of-light performance on target models on the latest hardware within a tight timeline requires manual optimization. Compiler-driven optimization alone has not been sufficient, and timely performance gains are critical to vLLM's adoption.
• Developer experience. torch.compile increases the vLLM startup time and imposes the opacity and constraints inherent to any compiler, both of which slow iteration on model code.

We are therefore removing vLLM's reliance on full-graph torch.compile and restructuring model definitions accordingly. The migration must preserve the current user-facing UX and introduce no performance regressions.

NOTE: we will keep using torch.compile to fuse adjacent ops locally, whenever we find it to be beneficial. The removal is only about full-graph compilation.

1-1. Compiler fusion → manual fusion

Historically, vLLM kept model code minimal (pure PyTorch ops) and relied on compiler passes to fuse them. We are reversing this convention: fusion should be expressed directly in model code. Concretely, model code should look like:

class Attention(nn.Module):
    def __init__(self, ...):
        ...
        self.o_proj = RowParallelLinear(..., reduce_results=False)

    def forward(self, ...):
        ...
        hidden_states = self.o_proj(attn_out)  # pre-allreduce
        return hidden_states


class Layer(nn.Module):
    def __init__(self, ...):
        self.attn = Attention(...)
        self.rms_norm = RMSNorm(...)
        self.mlp = MLP(...)

    def forward(self, hidden_states):
        ...
        hidden_states = self.attn(hidden_states)  # pre-allreduce
        # Call the fused kernel here.
        residual, quantized_inputs = add_rms_norm_quant(
            hidden_states,
            residual,
            self.rms_norm.weights,
            self.rms_norm.eps,
            # act_quant_kwargs depends on the quantization scheme
            self.mlp.act_quant_kwargs,
            do_allreduce=self.tp_size > 1,
        )
        hidden_states = self.mlp(quantized_inputs)  # pre-allreduce
        return hidden_states

Here, add_rms_norm_quant is the call site for the fused kernel. We will introduce such "big fused ops" directly in model code wherever manual fusion is beneficial. This way, we can replace all fusions in vLLM’s compiler passes and even further improve the performance in complex cases such as DeepSeek V3.2 that the current compiler optimizations do not handle well.

1-2. CustomOp / vLLM IR

With manual fusion in place — and the hardware-backend code separation discussed below — there is little reason to further develop the CustomOp and vLLM IR abstractions. We will deprecate both and only keep the minimal part of it We will keep them for legacy models, but more recent and complex models will instead have fused ops with a custom dispatching logic (for example, when different kernels must be launched on Hopper vs. Blackwell) possibly with some shared utility methods. For example, the code may look like:

def add_rms_norm_quant(
	hidden_states: torch.Tensor,
	residual: torch.Tensor,
	rms_weights: torch.Tensor,
	rms_eps: float,
	mlp_act_quant_kwargs: dict[str, Any],
	do_allreduce: bool,
) -> tuple[torch.Tensor, Any]:
	if not do_allreduce:
		# No comm. Launch a simple fused kernel.
		return ...

	device_info = get_device_info()
	if device_info.name in ("GB200", "GB300"):
		# Launch FlashInfer kernel
		...
	elif device_info.name in ("H100", "H200"):
		# Launch some other kernels
 		...

1-3. Piecewise CUDA graph

To preserve the piecewise CUDA graph feature without relying on torch.compile, we can adopt the "breakable CUDA graph" approach from SGL. @ZJY0516 already implemented a prototype in #42304 We expect this to match the performance of the current compiler-based approach without imposing many constraints on model code.

1-4. Hacks for compiler compatibility

vLLM has accumulated several hacks to keep model code compatible with torch.compile. One example is forward_context, a global variable used to bypass the compiler's argument checking. These workarounds have hurt readability and introduced rough edges — model inputs are scattered across multiple places, and some global objects are not freed when the vLLM engine is deleted. Removing these hacks will improve the project's maintainability.

2. Hardware separation and model isolation

vLLM currently uses a single model definition for all in-tree hardware backends and relies on CustomOp and torch.compile to apply different fusions and dispatch kernels. This has made the code unnecessarily complex and fragile: a change made for one backend can easily break or regress another.

Going forward, we will separate model code by hardware vendor. Each model definition can then evolve independently for its target hardware without risk of breaking the others. Just as importantly, each vendor can design fusions around the kernels available to them, rather than conforming to a shared abstraction that fits no one perfectly.

Concretely, we expect the model code to be organized as follows:

models/
    deepseek_v4/
        nvidia/
            model.py
            kernels/
                fused_q_kv_norm.py
            tests/
                test_model.py
                test_fused_q_kv_norm.py
        amd/
            model.py
    deepseek_v3_2/
        nvidia/
            model.py
            kernels/
                fused_q.py
            tests/
                test_model.py
                test_fused_q.py
        amd/
            ...

Here, kernels contains model-specific JIT kernels (and maybe AoT kernels depending on our build structure); commonly used kernels remain in their existing locations. Each model directory also includes unit tests, so we can easily verify basic correctness of the model and its kernels. The code organization will help aggressively leverage AI to optimize models, as well as cleanly deprecate old models from the vLLM codebase.

3. Model interface

vLLM relies on many implicit assumptions about models and does not cleanly separate model-agnostic logic from model-dependent logic. This has produced a poor developer experience for users bringing custom in-house models. To address this, we are introducing explicit interfaces for both model definition and configuration.

First, we will define vLLM's own ModelConfig to hold all the information vLLM needs about a model. It is essentially a union of the typical HF config and vLLM's model protocol. The idea is similar to #24384.

@dataclass
class ModelConfig(ABC):
    num_layers: int
    max_model_len: int
    vocab_size: int
    hidden_states_dtype: torch.dtype
    num_kv_heads: int
    num_query_heads: int
    is_moe: bool
    num_shared_experts: int | None
    num_routed_experts: int | None
    experts_per_tok: int | None
    supports_mm: bool
    supports_pp: bool
    supports_lora: bool
    ...
    model_kwargs: dict[str, Any]  # e.g., "rms_norm_eps": 1e-05

This brings three benefits:

1. A single source of truth for all model information.
2. We currently spawn a subprocess to read certain fields from the model definition; this config eliminates that overhead, which is often dominated by import time.
3. Model developers no longer need to port their models into HF format to use vLLM. In an extreme case, they can simply hardcode the config with values they already know.

Second, we will introduce vLLM's own model interface to make all contracts explicit. It should look like:

class Model(ABC):
    def __init__(
        self,
        config: ModelConfig,
        parallel_state: ParallelState,
        **kwargs,
    ) -> None:
        # Allows the model to manage its own state.
        raise NotImplementedError

    @abstractmethod
    def load_weights(self, use_dummy: bool) -> None:
        ...

    @abstractmethod
    def init_kv_cache(self) -> None:
        # KV cache-related API (Needs more thought).
        ...

    def optimize_and_warmup(self) -> None:
        return None

    def add_requests(self, req_indices: list[int], request_data: list[Any]) -> None:
        # For initializing model-specific state (if any) when a request first arrives.
        return None

    @abstractmethod
    def prepare_inputs(self, input_batch: InputBatch) -> dict[str, Any]:
        # Prepare model inputs for a step. May include building attention metadata.
        ...

    @abstractmethod
    def forward(self, **kwargs) -> torch.Tensor:
        # Run the model and return the hidden states.
        ...

    def compute_logits(self, hidden_states: torch.Tensor) -> torch.Tensor:
        raise NotImplementedError

    # More optional APIs for LoRA, spec decoding, and multimodal inputs.
    ...

The key idea is to let the model manage its own state when needed. For example, Qwen-VL models can manage their pre-computed M-RoPE values within this framework without any changes to the upper-level model runner. Similarly, prepare_inputs lets each model define its own inputs. With this design, the model runner handles only model-agnostic work — bookkeeping for continuous batching, sampling, CUDA graphs, and parts of model parallelism — while all model-dependent work lives in the model layer. Part of this idea was already implemented as ModelState in model runner V2.

Any model definition that implements both the Model and ModelConfig interfaces should be fully supported by vLLM, with no additional requirements.

Proposed Change.

In-Progress

• @ZJY0516 implements PW CUDA graph without relying on torch.compile [Experimental] Breakable CUDA graph #42304
• @WoosukKwon implements the initial hardware separation for DeepSeek V4 & @tjtanaa will handle the AMD implementation
• @mgoin tracks and ports the fusions currently implemented through the compiler-based approach [RFC]: Porting compiler fusions to manual fusion #43224

TODOs

1. When PW CUDA graph & manual fusions are ready, remove @supports_torch_compile for DeepSeek V4
2. Port more models (e.g., Kimi, GLM, Qwen, Minimax, Nemotron, etc.) to manual fusion
3. When all models are ported with manual fusion, start deprecate vLLM's torch.compile integration
4. (Orthogonally to the above) Define model interface & apply it to all models

Detailed timeline will be updated soon.

Feedback Period.

No response

CC List.

No response

Any Other Things.

No response

Before submitting a new issue...

• Make sure you already searched for relevant issues, and asked the chatbot living at the bottom right corner of the documentation page, which can answer lots of frequently asked questions.

[ProExpertProg]

Thanks for the detailed RFC! I know the decision has been made, but adding some comments for completeness. Comments below are opinions, which are fully my own. Tl;dr: I support the model definition overhaul and not requiring model-level torch.compile for head models. I also believe model-level torch.compile has fundamental benefits and addressable downsides. There are better middle-of-the-road solutions, with order-of-magnitude lower migration costs; most models are simple enough (and stabilized) to support a single, portable, and compile-friendly definition.

The new model abstraction is really good. Moving model state & input preparation from the model runner to the model is the correct approach with better separation of concerns, allowing cleaner code and better extensibility. Different models deserve different levels of portability & abstraction: DSV4 might have a specialized nvfp4 B200 implementation with more manual fusion while Qwen 2.5 can have a single model definition for all hardware platforms and quantization schemes. However, this restructuring is orthogonal to removing model-level torch.compile.

None of the known issues with torch.compile (some mentioned above) are insurmountable. The forward_context is an artifact of poor model abstraction and bookkeeping, and mostly orthogonal to torch.compile. Torch.compile did make it worse, but some rough edges were already fixed in 2.11, and more could be done to reduce the impact. Moving the model state from the context to the model fully resolves this issue. Other known issues:

• Startup time: on track (with 2.12) to <1s warm start, <10-20s cold start
• Multi-stream support: fixed in 2.12
• nvtx ranges: only mentioned recently, could be resolved
• Custom op wrapping: can be improved

In general, new issues might pop up and require upstream torch fixes. But custom op wrapping provides a feature-complete escape hatch that could be improved further, and the cited issues are trending towards resolution. While the accumulated frustration is understandable, it should not be mistaken for inevitability.

Conversely, model-level torch.compile brings many benefits, especially for tail models:

• Portable & configurable fusions across abstractions that compose
• Many automatic fusions, memory planning, copy elimination, lower CPU overhead
• Built-in kernel autotuning for the explosion of tunable kernel DSLs and hardwares
• Automatic support for platforms that require torch.compile (TPU, CPU, Spyre, …)

While some of the fusions can be done manually, they will not scale across quantization schemes and hardware platforms, and will be harder to configure and combine, reducing end-user control. For example, there’s no plan for SP+AsyncTP, a very intrusive transformation in eager mode. Finally, a standardized linear representation of operations is fundamentally valuable for inspection (#39215), and OOT compiler backends like megakernels or LLM compilers.

I believe there exist middle-of-the-road solutions, allowing eager and compiled implementations of fusions and piecewise cudagraph to live alongside each other. Specifically, vLLM IR is making huge strides allowing fusions to be equivalent in eager and compile modes while keeping the portability benefit. Alternatives include:

• Remove support_torch_compile from head models
• Turn compilation off by default
• Simplify/remove customizations to torch.compile
• Turn off fullgraph and allow Dynamo graph breaks

The only argument for removing torch.compile in portable model definitions is that shared abstractions like attention, moe, linear, etc. have to conform to torch.compile requirements. But this contradicts the notion that model definitions will be more separate: if there is less reuse within head models, shared abstractions can remain compile-compatible without a burden on head models. This is a small price to pay for portability across both quantization schemes and hardware generations and platforms.

The RFC also cites coding agents’ ability to write kernels and manual fusions. But agents lower the cost of writing code, not maintaining it (and there’s more to maintain). Duplicated definitions will accumulate bugs independently, and maintenance costs will increase over time. For head models with specialized hardware implementations, agent-assisted manual fusion makes sense. But for the long tail of portable models, the maintenance burden outweighs the benefit. Agents also tend to perform much worse on non-CUDA hardware.

At the end of the day, we are choosing to incur a massive migration cost: reimplementing all fusions manually and duplicating model definitions. Instead, we could just disable torch.compile for head models, refactor model definitions, and achieve the same goals without the multiple person-years in migration costs.

[noooop]

This is the best time to send this meme.

vLLM: We can handwritten kernel faster than torch.compile, so we removed it because it caused us too much trouble. LOL.

In fact, we are only going to remove full-graph torch.compile for now.

---

Split prefill and decode attention(v0) vs unified_attention(v1):

When PD separation is not used, the same batch may contain both prefill tokens and decode tokens.

One example is forward_context, a global variable used to bypass the compiler's argument checking.

In v1, we use unified_attention, which in fact only uses flash_attn_varlen_func.

• [RFC]: hide continuous batching complexity through forward context #9098
• [Kernel] Attention.forward with unified_attention when use_direct_call=True #11967
• https://github.com/vllm-project/vllm/commits/main/vllm/vllm_flash_attn/flash_attn_interface.py

In v0, we split prefill and decode

• prefill uses flash_attn_varlen_func for compute-intensive tasks.
• decode uses flash_attn_with_kvcache, which can benefit from flash decoding.

vllm/vllm/attention/backends/flash_attn.py

Lines 723 to 820 in f49777b

	if prefill_meta := attn_metadata.prefill_metadata:
	# Prompt run.
	if (kv_cache.numel() == 0 or prefill_meta.block_tables is None
	or prefill_meta.block_tables.numel() == 0):
	# normal attention
	# When block_tables are not filled, it means q and k are the
	# prompt, and they have the same length.
	q_seq_start_loc, q_seq_len, k_seq_start_loc, k_seq_len = \
	_get_query_key_seq_metadata(prefill_meta, True, attn_type)
	key = key[:num_prefill_kv_tokens]
	value = value[:num_prefill_kv_tokens]
	flash_attn_varlen_func(
	q=query,
	k=key,
	v=value,
	cu_seqlens_q=q_seq_start_loc,
	cu_seqlens_k=k_seq_start_loc,
	max_seqlen_q=q_seq_len,
	max_seqlen_k=k_seq_len,
	softmax_scale=softmax_scale,
	causal=_get_causal_option(attn_type),
	window_size=window_size,
	alibi_slopes=alibi_slopes,
	softcap=logits_soft_cap,
	out=prefill_output,
	)
	else:
	# prefix-enabled attention
	assert attn_type == AttentionType.DECODER, (
	"Only decoder-only models support prefix caching")
	assert prefill_meta.seq_lens is not None
	max_seq_len = max(prefill_meta.seq_lens)
	flash_attn_varlen_func( # noqa
	q=query,
	k=key_cache,
	v=value_cache,
	cu_seqlens_q=prefill_meta.query_start_loc,
	max_seqlen_q=prefill_meta.max_query_len,
	cu_seqlens_k=prefill_meta.seq_start_loc,
	max_seqlen_k=max_seq_len,
	softmax_scale=softmax_scale,
	causal=True,
	window_size=window_size,
	alibi_slopes=alibi_slopes,
	block_table=prefill_meta.block_tables,
	softcap=logits_soft_cap,
	out=prefill_output,
	)
	if decode_meta := attn_metadata.decode_metadata:
	# Decoding run.
	# Use flash_attn_varlen_func kernel for speculative decoding
	# because different queries might have different lengths.
	assert decode_meta.max_decode_query_len is not None
	# use only for actual varlen decoding
	if decode_meta.max_decode_query_len > 1:
	assert attn_type == AttentionType.DECODER, (
	"Only decoder-only models support max_decode_query_len > 1"
	)
	flash_attn_varlen_func(
	q=decode_query,
	k=key_cache,
	v=value_cache,
	cu_seqlens_q=decode_meta.query_start_loc,
	max_seqlen_q=decode_meta.max_decode_query_len,
	cu_seqlens_k=decode_meta.seq_start_loc,
	max_seqlen_k=decode_meta.max_decode_seq_len,
	softmax_scale=softmax_scale,
	causal=True,
	window_size=window_size,
	alibi_slopes=alibi_slopes,
	softcap=logits_soft_cap,
	block_table=decode_meta.block_tables,
	out=decode_output,
	)
	else:
	# Use flash_attn_with_kvcache for normal decoding.
	(
	seq_lens_arg,
	_,
	block_tables_arg,
	) = get_seq_len_block_table_args(decode_meta, False, attn_type)
	flash_attn_with_kvcache(
	q=decode_query.unsqueeze(1),
	k_cache=key_cache,
	v_cache=value_cache,
	block_table=block_tables_arg,
	cache_seqlens=seq_lens_arg,
	softmax_scale=softmax_scale,
	causal=True,
	window_size=window_size,
	alibi_slopes=alibi_slopes,
	softcap=logits_soft_cap,
	out=decode_output.unsqueeze(1),
	)

This yields a 10% improvement on some hardware and certain batch sizes.

Perhaps we can reconsider now: Split prefill and decode attention vs unified_attention.

PTAL: #42826

[Alex-ai-future]

1. Does the VLLM community guarantee that the implementation is optimal for a given hardware, and that other VLLM-using companies have no alternative channels, at the model, op, and kernel levels?

2. Some stable ops maintain their own internal dispatch logic. However, can it be guaranteed that this dispatch is reasonable under all hardware and network conditions? Without OOT channels (custom ops or IRs), it's difficult to verify in a private environment. Can external systems be given some autonomy?

3. For new, flexible, performance-bottleneck ops, are dispatches placed within the model (along with the integration logic)? If so, must the model be reimplemented for optimization and verification in a private environment?

4. If each submodule can independently manage its own persistent information, wouldn't this benefit users in maintaining their own models?

[aditew01]

Is torch_compile_models supposed to be the legacy support for CPU backend?

[ProExpertProg]

@noooop we're actually on a different spot on the graphic in this situation, almost all torch.compile generated single-op kernels are faster than our in-tree cuda or triton kernels. This RFC doesn't remove torch.compile for the kernels. This is about the model-level compilation, which allows automatic fusion across abstractions.

Also, splitting prefill-decode paths is orthogonal to torch.compile, and already done today by FlashInfer - it's a cudagraph restriction not torch.compile restriction, and we already use piecewise cudagraphs for mixed batches even for FA so the change should be trivial.

[noooop]

@noooop we're actually on a different spot on the graphic in this situation, almost all torch.compile generated single-op kernels are faster than our in-tree cuda or triton kernels. This RFC doesn't remove torch.compile for the kernels. This is about the model-level compilation, which allows automatic fusion across abstractions.

Yes, this meme doesn’t seem to capture the subtle semantics of RFC... Because things are not that simple.

[ZJY0516]

almost all torch.compile generated single-op kernels are faster than our in-tree cuda or triton kernels

@ProExpertProg Any idea if removing full graph compile and dynamic dim hints will hurt kernel performance? And what if we manually add hints instead?

[PatrykWo]

@WoosukKwon
"One shared model + CustomOp dispatch" is exactly what forces us to monkey-patch but 2 things worry me though. Writing this from perspective of plugin -> HPU/Gaudi. Please don't make OOT plugins fork every upstream model. Today vllm-gaudi ships 19 models in-plugin and relies on CustomOp.register_oot / PluggableLayer.register_oot as the only seam to inject HPU kernels into the upstream shared model files. If CustomOp is removed before there's a replacement seam for OOT platforms, our only option is to fork all upstream models we care about. That's a maintenance cliff for us and a fragmentation cost for users.
Concretely, 26 register_oot sites cover: RMSNorm, 9 RoPE variants, MoE methods, compressed-tensors FP8/INT8 + 2 MoE quant methods, Mamba mixer, Conv2d/3d, RowParallelLinear, plus MLA wrapper and MM encoder attention via PluggableLayer. None of these live in our model files — they hook into upstream layers.
A Few asks:

1. Per-vendor model dirs: section 2 lists nvidia/ cpu/ /torch_compile_models/. Where do HPU live? Is the intent that OOT plugins contribute models/<model>/<vendor>/ via ModelRegistry entry-points? If yes, great — but please make that path part of the RFC, not implied. If no, we end up forking torch_compile_models/ per vendor, which defeats the goal.
2. Replace CustomOp with something for OOT, even minimal. Even if CustomOp goes away in-tree, OOT platforms still need a way to swap an upstream layer class without rewriting the model. PluggableLayer is already 90% of that mechanism and is much smaller than CustomOp — could it stay?
3. Quantization methods. Our heaviest path (HPU FP8/INT8) goes through LinearMethodBase / FusedMoEMethodBase, not directly from model.forward. The add_rms_norm_quant(...) example works for things called from model code, but not for layer-internal quant dispatch. What replaces register_oot on the method-base classes?

Could we please get: per-vendor layout + OOT extension story landed first, then a deprecation window (one minor release with @deprecated warnings) onregister_oot, then removal? A hard break in a single release forces every OOT plugin into a flag-day rewrite.

[ProExpertProg]

@ProExpertProg Any idea if removing full graph compile and dynamic dim hints will hurt kernel performance? And what if we manually add hints instead?

It's possible, we'll need to manually add hints which shouldn't be hard & should recover the perf (for unfused ops)

[yiz-liu]

I agree that some optimizations can be made explicit in the model code or kernels, but I also think that other optimizations are more naturally handled via compilation, and the backend’s reliance on compilation may vary.

For example, on Ascend, one practical benefit of compilation is that certain scalar or logical operations can be resolved ahead of execution, which improves performance. This doesn’t mean that model-level compilation is always the right approach, but it does suggest that the value of a compile path can be backend- and workload-dependent.

I’m curious whether other backends have similar cases, where compilation provides irreplaceable benefits, even if they appear through different mechanisms. If so, it could be useful to explore ways to let backends decide for themselves whether to use compilation?

[tjtanaa]

I think the hardware vendor specific implementation is a good way to go as it is harder to achieve aggressive of fusion through compilation pass as there could be infinitely many patterns; and it does reduce the friction between delivering optimization for different hardware vendors.

However, like @ProExpertProg mentioned, both could be orthogonal. For head models we go for hardware specific optimization and at the same time we should allow support of tail models that could be easily support on multiple platforms to help them to get wider adoption. For tail models we can still help them to optimize the model through fusion passes.

I think we should keep some part of IR Ops as it allows users and hardware vendors to easily perform parameter sweep without modifying the code. Users just need to change the implementation though, for example, --kernel-config.ir_op_priority.rms_norm=vllm_c, --kernel-config.ir_op_priority.rms_norm=tilelang, to find the best kernel implementation (cuda, cutlass, tilelang etc) for their hardware.

@WoosukKwon "One shared model + CustomOp dispatch" is exactly what forces us to monkey-patch but 2 things worry me though. Writing this from perspective of plugin -> HPU/Gaudi. Please don't make OOT plugins fork every upstream model. Today vllm-gaudi ships 19 models in-plugin and relies on CustomOp.register_oot / PluggableLayer.register_oot as the only seam to inject HPU kernels into the upstream shared model files. If CustomOp is removed before there's a replacement seam for OOT platforms, our only option is to fork all upstream models we care about. That's a maintenance cliff for us and a fragmentation cost for users. Concretely, 26 register_oot sites cover: RMSNorm, 9 RoPE variants, MoE methods, compressed-tensors FP8/INT8 + 2 MoE quant methods, Mamba mixer, Conv2d/3d, RowParallelLinear, plus MLA wrapper and MM encoder attention via PluggableLayer. None of these live in our model files — they hook into upstream layers. A Few asks:

1. Per-vendor model dirs: section 2 lists nvidia/ cpu/ /torch_compile_models/. Where do HPU live? Is the intent that OOT plugins contribute models/<model>/<vendor>/ via ModelRegistry entry-points? If yes, great — but please make that path part of the RFC, not implied. If no, we end up forking torch_compile_models/ per vendor, which defeats the goal.
2. Replace CustomOp with something for OOT, even minimal. Even if CustomOp goes away in-tree, OOT platforms still need a way to swap an upstream layer class without rewriting the model. PluggableLayer is already 90% of that mechanism and is much smaller than CustomOp — could it stay?
3. Quantization methods. Our heaviest path (HPU FP8/INT8) goes through LinearMethodBase / FusedMoEMethodBase, not directly from model.forward. The add_rms_norm_quant(...) example works for things called from model code, but not for layer-internal quant dispatch. What replaces register_oot on the method-base classes?

Could we please get: per-vendor layout + OOT extension story landed first, then a deprecation window (one minor release with @deprecated warnings) onregister_oot, then removal? A hard break in a single release forces every OOT plugin into a flag-day rewrite.

@PatrykWo indeed mention a very important point how will OOT be supported. Would be interested to know more about the approach to support this.