feat: pack 2 fc layers in mlp with swiglu - #218
Open
Chamberlain0w0 wants to merge 3 commits into
Open
Chamberlain0w0 wants to merge 3 commits into
Chamberlain0w0 wants to merge 3 commits into
Conversation
JYMiracle305
reviewed
Sep 8, 2026
JYMiracle305
reviewed
Sep 8, 2026
JYMiracle305
approved these changes
Sep 15, 2026
kilinchange
requested changes
Sep 17, 2026
| } | ||
| } | ||
|
|
||
| inline size_t ChooseBlockSize(size_t num_elements) { |
Collaborator
There was a problem hiding this comment.
这个抽到 infini_train/include/common/cuda/common_cuda.h 里吧,elementwise.cu 和 swiglu.cu 统一调同一个公共函数。
| using namespace infini_train::common::cuda; | ||
|
|
||
| template <typename T> | ||
| __global__ void SwiGLUForwardKernel(T *__restrict__ output, const T *__restrict__ input, int64_t hidden, |
Collaborator
There was a problem hiding this comment.
留个 TODO,后续考虑新增向量化访存优化。
| } | ||
|
|
||
| // For packed SwiGLU: FC1 projection is stored locally as [gate_i | up_i] on each TP rank. | ||
| for (const auto &[module_name, module] : named_modules) { |
Collaborator
There was a problem hiding this comment.
这部分代码和上面 QKV 的处理逻辑上相似度比较高,能不能抽出通用逻辑,合并成一个循环:
for (const auto &[module_name, module] : named_modules) {
if (dynamic_cast<CausalSelfAttention *>(module.get())) {
MarkPackedQKVLoRASharding(module_name, module, shardings);
} else if (dynamic_cast<MLP *>(module.get())) {
MarkPackedSwiGLULoRASharding(module_name, module, shardings);
}
}
给指定 projection 的 LoRA-B 设置 sharding 的通用 helper:
void MarkLoRAColumnParallelBSharding(
const std::string &module_name,
const std::shared_ptr<Module> &module,
const std::string &projection_name,
LoRATensorSharding sharding,
std::unordered_map<std::string, LoRATensorSharding> &shardings) {
auto projection = module->mutable_module(projection_name);
if (!dynamic_cast<LoRAColumnParallelLinear *>(projection.get())) {
return;
}
const auto projection_module_name =
QualifiedParamName(module_name, projection_name);
shardings[QualifiedParamName(
projection_module_name,
LoRAColumnParallelLinear::kParamLoraBName)] = sharding;
}
QKV helper:
void MarkPackedQKVLoRASharding(
const std::string &module_name,
const std::shared_ptr<Module> &module,
std::unordered_map<std::string, LoRATensorSharding> &shardings) {
MarkLoRAColumnParallelBSharding(
module_name,
module,
CausalSelfAttention::kCAttnLayerName,
LoRATensorSharding::kPackedQKVColumnParallelDim0,
shardings);
}
SwiGLU helper:
void MarkPackedSwiGLULoRASharding(
const std::string &module_name,
const std::shared_ptr<Module> &module,
std::unordered_map<std::string, LoRATensorSharding> &shardings) {
bool is_swiglu = false;
for (const auto &child : module->modules()) {
if (dynamic_cast<SwiGLU *>(child.get())) {
is_swiglu = true;
break;
}
}
if (!is_swiglu) {
return;
}
MarkLoRAColumnParallelBSharding(
module_name,
module,
MLP::kFcLayerName,
LoRATensorSharding::kPackedSwiGLUColumnParallelDim0,
shardings);
}
| return nn::function::Concat(reordered_shards, 0); | ||
| } | ||
|
|
||
| std::shared_ptr<Tensor> SlicePackedSwiGLURowsForTensorParallel(const std::shared_ptr<Tensor> &full_tensor, int tp_rank, |
Collaborator
There was a problem hiding this comment.
这两个 slice/restore 函数跟 QKV 的处理逻辑也比较类似,可以看下是否方便抽出通用逻辑,避免后续 fused projection 越来越多,都各自写一套重复的逻辑。
| } | ||
|
|
||
| TEST_P(AutogradElementwiseBackwardTest, SwiGLUAutocastBackward) { | ||
| ONLY_CUDA(); |
Collaborator
There was a problem hiding this comment.
这里可能需要适配下 #206 的改动,不能再在测试里使用 ONLY_CUDA 了,应该使用 SKIP_CPU 跳过 cpu 执行,麻烦 @chen2021673 确认下。
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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
背景
当前 SwiGLU MLP 使用两个独立的
ColumnParallelLinear分别计算 gate projection 和 up projection,随后执行 SiLU 和逐元素乘法:该实现会产生两次独立 GEMM 调用,并分别保存中间结果。对于 MoE,大 batch 下各 expert 的中间激活会带来较明显的性能和显存开销。Megatron 中常将两个 FC1 投影打包为单个 GEMM,本 PR 也做了类似的修改,并新增融合的 SwiGLU kernel。
主要改动
Packed FC1
将 SwiGLU 的 gate/up projection 合并为一个输出维度为
2H的ColumnParallelLinear:packed tensor 采用 Megatron-LM 的
[gate, up]布局,便于后续对齐 Megatron checkpoint 和相关 fused kernel。Dense MLP 和 MoE expert 共用该实现,因此均切换为 packed SwiGLU。GELU MLP 路径保持不变。
SwiGLU kernel
新增独立的 SwiGLU kernel 文件:
infini_train/src/kernels/cpu/swiglu.ccinfini_train/src/kernels/cuda/swiglu.cukernel 融合执行 SiLU 和逐元素乘法,并实现对应反向传播:
支持范围:
本 PR 合并了两个 FC1 GEMM,并融合了
SiLU + multiply。Autograd 与 Module
[gate, up]的 packed tensor。[gate, up]顺序写回。Checkpoint loader
同步调整 LLMC checkpoint loader:
gate_proj写入 packed FC1 前半部分,将up_proj写入后半部分。w1/gate_proj写入前半部分,将w3/up_proj写入后半部分。[2H_local]布局分别加载 gate/up shard。参数布局变化
SwiGLU MLP 的 state dict 从:
调整为:
c_fc2.weight不再存在。并且将 Packed SwiGLU 设置为 SwiGLU MLP 的默认实现,不需要额外配置开关。
MoE 性能提升
统计时排除 step 1 warm-up,使用 step 2–10 的平均耗时。baseline 和 packed 使用相同启动参数及构建配置。
结果表明: