Skip to content

Commit d863a27

Browse files
authored
[ExecuTorch][WebGPU] Repair stacked merge artifacts (pytorch#21669)
The squash merge for pytorch#21646 duplicated the fp16 upload and BOOL conversion paths and dropped three native route tests. Restore the source that matches the reviewed fbsource stack, including compile-gating timestamp profiling. Test Plan: exact tree comparison against the reconstructed 21-diff stack; git diff --check; GitHub CI. Co-authored-with: OpenAI Codex. Signed-off-by: Julian Ng-Thow-Hing <juliannth@meta.com>
1 parent ceca90f commit d863a27

3 files changed

Lines changed: 62 additions & 115 deletions

File tree

‎backends/webgpu/runtime/WebGPUGraph.cpp‎

Lines changed: 1 addition & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -1482,21 +1482,6 @@ void WebGPUGraph::copy_inputs(const std::vector<InputData>& inputs) {
14821482
continue;
14831483
}
14841484

1485-
// Require an explicit fp32 host dtype, not merely "not int64": inferring
1486-
// the narrow from the 2:1 byte ratio alone would silently reinterpret a
1487-
// same-sized non-fp32 host buffer (e.g. a stale int32) as fp32.
1488-
if (in.host_is_fp32 && buffer_is_fp16 && in.nbytes == live_nbytes * 2) {
1489-
const size_t numel = live_nbytes / sizeof(uint16_t);
1490-
const float* src = static_cast<const float*>(in.data);
1491-
std::vector<executorch::runtime::etensor::Half> narrowed(numel);
1492-
for (size_t e = 0; e < numel; e++) {
1493-
narrowed[e] = executorch::runtime::etensor::Half(src[e]);
1494-
}
1495-
wgpuQueueWriteBuffer(
1496-
queue_, tensor.buffer, 0, narrowed.data(), live_nbytes);
1497-
continue;
1498-
}
1499-
15001485
throw std::runtime_error(
15011486
"WebGPU: unsupported input copy for input " + std::to_string(i) +
15021487
" (host " + std::to_string(in.nbytes) + " bytes" +
@@ -1533,17 +1518,10 @@ constexpr uint32_t kRouteK16CausalBound = 1u << 11;
15331518
constexpr uint32_t kRouteBicolSubgroup = 1u << 12;
15341519
constexpr uint32_t kRouteQwen3Q16K16 = 1u << 13;
15351520
constexpr uint32_t kRouteQwen3Q32K16 = 1u << 14;
1536-
#endif // WGPU_BACKEND_ENABLE_PROFILING
1537-
1538-
// Bench gate: compiled out unless WGPU_BACKEND_ENABLE_PROFILING; then the
1539-
// WEBGPU_TIMESTAMP_QUERY env var enables per-pass GPU timestamp queries.
15401521
bool should_timestamp_query() {
1541-
#ifdef WGPU_BACKEND_ENABLE_PROFILING
15421522
return std::getenv("WEBGPU_TIMESTAMP_QUERY") != nullptr;
1543-
#else
1544-
return false;
1545-
#endif
15461523
}
1524+
#endif // WGPU_BACKEND_ENABLE_PROFILING
15471525
} // namespace
15481526

15491527
#ifdef WGPU_BACKEND_ENABLE_PROFILING

‎backends/webgpu/runtime/ops/to_copy/ToCopy.cpp‎

Lines changed: 0 additions & 92 deletions
Original file line numberDiff line numberDiff line change
@@ -203,98 +203,6 @@ void add_bool_to_float_op(WebGPUGraph& graph, int in_id, int out_id) {
203203
});
204204
}
205205

