Skip to content

mega: wire block-scaled FP8 weights into the fused MoE slots #1

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
wseaton wants to merge 1 commit into
base: nv_dev
Choose a base branch
from
Draft

mega: wire block-scaled FP8 weights into the fused MoE slots #1

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
96 changes: 96 additions & 0 deletions csrc/apis/mega.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -226,11 +226,107 @@ static void fp8_fp4_mega_moe(
sym_buffer.zero_();
}

static void fp8_mega_moe(
const torch::Tensor& y,
const std::tuple<torch::Tensor, torch::Tensor>& l1_weights_tuple,
const std::tuple<torch::Tensor, torch::Tensor>& l2_weights_tuple,
const std::optional<torch::Tensor>& cumulative_local_expert_recv_stats,
const torch::Tensor& sym_buffer,
const std::vector<int64_t>& sym_buffer_ptrs, const int& rank_idx,
const int& num_max_tokens_per_rank,
const int& num_experts, const int& num_topk,
const std::tuple<int, int, int>& recipe,
const std::string& activation,
const std::optional<float>& activation_clamp_opt,
const bool& fast_math
) {
const auto [l1_weights, l1_weights_sf] = l1_weights_tuple;
const auto [l2_weights, l2_weights_sf] = l2_weights_tuple;

// Config checks
const auto num_tokens = static_cast<int>(y.size(0));
const auto [rm, rn, rk] = recipe;
DG_HOST_ASSERT(rm == 1 and rn == 1 and rk == 32);
DG_HOST_ASSERT(activation == "swiglu");

// Activation checks
const auto activation_clamp =
activation_clamp_opt.value_or(std::numeric_limits<float>::infinity());
DG_HOST_ASSERT(activation_clamp >= 0);

// FP8 weight checks: weights must be FP8 E4M3
DG_HOST_ASSERT(l1_weights.scalar_type() == torch::kFloat8_e4m3fn);
DG_HOST_ASSERT(l2_weights.scalar_type() == torch::kFloat8_e4m3fn);
DG_HOST_ASSERT(get_major_type_ab(l1_weights) == cute::UMMA::Major::K);
DG_HOST_ASSERT(get_major_type_ab(l2_weights) == cute::UMMA::Major::K);
const auto arch_major = device_runtime->get_arch_major();
const auto [num_experts_per_rank, intermediate_hidden_before_act, hidden] =
check_grouped_ab_fp8_fp4(l1_weights, cute::UMMA::Major::K, arch_major);
const auto [num_experts_per_rank_, hidden_, intermediate_hidden] =
check_grouped_ab_fp8_fp4(l2_weights, cute::UMMA::Major::K, arch_major);
DG_HOST_ASSERT(num_tokens <= num_max_tokens_per_rank);
DG_HOST_ASSERT(num_experts_per_rank == num_experts_per_rank_);
DG_HOST_ASSERT(hidden == hidden_);
DG_HOST_ASSERT(intermediate_hidden_before_act == 2 * intermediate_hidden);
DG_HOST_ASSERT(l1_weights.is_contiguous() and l2_weights.is_contiguous());

// Check weight SF layout for UE8M0 packing, MN-major, and TMA alignment
constexpr int kGranMN = 1, kGranK = 32;
check_sf_layout(l1_weights_sf, intermediate_hidden * 2, hidden, kGranMN, kGranK,
num_experts_per_rank, true, false, torch::kInt);
check_sf_layout(l2_weights_sf, hidden, intermediate_hidden, kGranMN, kGranK,
num_experts_per_rank, true, false, torch::kInt);

// Check stats counter
if (cumulative_local_expert_recv_stats.has_value()) {
DG_HOST_ASSERT(cumulative_local_expert_recv_stats->scalar_type() == torch::kInt);
DG_HOST_ASSERT(cumulative_local_expert_recv_stats->numel() == num_experts_per_rank);
DG_HOST_ASSERT(cumulative_local_expert_recv_stats->is_contiguous());
}

// Check buffer bytes
const auto num_ranks = static_cast<int>(sym_buffer_ptrs.size());
const auto num_experts_ = num_experts_per_rank * num_ranks;
const auto [num_required_bytes, slice] = get_symm_buffer_size_for_mega_moe(
num_ranks, num_experts,
num_max_tokens_per_rank, num_topk,
hidden, intermediate_hidden,
true, activation);
DG_HOST_ASSERT(sym_buffer.nbytes() >= static_cast<size_t>(num_required_bytes));
DG_HOST_ASSERT(num_experts == num_experts_);

// Already registered tensors
const auto [x, x_sf, topk_idx, topk_weights, l1_acts, l1_acts_sf, l2_acts, l2_acts_sf] = slice(sym_buffer);

// Dispatch into different architectures
if (arch_major == 10) {
sm100_fp8_mega_moe(y,
l1_acts, l1_acts_sf,
l2_acts, l2_acts_sf,
l1_weights, l2_weights,
l1_weights_sf, l2_weights_sf,
cumulative_local_expert_recv_stats,
sym_buffer_ptrs,
rank_idx, num_max_tokens_per_rank,
num_experts_per_rank,
num_tokens, num_topk,
hidden, intermediate_hidden,
activation_clamp, fast_math);
} else {
DG_HOST_UNREACHABLE("Unsupported architecture for FP8 MegaMoE");
}

// Zero the entire symmetric buffer for debug mode
if (get_env<int>("DG_COMM_KERNEL_DEBUG"))
sym_buffer.zero_();
}

