馃悰 Describe the bug
conv2d_local_wg_size() in backends/vulkan/runtime/graph/ops/impl/Convolution.cpp classifies the convolution method by shader name only, and the condition it uses matches every conv2d shader. As a result its SlidingWindow branch is unreachable for all conv2d variants, and they all receive the pointwise local workgroup shape.
// conv2d_local_wg_size()
Conv2dMethod method;
if (shader.kernel_name.find("conv2d_pw") != std::string::npos ||
(shader.kernel_name.find("conv2d") != std::string::npos &&
shader.kernel_name.find("conv_transpose2d") == std::string::npos)) {
method = Conv2dMethod::Pointwise;
} else {
method = Conv2dMethod::SlidingWindow;
}
The sliding window shader is itself named conv2d, so it matches find("conv2d") != npos and is classified as Pointwise. The same applies to conv2d_dw. Only conv_transpose2d reaches the else, and it is labelled SlidingWindow.
The intent is visible in the sibling function directly above, conv2d_global_wg_size(), which uses the identical outer name test but then disambiguates properly by inspecting the weights:
const auto& weight_sizes = graph->get_tref(weight_data)->sizes;
if (weight_sizes.at(2) == 1 && weight_sizes.at(3) == 1) {
method = Conv2dMethod::Pointwise;
} else {
method = Conv2dMethod::SlidingWindow;
}
conv2d_local_wg_size() is missing that second step, so the two functions can disagree about the method for the same dispatch: the global size is computed as sliding window while the local size is computed as pointwise.
Impact
Every non-pointwise conv2d takes the {64 / local_wg_size_y, local_wg_size_y, 1} path intended for pointwise convolutions, instead of graph->create_local_wg_size(global_workgroup_size). For the shapes I checked the resulting local size works out to {8, 8, 1} either way, so I have not observed a correctness or performance difference. Filing it as a latent bug rather than an active one: the dead branch means any future change to the sliding window local size selection would silently have no effect.
Found while investigating #21938 and #21939. It is not the cause of either.
Versions
ExecuTorch 1.4.1, Vulkan backend, Adreno 840v2 (Galaxy S26 Ultra, SM-S948B).
cc @SS-JIA @manuelcandales @digantdesai @cbilgin
馃悰 Describe the bug
conv2d_local_wg_size()inbackends/vulkan/runtime/graph/ops/impl/Convolution.cppclassifies the convolution method by shader name only, and the condition it uses matches everyconv2dshader. As a result itsSlidingWindowbranch is unreachable for all conv2d variants, and they all receive the pointwise local workgroup shape.The sliding window shader is itself named
conv2d, so it matchesfind("conv2d") != nposand is classified asPointwise. The same applies toconv2d_dw. Onlyconv_transpose2dreaches theelse, and it is labelledSlidingWindow.The intent is visible in the sibling function directly above,
conv2d_global_wg_size(), which uses the identical outer name test but then disambiguates properly by inspecting the weights:conv2d_local_wg_size()is missing that second step, so the two functions can disagree about the method for the same dispatch: the global size is computed as sliding window while the local size is computed as pointwise.Impact
Every non-pointwise conv2d takes the
{64 / local_wg_size_y, local_wg_size_y, 1}path intended for pointwise convolutions, instead ofgraph->create_local_wg_size(global_workgroup_size). For the shapes I checked the resulting local size works out to{8, 8, 1}either way, so I have not observed a correctness or performance difference. Filing it as a latent bug rather than an active one: the dead branch means any future change to the sliding window local size selection would silently have no effect.Found while investigating #21938 and #21939. It is not the cause of either.
Versions
ExecuTorch 1.4.1, Vulkan backend, Adreno 840v2 (Galaxy S26 Ultra, SM-S948B).
cc @SS-JIA @manuelcandales @digantdesai @cbilgin