Skip to content

Commit da333b2

Browse files
committed
[None][refactor] Dispatch both Sage smoothing flags in one place
launchWithSmoothK took the V-smoothing flag from its caller, which dispatched it per stage. launchWithSmoothKV dispatches both flags for a stage and instantiates only the combinations that exist: K smoothing needs V staging, and V smoothing only changes the statistics launch. Signed-off-by: yuhangh <58161490+heyuhhh@users.noreply.github.com>
1 parent 6069637 commit da333b2

1 file changed

Lines changed: 25 additions & 19 deletions

File tree

‎cpp/tensorrt_llm/common/sageQuant.cu‎

Lines changed: 25 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -726,33 +726,39 @@ void invokeSageQuantQkvImpl(SageQuantParams const& params)
726726
TLLM_CHECK_WITH_INFO(status == cudaSuccess, "%s", cudaGetErrorString(status));
727727
};
728728

729-
auto launchWithSmoothK = [&](auto smoothVStatic, auto vStageStatic)
729+
// K smoothing needs V staging, and V smoothing only changes what the VStage==1 launch
730+
// collects, so only those combinations are instantiated.
731+
auto launchWithSmoothKV = [&](auto vStageStatic)
730732
{
731-
if (params.kSmooth)
733+
constexpr int VStage_ = vStageStatic;
734+
auto launchWithSmoothV = [&](auto smoothKStatic)
732735
{
733-
launchKernel(cute::true_type{}, smoothVStatic, vStageStatic);
734-
}
735-
else
736+
if constexpr (VStage_ == 1)
737+
{
738+
if (params.ptrVMean != nullptr)
739+
{
740+
launchKernel(smoothKStatic, cute::true_type{}, vStageStatic);
741+
return;
742+
}
743+
}
744+
launchKernel(smoothKStatic, cute::false_type{}, vStageStatic);
745+
};
746+
if constexpr (VStage_ != 0)
736747
{
737-
launchKernel(cute::false_type{}, smoothVStatic, vStageStatic);
748+
if (params.kSmooth)
749+
{
750+
launchWithSmoothV(cute::true_type{});
751+
return;
752+
}
738753
}
754+
launchWithSmoothV(cute::false_type{});
739755
};
740756

741-
// V smoothing only changes what the VStage==1 launch collects.
742757
switch (params.vStage)
743758
{
744-
case 0: launchKernel(cute::false_type{}, cute::false_type{}, Int<0>{}); return;
745-
case 1:
746-
if (params.ptrVMean != nullptr)
747-
{
748-
launchWithSmoothK(cute::true_type{}, Int<1>{});
749-
}
750-
else
751-
{
752-
launchWithSmoothK(cute::false_type{}, Int<1>{});
753-
}
754-
return;
755-
case 2: launchWithSmoothK(cute::false_type{}, Int<2>{}); return;
759+
case 0: launchWithSmoothKV(Int<0>{}); return;
760+
case 1: launchWithSmoothKV(Int<1>{}); return;
761+
case 2: launchWithSmoothKV(Int<2>{}); return;
756762
default: TLLM_THROW("Unsupported SageQuantV stage %d", params.vStage);
757763
}
758764
};

0 commit comments

Comments
 (0)