206-
// Decode byte-packed bool storage into numeric fp32 values.
207-
void add_bool_to_float_op(WebGPUGraph& graph, int in_id, int out_id) {
208-
WGPUDevice device = graph.device();
209-
const auto& in_tensor = graph.get_tensor(in_id);
210-
const auto& out_tensor = graph.get_tensor(out_id);
211-
if (in_tensor.buffer == nullptr || out_tensor.buffer == nullptr) {
212-
throw std::runtime_error("to_copy_bool_to_float: null buffer binding");
213-
}
214-
if (!in_tensor.is_bool || in_tensor.elem_size != 1 || out_tensor.is_int ||
215-
out_tensor.elem_size != sizeof(float) ||
216-
out_tensor.nbytes % sizeof(float) != 0 ||
217-
out_tensor.nbytes / sizeof(float) != in_tensor.nbytes) {
218-
throw std::runtime_error("to_copy_bool_to_float: dtype/numel mismatch");
219-
}
220-
if (in_tensor.nbytes == 0u ||
221-
in_tensor.nbytes > std::numeric_limits<uint32_t>::max()) {
222-
throw std::runtime_error(
223-
"to_copy_bool_to_float: numel must be nonzero and fit u32");
224-
}
225-
226-
const uint32_t num_elements = static_cast<uint32_t>(in_tensor.nbytes);
227-
const uint64_t input_bind_size_u64 =
228-
(static_cast<uint64_t>(in_tensor.nbytes) + 3u) & ~uint64_t(3);
229-
if (input_bind_size_u64 > std::numeric_limits<size_t>::max()) {
230-
throw std::runtime_error(
231-
"to_copy_bool_to_float: input binding size overflows");
232-
}
233-
const size_t input_bind_size = static_cast<size_t>(input_bind_size_u64);
234-
235-
const uint32_t wg_size =
236-
utils::clamp_workgroup_size(device, kToCopyBoolToFloatWorkgroupSizeX);
237-
const uint32_t workgroup_count = utils::compute_1d_workgroup_count(
238-
device, num_elements, wg_size, "to_copy_bool_to_float");
239-
240-
WGPUConstantEntry wg_size_constant = {};
241-
wg_size_constant.key = {"wg_size", WGPU_STRLEN};
242-
wg_size_constant.value = static_cast<double>(wg_size);
243-
244-
ConvertParams params = {};
245-
params.num_elements = num_elements;
246-
WGPUBuffer uniform_buffer =
247-
utils::make_uniform(device, &params, sizeof(ConvertParams));
248-
graph.add_uniform_buffer_bytes(sizeof(ConvertParams));
249-
250-
utils::ComputePipelineBundle bundle = utils::make_compute_pipeline(
251-
device,
252-
kToCopyBoolToFloatWGSL,
253-
{
254-
{0,
255-
WGPUBufferBindingType_ReadOnlyStorage,
256-
in_tensor.buffer,
257-
input_bind_size},
258-
{1,
259-
WGPUBufferBindingType_Storage,
260-
out_tensor.buffer,
261-
out_tensor.nbytes},
262-
{2,
263-
WGPUBufferBindingType_Uniform,
264-
uniform_buffer,
265-
sizeof(ConvertParams)},
266-
},
267-
&wg_size_constant,
268-
1);
269-
270-
const size_t dispatch_idx =
271-
graph.add_dispatch({bundle.pipeline, bundle.bind_group, workgroup_count});
272-
273-
WGPUBuffer params_buf = uniform_buffer;
274-
graph.add_tensor_resize_hook(
275-
in_id,
276-
[in_id, out_id, wg_size, dispatch_idx, params_buf](WebGPUGraph& g) {
277-
const auto& dims = g.cur_dims(in_id);
278-
const uint64_t numel = utils::numel_of(dims);
279-
if (numel == 0u || numel > std::numeric_limits<uint32_t>::max()) {
280-
throw std::runtime_error(
281-
"to_copy_bool_to_float(resize): invalid numel");
282-
}
283-
g.set_cur_dims(out_id, dims);
284-
ConvertParams p = {};
285-
p.num_elements = static_cast<uint32_t>(numel);
286-
wgpuQueueWriteBuffer(g.queue(), params_buf, 0, &p, sizeof(p));
287-
g.dispatch_at(dispatch_idx).workgroup_count_x =
288-
utils::compute_1d_workgroup_count(
289-
g.device(),
290-
static_cast<uint32_t>(numel),
291-
wg_size,
292-
"to_copy_bool_to_float(resize)");
293-
});
294-
295-
graph.own_uniform_buffer(uniform_buffer);
296-
}
297-
298206
void to_copy_impl(WebGPUGraph& graph, const std::vector<int>& args) {
299207
// aten._to_copy.default args: [self, ...kwargs, out]; out = last value id.
300208
add_to_copy_node(graph, args.at(0), args.at(args.size() - 1));

‎backends/webgpu/test/native/test_compute_dispatch.cpp‎

Lines changed: 61 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -239,6 +239,67 @@ void build_conv1d_route_graph(
239239
graph.build(fbb.GetBufferPointer(), nullptr, 0, nullptr);
240240
}
241241

242+
TEST(WebGPUShaderRegistry, FindsKnownShaderAndRejectsUnknownName) {
243+
const WebGPUShaderInfo& sigmoid = get_webgpu_shader_info("sigmoid");
244+
EXPECT_EQ(sigmoid.name, "sigmoid");
245+
EXPECT_NE(sigmoid.source, nullptr);
246+
EXPECT_GT(sigmoid.workgroup_size_x, 0u);
247+
EXPECT_THROW(
248+
get_webgpu_shader_info("not_a_registered_shader"), std::runtime_error);
249+
}
250+
251+
TEST(WebGPUQ4RouteSignal, PreservesStaticAndRecordsBothDynamicSignals) {
252+
WebGPUGraph static_graph;
253+
static_graph.set_device(g_device);
254+
build_q4_route_graph(static_graph, Q4RouteSignal::Static);
255+
const auto static_dispatches = q4_dispatches(static_graph);
256+
ASSERT_EQ(static_dispatches.size(), 1);
257+
EXPECT_EQ(static_graph.num_dispatches(), 1);
258+
EXPECT_NE(static_dispatches[0]->pipeline, nullptr);
259+
EXPECT_NE(static_dispatches[0]->bind_group, nullptr);
260+
EXPECT_FALSE(static_graph.has_dynamic_shapes());
261+
EXPECT_FALSE(static_graph.config().record_q4gsw_decode_route);
262+
263+
WebGPUGraph legacy_graph;
264+
legacy_graph.set_device(g_device);
265+
build_q4_route_graph(legacy_graph, Q4RouteSignal::LegacyGraphMarker);
266+
EXPECT_TRUE(legacy_graph.has_dynamic_shapes());
267+
EXPECT_FALSE(legacy_graph.config().record_q4gsw_decode_route);
268+
EXPECT_EQ(legacy_graph.num_dispatches(), 3);
269+
expect_dual_q4_topology(legacy_graph);
270+
271+
WebGPUGraph explicit_graph;
272+
explicit_graph.set_device(g_device);
273+
build_q4_route_graph(explicit_graph, Q4RouteSignal::ExplicitOption);
274+
EXPECT_FALSE(explicit_graph.has_dynamic_shapes());
275+
EXPECT_TRUE(explicit_graph.config().record_q4gsw_decode_route);
276+
EXPECT_EQ(explicit_graph.num_dispatches(), 2);
277+
expect_dual_q4_topology(explicit_graph);
278+
}
279+
280+
TEST(WebGPUConv1dRoute, SelectsPointwiseGeneralAndSingleChannelDepthwise) {
281+
const Conv1dRouteCase cases[] = {
282+
{"pointwise", {1, 4, 8}, {6, 4, 1}, {1, 6, 8}, 1, 0, 1, 1, "conv1d_pw"},
283+
{"general", {1, 4, 10}, {6, 4, 3}, {1, 6, 4}, 2, 0, 1, 1, "conv1d"},
284+
{"single-channel depthwise",
285+
{1, 1, 8},
286+
{1, 1, 3},
287+
{1, 1, 8},
288+
1,
289+
1,
290+
1,
291+
1,
292+
"conv1d_dw"},
293+
};
294+
for (const Conv1dRouteCase& test_case : cases) {
295+
SCOPED_TRACE(test_case.name);
296+
WebGPUGraph graph;
297+
build_conv1d_route_graph(graph, test_case);
298+
ASSERT_EQ(graph.num_dispatches(), 1);
299+
EXPECT_EQ(graph.dispatch_at(0).kernel_name, test_case.expected_kernel);
300+
}
301+
}
302+
242303
TEST(WebGPUComputeDispatch, PipelineKeyCanonicalizesConstants) {
243304
WebGPUComputeDispatchDescriptor first;
244305
first.shader_name = "sigmoid";

0 commit comments

Comments
 (0)