Repository navigation
[None][feat] Kimi K3 on modeling_v2: MXFP4 target, MNNVL and collective decode kernels, DSpark drafter #19991
+83,148
−408
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -336,6 +336,7 @@ struct MnnvlAllReduceKernelParams | |
| uint32_t* bufferFlags; | ||
| bool waitForResults; | ||
| QuantizationSFLayout layout; | ||
| bool earlyTrigger; | ||
| }; | ||
|
|
||
| template <typename PackedType, typename T> | ||
|
|
@@ -590,6 +591,10 @@ __global__ void __launch_bounds__(1024) oneshotAllreduceFusionKernel(MnnvlAllRed | |
| int threadOffset = token * params.tokenDim + packedIdx * kELTS_PER_THREAD; | ||
|
|
||
| cudaGridDependencySynchronize(); | ||
| if (params.earlyTrigger) | ||
|
Collaborator
There was a problem hiding this comment. Even though the changes to these communication kernels only add a single parameter, they could potentially impact all existing models that rely on them. My recommendation is not to make incremental changes to this shared kernel in this PR, but instead copy it over to kimi_k3_vl's own directory to modify it there.
Collaborator
There was a problem hiding this comment. I agree with the principle in general. For this specific case, I think the existing allreduce kernel shall already supported the PDL, I am not sure why do we need this new parameter? Can existing PDL infra support this? @vsabavat |
||
| { | ||
| cudaTriggerProgrammaticLaunchCompletion(); | ||
| } | ||
| #else | ||
| int packedIdx = blockIdx.y * blockDim.x + threadIdx.x; | ||
| int token = blockIdx.x; | ||
|
|
@@ -843,7 +848,7 @@ void oneshotAllreduceFusionOp(AllReduceFusionParams const& params) | |
| MnnvlAllReduceKernelParams<T> kernelParams{output, residualOut, input, residualIn, gamma, ucPtrs, | ||
| reinterpret_cast<T*>(params.bufferPtrLocal), mcPtr, params.quantOut, params.scaleOut, params.scaleFactor, | ||
| numTokens, tokenDim, params.nRanks, params.rank, static_cast<float>(params.epsilon), params.bufferFlags, | ||
| false, params.layout, params.earlyTrigger}; | ||
|
|
||
| switch (params.nRanks) | ||
| { | ||
|
|
||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,165 @@ | ||
| /* | ||
| * Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved. | ||
| * | ||
| * Licensed under the Apache License, Version 2.0 (the "License"); | ||
| * you may not use this file except in compliance with the License. | ||
| * You may obtain a copy of the License at | ||
| * | ||
| * http://www.apache.org/licenses/LICENSE-2.0 | ||
| * | ||
| * Unless required by applicable law or agreed to in writing, software | ||
| * distributed under the License is distributed on an "AS IS" BASIS, | ||
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| * See the License for the specific language governing permissions and | ||
| * limitations under the License. | ||
| */ | ||
|
|
||
| #include "tensorrt_llm/kernels/kimiK3Mnnvl/mnnvlAllGatherKernels.h" | ||
|
|
||
| #include "tensorrt_llm/common/cudaUtils.h" | ||
| #include "tensorrt_llm/common/envUtils.h" | ||
| #include "tensorrt_llm/common/lamportUtils.cuh" | ||
|
|
||
| TRTLLM_NAMESPACE_BEGIN | ||
|
|
||
| namespace kernels::kimiK3Mnnvl | ||
| { | ||
|
|
||
| using tensorrt_llm::common::isLamportDirty; | ||
| using tensorrt_llm::common::LamportFlags; | ||
| using tensorrt_llm::common::loadPackedVolatile; | ||
| using tensorrt_llm::common::VolatilePackedLoad; | ||
|
|
||
| namespace | ||
| { | ||
|
|
||
| constexpr int kThreads = 256; | ||
|
|
||
| // Payload vectors (16 bytes) of one rank's slot for one token: the bf16 part, then the fp32 part. | ||
| __host__ __device__ inline int slotVectors(int bf16Columns, int fp32Columns) | ||
| { | ||
| return bf16Columns / 8 + fp32Columns / 4; | ||
| } | ||
|
|
||
| // No payload word may equal the Lamport sentinel (the fp32 -0.0 word): -0.0 halves of a bf16 | ||
| // pair and -0.0 fp32 words become +0.0. | ||
| __device__ inline uint32_t sanitizeBf16Pair(uint32_t word) | ||
| { | ||
| if ((word & 0xffffu) == 0x8000u) | ||
| { | ||
| word &= 0xffff0000u; | ||
| } | ||
| if ((word >> 16) == 0x8000u) | ||
| { | ||
| word &= 0x0000ffffu; | ||
| } | ||
| return word; | ||
| } | ||
|
|
||
| __device__ inline uint32_t sanitizeFp32(uint32_t word) | ||
| { | ||
| return word == 0x80000000u ? 0u : word; | ||
| } | ||
|
|
||
| __device__ inline uint32_t packBf16Pair(float lo, float hi) | ||
| { | ||
| __nv_bfloat162 const pair = __floats2bfloat162_rn(lo, hi); | ||
| return sanitizeBf16Pair(*reinterpret_cast<uint32_t const*>(&pair)); | ||
| } | ||
|
|
||
| // One CTA per token: broadcast this rank's slot of the token, then poll every rank's slot and | ||
| // scatter it to the outputs. | ||
| __global__ void __launch_bounds__(kThreads) mnnvlAllGatherSplitKernel(AllGatherSplitParams const p) | ||
| { | ||
| #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900) | ||
| cudaGridDependencySynchronize(); | ||
| // Dependents read the outputs (and the Lamport flags) only after their own wait for this grid. | ||
| cudaTriggerProgrammaticLaunchCompletion(); | ||
| #endif | ||
| int const token = blockIdx.x; | ||
| int const bf16Vectors = p.bf16Columns / 8; | ||
| int const vectors = slotVectors(p.bf16Columns, p.fp32Columns); | ||
| int const inputColumns = p.bf16Columns + p.fp32Columns; | ||
|
|
||
| LamportFlags<float4> flag(p.bufferFlags, 1); | ||
| auto* lamportMcast = reinterpret_cast<uint4*>(flag.getCurLamportBuf(p.multicastPtr, 0)); | ||
| auto* lamportLocal = reinterpret_cast<uint4*>(flag.getCurLamportBuf(p.bufferPtrsDev[p.rank], 0)); | ||
|
|
||
| float const* row = p.input + static_cast<int64_t>(token) * inputColumns; | ||
| for (int v = threadIdx.x; v < vectors; v += blockDim.x) | ||
| { | ||
| uint4 packed; | ||
| if (v < bf16Vectors) | ||
| { | ||
| float4 const a = reinterpret_cast<float4 const*>(row)[2 * v]; | ||
| float4 const b = reinterpret_cast<float4 const*>(row)[2 * v + 1]; | ||
| packed = make_uint4( | ||
| packBf16Pair(a.x, a.y), packBf16Pair(a.z, a.w), packBf16Pair(b.x, b.y), packBf16Pair(b.z, b.w)); | ||
| } | ||
| else | ||
| { | ||
| uint4 const words = reinterpret_cast<uint4 const*>(row + p.bf16Columns)[v - bf16Vectors]; | ||
| packed = make_uint4( | ||
| sanitizeFp32(words.x), sanitizeFp32(words.y), sanitizeFp32(words.z), sanitizeFp32(words.w)); | ||
| } | ||
| lamportMcast[(static_cast<int64_t>(token) * p.nRanks + p.rank) * vectors + v] = packed; | ||
| } | ||
|
|
||
| flag.ctaArrive(); | ||
| flag.clearDirtyLamportBuf(p.bufferPtrsDev[p.rank], -1); | ||
|
|
||
| for (int i = threadIdx.x; i < p.nRanks * vectors; i += blockDim.x) | ||
| { | ||
| int const r = i / vectors; | ||
| int const v = i % vectors; | ||
| VolatilePackedLoad<float4> value; | ||
| do | ||
| { | ||
| value | ||
| = loadPackedVolatile<float4>(&lamportLocal[(static_cast<int64_t>(token) * p.nRanks + r) * vectors + v]); | ||
| } while (isLamportDirty(value)); | ||
| uint4 const words = make_uint4(value.words[0], value.words[1], value.words[2], value.words[3]); | ||
| if (v < bf16Vectors) | ||
| { | ||
| reinterpret_cast<uint4*>( | ||
| p.bf16Output + static_cast<int64_t>(token) * p.nRanks * p.bf16Columns + r * p.bf16Columns)[v] | ||
| = words; | ||
| } | ||
| else | ||
| { | ||
| reinterpret_cast<uint4*>(p.fp32Output + static_cast<int64_t>(token) * p.nRanks * p.fp32Columns | ||
| + r * p.fp32Columns)[v - bf16Vectors] | ||
| = words; | ||
| } | ||
| } | ||
|
|
||
| flag.waitAndUpdate({static_cast<uint32_t>(p.numTokens * p.nRanks * vectors * sizeof(uint4)), 0, 0, 0}); | ||
| } | ||
|
|
||
| } // namespace | ||
|
|
||
| int64_t mnnvlAllGatherSplitFootprint(int numTokens, int bf16Columns, int fp32Columns, int nRanks) | ||
| { | ||
| return static_cast<int64_t>(numTokens) * nRanks * slotVectors(bf16Columns, fp32Columns) * sizeof(uint4); | ||
| } | ||
|
|
||
| void mnnvlAllGatherSplitOp(AllGatherSplitParams const& params) | ||
| { | ||
| TLLM_CHECK_WITH_INFO(params.bf16Columns % 8 == 0 && params.fp32Columns % 4 == 0, | ||
| "[mnnvlAllGatherSplit] needs bf16 columns in multiples of 8 and fp32 columns in multiples of 4"); | ||
| TLLM_CHECK_WITH_INFO(params.numTokens > 0, "[mnnvlAllGatherSplit] needs at least one token"); | ||
| cudaLaunchAttribute attrs[1]; | ||
| attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization; | ||
| attrs[0].val.programmaticStreamSerializationAllowed = tensorrt_llm::common::getEnvEnablePDL() ? 1 : 0; | ||
| cudaLaunchConfig_t config{}; | ||
| config.gridDim = dim3(params.numTokens); | ||
| config.blockDim = dim3(kThreads); | ||
| config.stream = params.stream; | ||
| config.attrs = attrs; | ||
| config.numAttrs = 1; | ||
| TLLM_CUDA_CHECK(cudaLaunchKernelEx(&config, mnnvlAllGatherSplitKernel, params)); | ||
| } | ||
|
|
||
| } // namespace kernels::kimiK3Mnnvl | ||
|
|
||
| TRTLLM_NAMESPACE_END |
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Uh oh!
There was an error while loading. Please reload this page.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
For the new kernels, please move them to the kimi directory under modeling_v2 (you can add a csrc/ folder for the C++/CUDA sources). Shared kernels for TP16/TP4 belong in models/kimi_k3_vl, while target-specific kernels should reside in their respective target folders. In short, new ops should be scoped to the level where they take effect rather than dumped into generic modules.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
The K3 C++ kernels are already isolated under
cpp/tensorrt_llm/kernels/kimiK3Mnnvl/; the remaining sharedearlyTriggeroption is default-off. Moving the sources undermodels/kimi_k3_vl/csrc/would also need explicit build wiring, since the current kernel build discovers them undercpp/.Do you want that model-local C++/CUDA build-layout change as well, or is keeping the K3-specific compiled kernels in
cpp/acceptable while moving the Python ops and CuTe DSL kernels into the K3 family/target directories under modeling_v2?