kExl3MoeMlp on Vulkan is owed and unowned: its group barrier needs a grid sync Vulkan has at no version, and its tiles need 28-38 KiB against a 16 KiB guarantee

Row: BACKEND-VULKAN-EXL3

.agents/specs/backend-vulkan-exl3.md ## Owed names this and says "No issue owns it yet". This issue is that owner. It adds the numbers, because the reason is structural and the numbers are what make it checkable rather than a judgement.

kExl3MoeMlp is registered for CPU and CUDA only. On a Vulkan queue an EXL3 MoE runs on the portable CPU reference tier. kExl3Gemm and kCastF16 are native on Vulkan as of #2530 and decode every width the host does, so the fused MoE is the one remaining EXL3 op.

Why it is a rewrite, not a transcription

vt_exl3_gemm.comp could be transcribed from the portable CPU reference because a GEMM is embarrassingly parallel over output elements. The fused MoE is not that shape.

1. It needs a grid-wide barrier, and Vulkan has none at any version. exl3_moe_kernel is a persistent cooperative launch: num_groups groups of group_size blocks draw expert tickets from a self-resetting scheduler in the lock buffer and meet at group_barrier(barrier_cs, group_idx, group_size) between each of the five stages. That spin barrier is legal only because cudaLaunchCooperativeKernel guarantees every block of the grid is co-resident; a spin barrier over blocks that are not co-resident deadlocks. Vulkan's vkCmdDispatch makes no co-residency guarantee, has no cross-workgroup barrier, and nothing in VK_KHR_cooperative_matrix or the VK_NV_cooperative_* family supplies one. The port is therefore five dispatches with pipeline barriers between them, which discards the persistent scheduler and the temp-buffer reuse the fused arm exists for — a different algorithm with a different cost profile, not the same one in another language.

2. Shared memory. The launcher requests kSmemMax = 92160 bytes of dynamic shared memory. Vulkan's guaranteed maxComputeSharedMemorySize is 16384 bytes, and real implementations commonly expose 32768 or 49152. What the MoE shape actually needs, from the static_assert in exl3_gemm_kernel_inner at TILESIZE_M 16, TILESIZE_K 32, MOE_TILESIZE_N 256, SH_STAGES 3:

So even the narrowest arm is over the guarantee, and the two widest are over a common 32 KiB limit. This is the second blocker and not the first: a smaller tile could be chosen, whereas no tile choice creates a grid barrier.

3. The same four blockers rocm_exl3.hip records for the GEMM apply here too, because the MoE bands call exl3_gemm_kernel_inner: mma.sync.aligned.m16n8k16, ldmatrix.sync.aligned.m8n8.x4.shared.b16, cp.async with its commit/wait groups, and the 90 KiB budget. On Vulkan the donor would again be the portable CPU reference, as it was for vt_exl3_gemm.comp.

What would make it worth doing

Nothing routes to it yet. vt::Exl3MoeMlp has two production callers, both in Exl3FusedMoePass (deepseek_v4.cpp), and DeepSeek-V4 EXL3 is the only architecture that reaches the op. A Vulkan arm today would be a registered capability no production path can drive, which is the dead-code case. The trigger is a second EXL3 MoE architecture, or a DeepSeek-V4 EXL3 artifact small enough for a Vulkan device.

Related

• #2756 widened the CUDA arm from one instantiation to bits 3..6 at codebook 1, and records the codebook half as a loader slice.
• #2762 is why that arm could be gated at all: the device case had guarded on a predicate a static_assert pins false everywhere.
• kExl3MoeMlp on ROCm is owed on BACKEND-ROCM-EXL3 for a different reason — HIP does have hipLaunchCooperativeKernel, so the barrier is not the blocker there; the LDS budget, the matrix-core shapes and the absence of any AMD board on this fleet that holds the reaching artifact are.

Start with .agents/specs/backend-vulkan-exl3.md and inspect vt_exl3_gemm.comp, rocm_exl3.hip, and the Exl3FusedMoePass callers in deepseek_v4.cpp. Confirm whether a triggering EXL3 MoE architecture or suitable artifact exists, then define a Vulkan approach whose barriers and shared-memory use fit the device limits; done means the Vulkan arm is registered and exercised by a production path.