static void register_apis(pybind11::module_& m) {
#if DG_TENSORMAP_COMPATIBLE
m.def("get_token_alignment_for_mega_moe", &get_token_alignment_for_mega_moe);
m.def("get_symm_buffer_size_for_mega_moe", &get_symm_buffer_size_for_mega_moe);
m.def("fp8_fp4_mega_moe", &fp8_fp4_mega_moe);
m.def("fp8_mega_moe", &fp8_mega_moe);
#endif
}

Expand Down
60 changes: 56 additions & 4 deletions csrc/jit_kernels/impls/sm100_fp8_fp4_mega_moe.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,7 @@ class SM100FP8FP4MegaMoERuntime final : public LaunchRuntime<SM100FP8FP4MegaMoER
int num_ranks;
float activation_clamp;
bool fast_math;
bool use_fp8_weights;
MegaMoEConfig config;

// Runtime arguments
Expand Down Expand Up @@ -70,6 +71,7 @@ static void __instantiate_kernel() {{
{}, {},
{},
{},
{},
{}
>);
}};
Expand All @@ -87,7 +89,8 @@ static void __instantiate_kernel() {{
args.config.num_dispatch_threads, args.config.num_non_epilogue_threads, args.config.num_epilogue_threads,
args.launch_args.grid_dim.first, args.num_ranks,
to_string(args.activation_clamp),
args.fast_math ? "true" : "false");
args.fast_math ? "true" : "false",
args.use_fp8_weights ? "true" : "false");
}

