Skip to content

Commit 825bd30

Browse files
[MLX] Qwen3.5 MoE ondevice sampling (pytorch#20545)
Summary Lets the MLX-exported Qwen3.5 MoE model sample the next token on-device instead of returning logits for host-side sampling. Sampling is opt-in at export (--sample); temperature, top_p, and seed are runtime inputs, and the runner increments the seed per token. Measured ~17% higher decode throughput. Changes - export.py --sample flag wraps the model in SamplingHead so forward(tokens, input_pos, temperature, top_p, seed) → int64 token; records a use_sampling constant-method flag. Non-sample export unchanged. - qwen35_moe_engine.cpp reads use_sampling from metadata; when set, consumes the token id directly instead of logits_to_token, feeds the scalar inputs in prefill/decode, and manages the per-token seed schedule. top_k is still rejected; top_p/seed are rejected on non-sample models. - main.cpp --top_p / --seed flags wired into SamplingConfig. - run.py detects use_sampling, mirrors the same seed schedule for parity, and rejects top_p/seed on non-sample models (matching the C++ runner). --top-p / --seed flags added. Verification - Real model (C++ runner): coherent, constraint-following generation; same --seed reproducible; --temperature 0 deterministic. - +17% decode throughput vs host sampling; prefill unchanged.
1 parent 51729bb commit 825bd30

4 files changed

Lines changed: 270 additions & 30 deletions

File tree

‎examples/models/qwen3_5_moe/export.py‎

Lines changed: 49 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -733,6 +733,25 @@ def _clean_forward(self, tokens, input_pos):
733733
model.forward = types.MethodType(_clean_forward, model)
734734

735735

736+
class _MLXSampleWrapper(nn.Module):
737+
"""Wrap the logits-producing model so ``forward`` returns a sampled token.
738+
739+
Temperature, top_p, and seed are runtime scalar inputs so the same .pte
740+
serves any sampling request; the runner increments the seed per token.
741+
"""
742+
743+
def __init__(self, model: nn.Module):
744+
super().__init__()
745+
from executorch.backends.mlx.llm.sampling import SamplingHead
746+
747+
self.head = SamplingHead(model)
748+
749+
def forward(self, tokens, input_pos, temperature, top_p, seed):
750+
return self.head(
751+
tokens, input_pos, temperature=temperature, top_p=top_p, seed=seed
752+
)
753+
754+
736755
def _export_mlx(model, config, args):
737756
"""Export model to .pte via torch.export + MLX backend."""
738757
import gc
@@ -749,16 +768,33 @@ def _export_mlx(model, config, args):
749768

750769
_strip_sampler_from_forward(model)
751770

771+
sample = getattr(args, "sample", False)
772+
752773
example_tokens = torch.tensor([[0, 1]], dtype=torch.long)
753774
example_input_pos = torch.tensor([0, 1], dtype=torch.long)
754775
seq_dim = Dim("seq_len", min=1, max=config.max_seq_len - 1)
755-
dynamic_shapes = ({1: seq_dim}, {0: seq_dim})
776+
777+
if sample:
778+
# forward(tokens, input_pos, temperature, top_p, seed) -> token id.
779+
# Scalars are static (None in dynamic_shapes); only the seq dim is dynamic.
780+
model = _MLXSampleWrapper(model)
781+
example_args = (
782+
example_tokens,
783+
example_input_pos,
784+
torch.tensor(1.0, dtype=torch.float32),
785+
torch.tensor(1.0, dtype=torch.float32),
786+
torch.tensor(0, dtype=torch.int64),
787+
)
788+
dynamic_shapes = ({1: seq_dim}, {0: seq_dim}, None, None, None)
789+
else:
790+
example_args = (example_tokens, example_input_pos)
791+
dynamic_shapes = ({1: seq_dim}, {0: seq_dim})
756792

757793
print("Exporting with torch.export...")
758794
with torch.no_grad():
759795
exported = export(
760796
model,
761-
(example_tokens, example_input_pos),
797+
example_args,
762798
dynamic_shapes=dynamic_shapes,
763799
strict=True,
764800
)
@@ -781,6 +817,7 @@ def _export_mlx(model, config, args):
781817
"use_kv_cache": True,
782818
"use_sdpa_with_kv_cache": False,
783819
"enable_dynamic_shape": True,
820+
"use_sampling": sample,
784821
}
785822
et_prog = to_edge_transform_and_lower(
786823
exported,
@@ -1182,6 +1219,13 @@ def main(): # noqa: C901
11821219
action="store_true",
11831220
help="Disable split-K (flash-decoding) SDPA for decode; use tiled SDPA instead.",
11841221
)
1222+
parser.add_argument(
1223+
"--sample",
1224+
action="store_true",
1225+
help="MLX only: sample the next token on-device (Gumbel-max with "
1226+
"temperature/top_p/seed runtime inputs) instead of returning logits "
1227+
"for host-side sampling.",
1228+
)
11851229
parser.add_argument(
11861230
"--moe-activation-dtype",
11871231
choices=["bf16", "int8"],
@@ -1230,6 +1274,9 @@ def main(): # noqa: C901
12301274
if args.turboquant:
12311275
parser.error("--turboquant is not supported with --backend mlx")
12321276

1277+
if args.sample and args.backend != "mlx":
1278+
parser.error("--sample is only supported with --backend mlx")
1279+
12331280
if args.backend == "metal":
12341281
if args.turboquant:
12351282
parser.error("--turboquant is not supported with --backend metal")

‎examples/models/qwen3_5_moe/main.cpp‎

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,7 @@
2121
#include <algorithm>
2222
#include <cinttypes>
2323
#include <fstream>
24+
#include <random>
2425
#include <string>
2526
#include <vector>
2627

@@ -37,6 +38,17 @@ DEFINE_string(
3738
"",
3839
"Path to file containing prompt text (overrides --prompt).");
3940
DEFINE_double(temperature, 0.8, "Sampling temperature (0 = greedy).");
41+
DEFINE_double(
42+
top_p,
43+
1.0,
44+
"Nucleus sampling top_p in (0, 1]; 1.0 = off. Requires a model exported "
45+
"with --sample (MLX on-device sampling).");
46+
DEFINE_int64(
47+
seed,
48+
-1,
49+
"Base RNG seed for on-device sampling; the runner increments it per token. "
50+
"-1 (default) draws a random seed each run; set a value for reproducible "
51+
"output. Requires a model exported with --sample.");
4052
DEFINE_int32(max_new_tokens, 128, "Maximum tokens to generate.");
4153
DEFINE_int32(
4254
warmup,
@@ -149,6 +161,20 @@ int main(int argc, char** argv) {
149161
// printed only on the first iteration (coherence check).
150162
llm::SamplingConfig sampling;
151163
sampling.temperature = static_cast<float>(FLAGS_temperature);
164+
sampling.top_p = static_cast<float>(FLAGS_top_p);
165+
// Only a --sample model uses the seed; randomize an unset seed for those and
166+
// leave non-sample models at 0 so they don't trip the top_p/seed guard.
167+
const auto& md = engine->metadata();
168+
const auto us_it = md.find("use_sampling");
169+
const bool model_samples = us_it != md.end() && us_it->second != 0;
170+
uint64_t base_seed = FLAGS_seed < 0 ? 0 : static_cast<uint64_t>(FLAGS_seed);
171+
if (model_samples && FLAGS_seed < 0) {
172+
base_seed = static_cast<uint64_t>(std::random_device{}());
173+
}
174+
sampling.seed = base_seed;
175+
if (model_samples) {
176+
ET_LOG(Info, "Sampling base seed: %" PRIu64, base_seed);
177+
}
152178
const int total_iters = FLAGS_warmup + std::max(1, FLAGS_num_iters);
153179
std::vector<double> prefill_tps_samples;
154180
std::vector<double> decode_tps_samples;

‎examples/models/qwen3_5_moe/qwen35_moe_engine.cpp‎

Lines changed: 118 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -57,11 +57,18 @@ constexpr const char* kDecodeMethod = "decode";
5757
// `forward` method was compiled for. Read into the metadata map in create().
5858
constexpr const char* kMaxPrefillChunk = "get_max_prefill_chunk";
5959

60+
// Constant method (MLX --sample exports) signalling that `forward` samples a
61+
// token id on-device and returns it directly. Read into the metadata map in
62+
// create().
63+
constexpr const char* kUseSampling = "use_sampling";
64+
6065
Result<uint64_t> read_sampled_token(
6166
const executorch::aten::Tensor& output,
62-
float temperature) {
67+
float temperature,
68+
bool use_sampling) {
6369
#ifdef EXECUTORCH_BUILD_CUDA
6470
(void)temperature;
71+
(void)use_sampling;
6572
const void* ptr = output.const_data_ptr();
6673
cudaPointerAttributes attrs;
6774
const bool on_device = cudaPointerGetAttributes(&attrs, ptr) == cudaSuccess &&
@@ -78,6 +85,11 @@ Result<uint64_t> read_sampled_token(
7885
}
7986
return static_cast<uint64_t>(llrintf(val));
8087
#else
88+
// On-device sampling (MLX --sample): `forward` already returns the int64
89+
// token id, so read it directly instead of sampling host-side from logits.
90+
if (use_sampling) {
91+
return static_cast<uint64_t>(output.const_data_ptr<int64_t>()[0]);
92+
}
8193
return static_cast<uint64_t>(
8294
logits_to_token(output, temperature < 0.0f ? 0.0f : temperature));
8395
#endif
@@ -208,6 +220,19 @@ class Qwen35MoESession : public LLMSession {
208220
#ifdef EXECUTORCH_BUILD_CUDA
209221
temp_tensor_ =
210222
from_blob(&temp_val_, {1}, executorch::aten::ScalarType::Float);
223+
#endif
224+
#ifdef EXECUTORCH_BUILD_MLX
225+
if (auto it = metadata_.find(kUseSampling); it != metadata_.end()) {
226+
use_sampling_ = it->second != 0;
227+
}
228+
// 0-dim scalar inputs appended to `forward` when sampling on-device. Reused
229+
// across calls; only the backing scalar is rewritten per token.
230+
temp_tensor_mlx_ =
231+
from_blob(&temp_val_mlx_, {}, executorch::aten::ScalarType::Float);
232+
top_p_tensor_ =
233+
from_blob(&top_p_val_, {}, executorch::aten::ScalarType::Float);
234+
seed_tensor_ =
235+
from_blob(&seed_val_, {}, executorch::aten::ScalarType::Long);
211236
#endif
212237
}
213238

@@ -231,15 +256,27 @@ class Qwen35MoESession : public LLMSession {
231256
}
232257
float first_token_temp = temperature_;
233258
if (initial_sampling != nullptr) {
234-
if (initial_sampling->top_p != 1.0f || initial_sampling->top_k != 0 ||
235-
initial_sampling->seed != 0) {
259+
if (initial_sampling->top_k != 0) {
260+
ET_LOG(Error, "prefill_tokens: top_k is not implemented");
261+
return Error::NotSupported;
262+
}
263+
if (!use_sampling_ &&
264+
(initial_sampling->top_p != 1.0f || initial_sampling->seed != 0)) {
236265
ET_LOG(
237266
Error,
238-
"prefill_tokens: only temperature is supported; top_p/top_k/seed "
239-
"are not yet implemented");
267+
"prefill_tokens: top_p/seed require a sampling model "
268+
"(export with --sample); only temperature is supported otherwise");
240269
return Error::NotSupported;
241270
}
242271
first_token_temp = initial_sampling->temperature;
272+
if (use_sampling_) {
273+
if (!valid_top_p(initial_sampling->top_p)) {
274+
ET_LOG(Error, "prefill_tokens: top_p must be in (0, 1]");
275+
return Error::InvalidArgument;
276+
}
277+
top_p_ = initial_sampling->top_p;
278+
seed_ = initial_sampling->seed; // base seed for the kept (token 0) draw
279+
}
243280
}
244281
if (!valid_temperature(first_token_temp)) {
245282
ET_LOG(Error, "prefill_tokens: temperature must be -1 or in [0, 1]");
@@ -307,6 +344,14 @@ class Qwen35MoESession : public LLMSession {
307344
#ifdef EXECUTORCH_BUILD_CUDA
308345
set_temp(first_token_temp);
309346
inputs.push_back(EValue(temp_tensor_));
347+
#endif
348+
#ifdef EXECUTORCH_BUILD_MLX
349+
if (use_sampling_) {
350+
set_sampling_inputs(first_token_temp, top_p_, seed_);
351+
inputs.push_back(EValue(temp_tensor_mlx_));
352+
inputs.push_back(EValue(top_p_tensor_));
353+
inputs.push_back(EValue(seed_tensor_));
354+
}
310355
#endif
311356
auto sampled =
312357
run_locked(method, inputs, first_token_temp, /*sync_after=*/true);
@@ -316,15 +361,24 @@ class Qwen35MoESession : public LLMSession {
316361
}
317362
pending_ = sampled_token;
318363
prev_decode_token_.reset();
364+
#ifdef EXECUTORCH_BUILD_MLX
365+
if (use_sampling_) {
366+
seed_ += 1; // first decode token draws with base+1
367+
}
368+
#endif
319369
return Error::Ok;
320370
}
321371

322372
Result<DecodeResult> decode_one(const SamplingConfig& sampling) override {
323-
if (sampling.top_p != 1.0f || sampling.top_k != 0 || sampling.seed != 0) {
373+
if (sampling.top_k != 0) {
374+
ET_LOG(Error, "Qwen35MoESession: top_k is not implemented");
375+
return Error::NotSupported;
376+
}
377+
if (!use_sampling_ && (sampling.top_p != 1.0f || sampling.seed != 0)) {
324378
ET_LOG(
325379
Error,
326-
"Qwen35MoESession: only temperature is supported; top_p/top_k/seed "
327-
"are not implemented");
380+
"Qwen35MoESession: top_p/seed require a sampling model "
381+
"(export with --sample); only temperature is supported otherwise");
328382
return Error::NotSupported;
329383
}
330384
if (!valid_temperature(sampling.temperature)) {
@@ -336,6 +390,13 @@ class Qwen35MoESession : public LLMSession {
336390
InvalidState,
337391
"decode_one requires a pending token; call prefill_tokens() first");
338392
temperature_ = sampling.temperature;
393+
if (use_sampling_) {
394+
if (!valid_top_p(sampling.top_p)) {
395+
ET_LOG(Error, "decode_one: top_p must be in (0, 1]");
396+
return Error::InvalidArgument;
397+
}
398+
top_p_ = sampling.top_p;
399+
}
339400

340401
if (stop_.load(std::memory_order_relaxed)) {
341402
return DecodeResult{0, "", /*is_eos=*/false, /*is_terminal=*/true};
@@ -380,13 +441,26 @@ class Qwen35MoESession : public LLMSession {
380441
#ifdef EXECUTORCH_BUILD_CUDA
381442
set_temp(temperature_);
382443
inputs.push_back(EValue(temp_tensor_));
444+
#endif
445+
#ifdef EXECUTORCH_BUILD_MLX
446+
if (use_sampling_) {
447+
set_sampling_inputs(temperature_, top_p_, seed_);
448+
inputs.push_back(EValue(temp_tensor_mlx_));
449+
inputs.push_back(EValue(top_p_tensor_));
450+
inputs.push_back(EValue(seed_tensor_));
451+
}
383452
#endif
384453
auto sampled =
385454
run_locked(kDecodeMethod, inputs, temperature_, /*sync_after=*/false);
386455
ET_CHECK_OK_OR_RETURN_ERROR(sampled.error());
387456
pending_ = sampled.get();
388457
prev_decode_token_ = token;
389458
pos_ += 1;
459+
#ifdef EXECUTORCH_BUILD_MLX
460+
if (use_sampling_) {
461+
seed_ += 1; // next token draws with the next seed in the schedule
462+
}
463+
#endif
390464
return DecodeResult{
391465
token, std::move(text_piece), /*is_eos=*/false, /*is_terminal=*/false};
392466
}
@@ -412,12 +486,26 @@ class Qwen35MoESession : public LLMSession {
412486
return temperature == -1.0f || (temperature >= 0.0f && temperature <= 1.0f);
413487
}
414488

489+
static bool valid_top_p(float top_p) {
490+
return top_p > 0.0f && top_p <= 1.0f;
491+
}
492+
415493
#ifdef EXECUTORCH_BUILD_CUDA
416494
void set_temp(float t) {
417495
temp_val_ = (t <= 0.0f) ? 1e-6f : t;
418496
}
419497
#endif
420498

499+
#ifdef EXECUTORCH_BUILD_MLX
500+
// Rewrite the backing scalars for the reused 0-dim sampling input tensors.
501+
// The -1 temperature sentinel maps to 0 (greedy).
502+
void set_sampling_inputs(float temp, float top_p, uint64_t seed) {
503+
temp_val_mlx_ = (temp < 0.0f) ? 0.0f : temp;
504+
top_p_val_ = top_p;
505+
seed_val_ = static_cast<int64_t>(seed);
506+
}
507+
#endif
508+
421509
Result<uint64_t> run_locked(
422510
const char* method,
423511
std::vector<EValue>& inputs,
@@ -434,7 +522,8 @@ class Qwen35MoESession : public LLMSession {
434522
auto res = module_->execute(method, inputs);
435523
#endif
436524
ET_CHECK_OK_OR_RETURN_ERROR(res.error());
437-
auto sampled = read_sampled_token(res.get()[0].toTensor(), temperature);
525+
auto sampled =
526+
read_sampled_token(res.get()[0].toTensor(), temperature, use_sampling_);
438527
ET_CHECK_OK_OR_RETURN_ERROR(sampled.error());
439528
#ifdef EXECUTORCH_BUILD_CUDA
440529
if (sync_after && cudaDeviceSynchronize() != cudaSuccess) {
@@ -460,6 +549,13 @@ class Qwen35MoESession : public LLMSession {
460549
float temperature_ = -1.0f;
461550
std::atomic<bool> stop_{false};
462551

552+
// On-device sampling state (MLX --sample). top_p_/seed_ are engine-managed:
553+
// the seed increments per generated token for decorrelated, reproducible
554+
// draws.
555+
bool use_sampling_ = false;
556+
float top_p_ = 1.0f;
557+
uint64_t seed_ = 0;
558+
463559
int64_t decode_token_data_[1] = {0};
464560
int64_t decode_pos_data_[1] = {0};
465561
TensorPtr decode_tokens_;
@@ -472,6 +568,14 @@ class Qwen35MoESession : public LLMSession {
472568
float temp_val_ = 1e-6f;
473569
TensorPtr temp_tensor_;
474570
#endif
571+
#ifdef EXECUTORCH_BUILD_MLX
572+
float temp_val_mlx_ = 0.0f;
573+
float top_p_val_ = 1.0f;
574+
int64_t seed_val_ = 0;
575+
TensorPtr temp_tensor_mlx_;
576+
TensorPtr top_p_tensor_;
577+
TensorPtr seed_tensor_;
578+
#endif
475579
};
476580

477581
} // namespace
@@ -513,6 +617,11 @@ Result<std::unique_ptr<Qwen35MoEEngine>> Qwen35MoEEngine::create(
513617
if (auto mpc = meta_module->get(kMaxPrefillChunk); mpc.ok()) {
514618
metadata_result.get()[kMaxPrefillChunk] = mpc->toScalar().to<int64_t>();
515619
}
620+
// get_llm_metadata doesn't harvest use_sampling; surface it so the session
621+
// knows `forward` returns a token id (on-device) rather than logits.
622+
if (auto us = meta_module->get(kUseSampling); us.ok()) {
623+
metadata_result.get()[kUseSampling] = us->toScalar().to<int64_t>();
624+
}
516625
#endif
517626
auto eos_ids = get_eos_ids(tokenizer.get(), meta_module.get());
518627
// This export's metadata doesn't carry the chat-turn EOS (config.json has no

0 commit comments

Comments
 (0)