Skip to content

Commit c8ce66c

Browse files
[executorch][muse] Verify CUDA DFlash drafts on device under any sampling (#24042)
This PR was created by the merge bot to help merge the original PR into the main branch. ghstack PR number: #23984 by @Gasoonjia ^ Please use this as the source of truth for the PR details, comments, and reviews ghstack PR base: https://github.com/pytorch/executorch/tree/gh/Gasoonjia/83/base ghstack PR head: https://github.com/pytorch/executorch/tree/gh/Gasoonjia/83/head Merge bot PR base: https://github.com/pytorch/executorch/tree/gh/Gasoonjia/82/orig Merge bot PR head: https://github.com/pytorch/executorch/tree/gh/Gasoonjia/83/orig Differential Revision: [D125705716](https://our.internmc.facebook.com/intern/diff/D125705716/) @diff-train-skip-merge Co-authored-by: gasoonjia <gasoonjia@icloud.com>
1 parent f6bb20a commit c8ce66c

11 files changed

Lines changed: 858 additions & 152 deletions

‎backends/cuda/batching/cuda_executor.cpp‎

Lines changed: 54 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -810,21 +810,40 @@ bool CudaExecutor::execute(const BatchInput& batch, BatchOutput& out) {
810810
});
811811
}
812812

813+
std::optional<std::vector<Token>> CudaExecutor::sample_rows(
814+
const ::executorch::aten::Tensor& logits,
815+
std::size_t rows,
816+
const std::vector<float>& params) {
817+
if (params.size() != rows * 4) {
818+
return std::nullopt;
819+
}
820+
bool all_greedy = true;
821+
for (std::size_t row = 0; row < rows; ++row) {
822+
all_greedy = all_greedy && params[row * 4] == 0.0f;
823+
}
824+
if (all_greedy) {
825+
return argmax_rows(logits, rows);
826+
}
827+
if (logits.device_type() != ::executorch::aten::DeviceType::CUDA) {
828+
ET_LOG(
829+
Error,
830+
"CudaExecutor: sampled rows need device logits and %s",
831+
kSampleMethod);
832+
return std::nullopt;
833+
}
834+
if (!initialized_ || !has_sampler_ || logits.dim() < 2 || !is_dense(logits) ||
835+
logits.size(logits.dim() - 1) != vocab_size_ ||
836+
logits.scalar_type() != ::executorch::aten::ScalarType::Float ||
837+
rows == 0 ||
838+
rows > static_cast<std::size_t>(logits.numel() / vocab_size_)) {
839+
return std::nullopt;
840+
}
841+
return run_device_sampler(logits, rows, params, false);
842+
}
843+
813844
std::optional<std::vector<Token>> CudaExecutor::sample_on_device(
814845
const ::executorch::aten::Tensor& logits,
815846
const std::vector<SessionId>& sessions) {
816-
using ::executorch::aten::ScalarType;
817-
using ::executorch::aten::SizesType;
818-
const ::executorch::aten::Device cuda(
819-
::executorch::aten::DeviceType::CUDA, logits.device_index());
820-
const auto rows = static_cast<SizesType>(sessions.size());
821-
// The rows that sample lead the logits; the rest are padding.
822-
auto row_logits = make_tensor_ptr(
823-
{rows, static_cast<SizesType>(vocab_size_)},
824-
logits.mutable_data_ptr(),
825-
logits.scalar_type(),
826-
cuda);
827-
828847
const bool all_greedy =
829848
std::all_of(sessions.begin(), sessions.end(), [this](SessionId session) {
830849
return sessions_.greedy(session);
@@ -840,6 +859,25 @@ std::optional<std::vector<Token>> CudaExecutor::sample_on_device(
840859
params.insert(
841860
params.end(), {row->temperature, row->top_p, row->top_k, row->coin});
842861
}
862+
return run_device_sampler(logits, sessions.size(), params, all_greedy);
863+
}
864+
865+
std::optional<std::vector<Token>> CudaExecutor::run_device_sampler(
866+
const ::executorch::aten::Tensor& logits,
867+
std::size_t row_count,
868+
const std::vector<float>& params,
869+
bool all_greedy) {
870+
using ::executorch::aten::ScalarType;
871+
using ::executorch::aten::SizesType;
872+
const ::executorch::aten::Device cuda(
873+
::executorch::aten::DeviceType::CUDA, logits.device_index());
874+
const auto rows = static_cast<SizesType>(row_count);
875+
// The rows that sample lead the logits; the rest are padding.
876+
auto row_logits = make_tensor_ptr(
877+
{rows, static_cast<SizesType>(vocab_size_)},
878+
logits.mutable_data_ptr(),
879+
logits.scalar_type(),
880+
cuda);
843881

844882
const cudaStream_t stream =
845883
::executorch::extension::cuda::getCallerStream().value_or(
@@ -851,18 +889,18 @@ std::optional<std::vector<Token>> CudaExecutor::sample_on_device(
851889
kArgmaxMethod,
852890
std::vector<::executorch::runtime::EValue>{row_logits});
853891
}
854-
if (sessions.size() > device_params_rows_) {
892+
if (row_count > device_params_rows_) {
855893
if (device_params_ != nullptr) {
856894
(void)cudaFree(device_params_);
857895
device_params_ = nullptr;
858896
device_params_rows_ = 0;
859897
}
860898
ET_CHECK_OR_RETURN_ERROR(
861-
cudaMalloc(&device_params_, sessions.size() * 4 * sizeof(float)) ==
899+
cudaMalloc(&device_params_, row_count * 4 * sizeof(float)) ==
862900
cudaSuccess,
863901
MemoryAllocationFailed,
864902
"CudaExecutor: could not stage the sampler's params");
865-
device_params_rows_ = sessions.size();
903+
device_params_rows_ = row_count;
866904
}
867905
// The backend runs every method on the caller's stream if one is set,
868906
// else the per-thread stream; ordering the upload there puts it ahead of
@@ -890,7 +928,7 @@ std::optional<std::vector<Token>> CudaExecutor::sample_on_device(
890928
ET_LOG(Error, "CudaExecutor: the device sampler returned a bad shape");
891929
return std::nullopt;
892930
}
893-
std::vector<std::int64_t> host(sessions.size());
931+
std::vector<std::int64_t> host(row_count);
894932
const auto kind = tokens.device_type() == ::executorch::aten::DeviceType::CUDA
895933
? cudaMemcpyDeviceToHost
896934
: cudaMemcpyHostToHost;

‎backends/cuda/batching/cuda_executor.h‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -174,6 +174,14 @@ class ET_EXPERIMENTAL CudaExecutor : public llm_batching::Executor {
174174
std::optional<std::vector<llm_batching::Token>> argmax_rows(
175175
const ::executorch::aten::Tensor& logits,
176176
std::size_t rows);
177+
// Samples device logits with the exported sampler. params is row-major
178+
// [rows, 4] float32: temperature, top_p, top_k and a uniform coin, as
179+
// extension/llm/batching/sampler.py consumes them. All-greedy rows take
180+
// sample_argmax.
181+
std::optional<std::vector<llm_batching::Token>> sample_rows(
182+
const ::executorch::aten::Tensor& logits,
183+
std::size_t rows,
184+
const std::vector<float>& params);
177185
::executorch::extension::Module& module() {
178186
return *module_;
179187
}
@@ -212,6 +220,11 @@ class ET_EXPERIMENTAL CudaExecutor : public llm_batching::Executor {
212220

213221
// Draws one token per row of `logits`, rows [0, sessions.size()), on the
214222
// device; returns them in row order.
223+
std::optional<std::vector<llm_batching::Token>> run_device_sampler(
224+
const ::executorch::aten::Tensor& logits,
225+
std::size_t rows,
226+
const std::vector<float>& params,
227+
bool all_greedy);
215228
std::optional<std::vector<llm_batching::Token>> sample_on_device(
216229
const ::executorch::aten::Tensor& logits,
217230
const std::vector<llm_batching::SessionId>& sessions);

‎examples/models/muse-glimmer/export/export_dflash.py‎

Lines changed: 46 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -95,6 +95,50 @@ def _export_cuda_sampler_methods(
9595
}
9696

9797

98+
VERIFY_METHOD = "dflash_verify"
99+
VERIFY_ARGMAX_METHOD = "dflash_verify_argmax"
100+
101+
102+
def cuda_owner_verifier_programs(
103+
vocab_size: int, max_owners: int, block: int
104+
) -> dict[str, torch.export.ExportedProgram]:
105+
"""Batched on-device verification of up to ``max_owners`` draft blocks."""
106+
from executorch.examples.models.muse_glimmer.model.dflash_token_sampler import (
107+
DFlashOwnerArgmaxVerifier,
108+
DFlashOwnerVerifier,
109+
)
110+
from torch.export import Dim, export
111+
112+
if block < 3:
113+
raise ValueError("CUDA DFlash verifiers need a block of at least 3")
114+
owners = (
115+
Dim("verify_owners", min=1, max=max_owners) if max_owners > 1 else Dim.STATIC
116+
)
117+
# A one-proposal width would specialize its size-1 proposal axis.
118+
width = Dim("verify_width", min=3, max=block)
119+
logits = torch.zeros(max_owners, block, vocab_size)
120+
candidates = torch.zeros(max_owners, block, dtype=torch.long)
121+
params = torch.zeros(max_owners, block + 3)
122+
return {
123+
VERIFY_METHOD: export(
124+
DFlashOwnerVerifier(),
125+
(logits, candidates, params),
126+
dynamic_shapes=(
127+
{0: owners, 1: width},
128+
{0: owners, 1: width},
129+
{0: owners, 1: width + 3},
130+
),
131+
strict=True,
132+
),
133+
VERIFY_ARGMAX_METHOD: export(
134+
DFlashOwnerArgmaxVerifier(),
135+
(logits, candidates),
136+
dynamic_shapes=({0: owners, 1: width}, {0: owners, 1: width}),
137+
strict=True,
138+
),
139+
}
140+
141+
98142
def _cuda_propagate_device_config() -> dict[str, PropagateDeviceConfig]:
99143
"""Keeps the speculative chain's vocabulary-wide tensors on the device."""
100144
from executorch.exir.passes.propagate_device_pass import PropagateDeviceConfig
@@ -337,6 +381,7 @@ def prepare_cuda_packed_draft(
337381
),
338382
strict=True,
339383
)
384+
programs = {"draft_forward": program}
340385
geometry = offgraph_kv_cache_geometry(draft_model)
341386
metadata = {
342387
key.replace("get_", "get_draft_", 1): value for key, value in geometry.items()
@@ -357,7 +402,7 @@ def prepare_cuda_packed_draft(
357402
"get_dflash_draft_cache_key": "muse_glimmer_draft",
358403
}
359404
)
360-
return {"draft_forward": program}, manifest, metadata
405+
return programs, manifest, metadata
361406

362407

363408
def _dflash_constant_methods(

‎examples/models/muse-glimmer/export/export_solo.py‎

Lines changed: 22 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -318,7 +318,7 @@ def sampler_programs(vocab_size: int, max_rows: int = MAX_SAMPLE_ROWS):
318318
}
319319

320320

321-
def device_residency(forward_methods, device_sampling: bool):
321+
def device_residency(forward_methods, device_sampling: bool, samplers=()):
322322
"""Per-method ``PropagateDeviceConfig`` for ``to_executorch``.
323323
324324
With device sampling the forwards keep their logits in device memory, and
@@ -331,7 +331,7 @@ def device_residency(forward_methods, device_sampling: bool):
331331
for name in forward_methods
332332
}
333333
if device_sampling:
334-
for name in (SAMPLE_METHOD, ARGMAX_METHOD):
334+
for name in (SAMPLE_METHOD, ARGMAX_METHOD, *samplers):
335335
device[name] = PropagateDeviceConfig(skip_h2d_for_method_inputs=True)
336336
return device
337337

@@ -469,19 +469,31 @@ def export_cuda_batching( # noqa: C901
469469
if dflash_draft is not None:
470470
programs.update(draft_programs)
471471
del draft_programs, dflash_draft
472-
cache_methods = forward_methods + (
473-
["draft_forward"] if "draft_forward" in programs else []
474-
)
472+
draft_methods = [name for name in programs if name.startswith("draft_forward")]
473+
cache_methods = forward_methods + draft_methods
474+
verifiers = {}
475475
if device_sampling:
476476
print(f"Exporting {SAMPLE_METHOD} and {ARGMAX_METHOD}...")
477477
programs.update(sampler_programs(config.vocab_size))
478+
if "draft_forward" in programs:
479+
from executorch.examples.models.muse_glimmer.export.export_dflash import (
480+
cuda_owner_verifier_programs,
481+
)
482+
483+
verifiers = cuda_owner_verifier_programs(
484+
config.vocab_size,
485+
max_draft_owners,
486+
draft_metadata["get_block_size"],
487+
)
488+
print(f"Exporting {', '.join(verifiers)}...")
489+
programs.update(verifiers)
478490

479491
def partitioner(name: str) -> CudaPartitioner:
480492
if name not in cache_methods:
481493
return CudaPartitioner(
482494
[CudaBackend.generate_method_name_compile_spec(name)]
483495
)
484-
is_draft = name == "draft_forward"
496+
is_draft = name in draft_methods
485497
specs = [
486498
CudaBackend.generate_method_name_compile_spec(name),
487499
CompileSpec("low_memory_mode", b"ON"),
@@ -494,7 +506,7 @@ def partitioner(name: str) -> CudaPartitioner:
494506
b"2:0" if is_draft else step_width_spec(name),
495507
),
496508
]
497-
if "draft_forward" in cache_methods:
509+
if draft_methods:
498510
specs.append(
499511
CompileSpec(
500512
"llm_cache_registry_key",
@@ -539,7 +551,9 @@ def partitioner(name: str) -> CudaPartitioner:
539551
extract_delegate_segments=True,
540552
do_quant_fusion_and_const_prop=True,
541553
memory_planning_pass=MemoryPlanningPass(alloc_graph_input=False),
542-
propagate_device_config=device_residency(cache_methods, device_sampling),
554+
propagate_device_config=device_residency(
555+
cache_methods, device_sampling, tuple(verifiers)
556+
),
543557
emit_mutable_buffer_names=True,
544558
),
545559
)

0 commit comments

Comments
 (0)