static void launch_impl(const KernelHandle& kernel, const LaunchConfigHandle& config, Args args) {
Expand All @@ -110,7 +113,7 @@ static void __instantiate_kernel() {{
}
};

static void sm100_fp8_fp4_mega_moe(
static void sm100_mega_moe_launch(
const torch::Tensor& y,
const torch::Tensor& l1_acts, const torch::Tensor& l1_acts_sf,
const torch::Tensor& l2_acts, const torch::Tensor& l2_acts_sf,
Expand All @@ -123,7 +126,8 @@ static void sm100_fp8_fp4_mega_moe(
const int& num_tokens, const int& num_topk,
const int& hidden, const int& intermediate_hidden,
const float& activation_clamp,
const bool& fast_math
const bool& fast_math,
const bool& use_fp8_weights
) {
const auto num_ranks = static_cast<int>(sym_buffer_ptrs.size());
const auto num_experts = num_experts_per_rank * num_ranks;
Expand Down Expand Up @@ -199,6 +203,7 @@ static void sm100_fp8_fp4_mega_moe(
.num_ranks = num_ranks,
.activation_clamp = activation_clamp,
.fast_math = fast_math,
.use_fp8_weights = use_fp8_weights,
.config = config,
.y = y.data_ptr(),
.cumulative_local_expert_recv_stats = cumulative_local_expert_recv_stats_ptr,
Expand All @@ -218,9 +223,56 @@ static void sm100_fp8_fp4_mega_moe(
config.smem_size, 2)
};

const auto kernel_name = use_fp8_weights ? "sm100_fp8_mega_moe" : "sm100_fp8_fp4_mega_moe";
const auto code = SM100FP8FP4MegaMoERuntime::generate(args);
const auto runtime = compiler->build("sm100_fp8_fp4_mega_moe", code);
const auto runtime = compiler->build(kernel_name, code);
SM100FP8FP4MegaMoERuntime::launch(runtime, args);
}

static void sm100_fp8_fp4_mega_moe(
const torch::Tensor& y,
const torch::Tensor& l1_acts, const torch::Tensor& l1_acts_sf,
const torch::Tensor& l2_acts, const torch::Tensor& l2_acts_sf,
const torch::Tensor& l1_weights, const torch::Tensor& l2_weights,
const torch::Tensor& l1_weights_sf, const torch::Tensor& l2_weights_sf,
const std::optional<torch::Tensor> cumulative_local_expert_recv_stats,
const std::vector<int64_t>& sym_buffer_ptrs,
const int& rank_idx, const int& num_max_tokens_per_rank,
const int& num_experts_per_rank,
const int& num_tokens, const int& num_topk,
const int& hidden, const int& intermediate_hidden,
const float& activation_clamp,
const bool& fast_math
) {
sm100_mega_moe_launch(y, l1_acts, l1_acts_sf, l2_acts, l2_acts_sf,
l1_weights, l2_weights, l1_weights_sf, l2_weights_sf,
cumulative_local_expert_recv_stats, sym_buffer_ptrs,
rank_idx, num_max_tokens_per_rank, num_experts_per_rank,
num_tokens, num_topk, hidden, intermediate_hidden,
activation_clamp, fast_math, false);
}

static void sm100_fp8_mega_moe(
const torch::Tensor& y,
const torch::Tensor& l1_acts, const torch::Tensor& l1_acts_sf,
const torch::Tensor& l2_acts, const torch::Tensor& l2_acts_sf,
const torch::Tensor& l1_weights, const torch::Tensor& l2_weights,
const torch::Tensor& l1_weights_sf, const torch::Tensor& l2_weights_sf,
const std::optional<torch::Tensor> cumulative_local_expert_recv_stats,
const std::vector<int64_t>& sym_buffer_ptrs,
const int& rank_idx, const int& num_max_tokens_per_rank,
const int& num_experts_per_rank,
const int& num_tokens, const int& num_topk,
const int& hidden, const int& intermediate_hidden,
const float& activation_clamp,
const bool& fast_math
) {
sm100_mega_moe_launch(y, l1_acts, l1_acts_sf, l2_acts, l2_acts_sf,
l1_weights, l2_weights, l1_weights_sf, l2_weights_sf,
cumulative_local_expert_recv_stats, sym_buffer_ptrs,
rank_idx, num_max_tokens_per_rank, num_experts_per_rank,
num_tokens, num_topk, hidden, intermediate_hidden,
activation_clamp, fast_math, true);
}

} // namespace deep_gemm
2 changes: 2 additions & 0 deletions deep_gemm/__init__.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -85,7 +85,9 @@
SymmBuffer,
get_symm_buffer_for_mega_moe,
transform_weights_for_mega_moe,
transform_weights_for_fp8_mega_moe,
fp8_fp4_mega_moe,
fp8_mega_moe,
)

# Some utils
Expand Down
5 changes: 3 additions & 2 deletions deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_mega_moe.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,6 +35,7 @@ template <
uint32_t kNumSMs, uint32_t kNumRanks,
float kActivationClamp,
bool kFastMath,
bool kUseFP8Weights = false,
uint32_t L1_SHAPE_N = kIntermediateHidden * 2,
uint32_t L1_SHAPE_K = kHidden,
uint32_t L2_SHAPE_N = kHidden,
Expand Down Expand Up @@ -160,9 +161,9 @@ sm100_fp8_fp4_mega_moe_impl(void* y,
);

// Data types
// NOTES: activations are FP8 (e4m3), weights are FP4 (e2m1)
// NOTES: activations are FP8 (e4m3), weights are FP4 (e2m1) or FP8 (e4m3)
using a_dtype_t = cutlass::float_e4m3_t;
using b_dtype_t = cutlass::detail::float_e2m1_unpacksmem_t;
using b_dtype_t = cute::conditional_t<kUseFP8Weights, cutlass::float_e4m3_t, cutlass::detail::float_e2m1_unpacksmem_t>;

// MMA configs
// NOTES: always swap A/B, 2-CTA MMA, and matrices are K-major
Expand Down
36 changes: 36 additions & 0 deletions deep_gemm/mega/__init__.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -108,6 +108,19 @@ def transform_weights_for_mega_moe(
return l1_transformed, l2_transformed


def transform_weights_for_fp8_mega_moe(
l1_weights: Tuple[torch.Tensor, torch.Tensor],
l2_weights: Tuple[torch.Tensor, torch.Tensor]
) -> Tuple[Tuple[torch.Tensor, torch.Tensor], Tuple[torch.Tensor, torch.Tensor]]:
# FP8 weights: same interleave gate/up layout as FP4, plus UTCCP SF transpose.
l1_w = _interleave_weights(l1_weights[0])
l1_sf = _transpose_sf_for_utccp(_interleave_weights(l1_weights[1]))
l1_transformed = (l1_w, l1_sf)
# L2: only transpose SF for UTCCP.
l2_transformed = (l2_weights[0], _transpose_sf_for_utccp(l2_weights[1]))
return l1_transformed, l2_transformed



def fp8_fp4_mega_moe(y: torch.Tensor,
l1_weights: Tuple[torch.Tensor, torch.Tensor],
Expand All @@ -130,3 +143,26 @@ def fp8_fp4_mega_moe(y: torch.Tensor,
activation, activation_clamp,
fast_math
)


def fp8_mega_moe(y: torch.Tensor,
l1_weights: Tuple[torch.Tensor, torch.Tensor],
l2_weights: Tuple[torch.Tensor, torch.Tensor],
sym_buffer: SymmBuffer,
cumulative_local_expert_recv_stats: Optional[torch.Tensor] = None,
recipe: Tuple[int, int, int] = (1, 1, 32),
activation: str = 'swiglu',
activation_clamp: Optional[float] = None,
fast_math: bool = True):
_C.fp8_mega_moe(
y,
l1_weights, l2_weights,
cumulative_local_expert_recv_stats,
sym_buffer.buffer,
sym_buffer.handle.buffer_ptrs, sym_buffer.group.rank(),
sym_buffer.num_max_tokens_per_rank,
sym_buffer.num_experts, sym_buffer.num_topk,
recipe,
activation, activation_clamp,
fast_math
)
Loading