From a90027fee0f86170ea8601640c97072e9729a47e Mon Sep 17 00:00:00 2001 From: Stephen Jia Date: Mon, 2 Nov 2026 07:35:38 -0800 Subject: [PATCH 1/2] Update [ghstack-poisoned] --- .../convolution/conv2d/conv2d_pw_tiled.glsl | 88 +++++++++------ .../convolution/conv2d/conv2d_pw_tiled.yaml | 16 +++ .../linear_fp_packed_weight_tile_load.glslh | 22 ++++ .../ops/impl/convolution/conv2d/Conv2dPW.cpp | 106 +++++++++++++++++- .../ops/benchmark/cases/conv2d/conv2d.cpp | 27 +++-- .../cases/conv2d/test_sets/pw_sets.cpp | 29 ++++- 6 files changed, 238 insertions(+), 50 deletions(-) diff --git a/backends/vulkan/runtime/graph/ops/glsl/convolution/conv2d/conv2d_pw_tiled.glsl b/backends/vulkan/runtime/graph/ops/glsl/convolution/conv2d/conv2d_pw_tiled.glsl index e4e8ae7a5ff..a746e4c3091 100644 --- a/backends/vulkan/runtime/graph/ops/glsl/convolution/conv2d/conv2d_pw_tiled.glsl +++ b/backends/vulkan/runtime/graph/ops/glsl/convolution/conv2d/conv2d_pw_tiled.glsl @@ -9,17 +9,23 @@ #version 450 core #define PRECISION ${PRECISION} -#define VEC4_T ${texel_load_type(DTYPE, STORAGE)} +#define VEC4_T ${buffer_gvec_type(DTYPE, 4)} #define TILE_M4 ${TILE_M4} #define TILE_K4 ${TILE_K4} #define TILE_N4 ${TILE_N4} #define TILE_M ${TILE_M} +#define TILE_H ${TILE_H} +#define TILE_W (TILE_M / TILE_H) #define TILE_K ${TILE_K4 * 4} #define TILE_N ${TILE_N4 * 4} ${define_required_extensions(STORAGE, DTYPE)} +$if DTYPE == "half": + #extension GL_EXT_shader_explicit_arithmetic_types_float16 : require +$if WEIGHT_STORAGE == "buffer": + ${define_required_extensions("buffer", DTYPE)} layout(std430) buffer; @@ -27,7 +33,11 @@ layout(std430) buffer; ${layout_declare_tensor(B, "w", "t_out", DTYPE, STORAGE)} ${layout_declare_tensor(B, "r", "t_in", DTYPE, STORAGE)} -${layout_declare_tensor(B, "r", "t_weight_packed", DTYPE, "texture2d")} +${layout_declare_tensor(B, "r", "t_weight_packed", DTYPE, WEIGHT_STORAGE, is_scalar_array=False)} +$if WEIGHT_STORAGE == "buffer": + #define WEIGHT_BUFFER +$if PIXEL_TILES_ON_X: + #define PIXEL_TILES_ON_X ${layout_declare_tensor(B, "r", "t_bias", DTYPE, "texture2d")} ${layout_declare_ubo(B, "ivec4", "in_sizes")} @@ -51,37 +61,40 @@ ${layout_declare_spec_const(C, "int", "activation_type", "0")} #include "gemm/tile_utils/linear_fp_packed_weight_tile_load.glslh" #include "gemm/tile_utils/linear_fp_output_tile_fp_compute.glslh" -// Texel coordinates for the TILE_M spatial positions this invocation covers. +// Texel coordinates for the TILE_W x TILE_H block of output pixels this +// invocation covers. // -// Recovering (x, y) from the flattened spatial index costs a division and a -// modulo by W_out, which is a runtime uniform, so the compiler cannot strength -// reduce either into a shift. Those coordinates depend only on m, so computing -// them once per invocation keeps the K loop free of integer division; done -// per (m, k4) instead, a conv with K4 = 4 issues four times as many divides as -// the arithmetic they feed. +// A block spanning two rows lets each warp read and write whole cache lines of +// the block-interleaved texture layout used by some GPUs. +// +// Locating the block costs a division and a modulo by a runtime uniform, so the +// coordinates are computed once per invocation to keep the K loop free of +// integer division. struct InCoords { ivec2 pos[TILE_M]; ivec2 out_pos[TILE_M]; bool valid[TILE_M]; + bool in_bounds[TILE_M]; }; void compute_in_coords( out InCoords coords, - const int m_start, - const int M, + const int tile_idx_m, + const int tiles_x, const int W_out, + const int H_out, const int W_in, const int H_in) { + const int x0 = (tile_idx_m % tiles_x) * TILE_W; + const int y0 = (tile_idx_m / tiles_x) * TILE_H; [[unroll]] for (int m = 0; m < TILE_M; ++m) { - const int out_spatial = m_start + m; - if (out_spatial >= M) { - coords.pos[m] = ivec2(0); - coords.out_pos[m] = ivec2(0); - coords.valid[m] = false; - continue; - } - const int out_x = out_spatial % W_out; - const int out_y = out_spatial / W_out; + const int x = x0 + m % TILE_W; + const int y = y0 + m / TILE_W; + coords.in_bounds[m] = x < W_out && y < H_out; + // Pixels outside the output reuse the nearest pixel inside it so the K loop + // needs no bounds check for them; their results are never stored. + const int out_x = min(x, W_out - 1); + const int out_y = min(y, H_out - 1); coords.out_pos[m] = ivec2(out_x, out_y); if (stride_1_padding_0 != 0) { coords.pos[m] = ivec2(out_x, out_y); @@ -102,9 +115,13 @@ void load_input_tile_with_checks( const int K4) { [[unroll]] for (int m = 0; m < TILE_M; ++m) { [[unroll]] for (int k4 = 0; k4 < TILE_K4; ++k4) { - if (k4_start + k4 < K4 && coords.valid[m]) { + // The k4 check is redundant when TILE_K4 == 1, but with no branch left in + // the K loop the Adreno compiler reads the weight buffer through a much + // slower path. + if (k4_start + k4 < K4 && + (stride_1_padding_0 != 0 || coords.valid[m])) { tile.data[m][k4] = - texelFetch(t_in, ivec3(coords.pos[m], k4_start + k4), 0); + VEC4_T(texelFetch(t_in, ivec3(coords.pos[m], k4_start + k4), 0)); } else { tile.data[m][k4] = VEC4_T(0.0); } @@ -116,32 +133,34 @@ void store_output_tile_with_checks( const FPOutTile out_tile, const InCoords coords, const int n4_start, - const int m_start, - const int N4, - const int M) { + const int N4) { // Reuse the coordinates computed for this invocation rather than dividing by // W_out again per output texel. [[unroll]] for (int m = 0; m < TILE_M; ++m) { [[unroll]] for (int n4 = 0; n4 < TILE_N4; ++n4) { - if (m_start + m < M && n4_start + n4 < N4) { + if (coords.in_bounds[m] && n4_start + n4 < N4) { VEC4_T texel = out_tile.data[m][n4]; if (activation_type == 1) { texel = max(texel, VEC4_T(0.0)); } else if (activation_type == 2) { texel = clamp(texel, VEC4_T(out_min), VEC4_T(out_max)); } - imageStore(t_out, ivec3(coords.out_pos[m], n4_start + n4), texel); + imageStore(t_out, ivec3(coords.out_pos[m], n4_start + n4), vec4(texel)); } } } } void main() { +#ifdef PIXEL_TILES_ON_X + const int tile_idx_m = int(gl_GlobalInvocationID.x); + const int tile_idx_n = int(gl_GlobalInvocationID.y); +#else const int tile_idx_n = int(gl_GlobalInvocationID.x); const int tile_idx_m = int(gl_GlobalInvocationID.y); +#endif const int n4_start = tile_idx_n * TILE_N4; - const int m_start = tile_idx_m * TILE_M; const int W_in = in_sizes.x; const int H_in = in_sizes.y; @@ -150,11 +169,11 @@ void main() { const int W_out = out_sizes.x; const int H_out = out_sizes.y; - const int M = W_out * H_out; + const int tiles_x = div_up(W_out, TILE_W); const int N = out_sizes.z; const int N4 = div_up_4(N); - if (n4_start >= N4 || m_start >= M) { + if (n4_start >= N4 || tile_idx_m >= tiles_x * div_up(H_out, TILE_H)) { return; } @@ -165,11 +184,12 @@ void main() { FPWeightTile w_tile; InCoords in_coords; - compute_in_coords(in_coords, m_start, M, W_out, W_in, H_in); + compute_in_coords( + in_coords, tile_idx_m, tiles_x, W_out, H_out, W_in, H_in); for (int k4 = 0; k4 < K4; k4++) { load_input_tile_with_checks(in_tile, in_coords, k4, K4); - load_packed_weight_tile_with_checks(w_tile, n4_start, k4, 0, N4, K4); + load_packed_weight_tile_clamped(w_tile, n4_start, k4, 0, N4, K4); fp_accumulate_with_fp_weight(out_tile, in_tile, w_tile); } @@ -178,10 +198,10 @@ void main() { [[unroll]] for (int n4 = 0; n4 < TILE_N4; ++n4) { if (n4_start + n4 < N4) { out_tile.data[m][n4] += - texelFetch(t_bias, ivec2(n4_start + n4, 0), 0); + VEC4_T(texelFetch(t_bias, ivec2(n4_start + n4, 0), 0)); } } } - store_output_tile_with_checks(out_tile, in_coords, n4_start, m_start, N4, M); + store_output_tile_with_checks(out_tile, in_coords, n4_start, N4); } diff --git a/backends/vulkan/runtime/graph/ops/glsl/convolution/conv2d/conv2d_pw_tiled.yaml b/backends/vulkan/runtime/graph/ops/glsl/convolution/conv2d/conv2d_pw_tiled.yaml index 82a6dcd24e2..80299851040 100644 --- a/backends/vulkan/runtime/graph/ops/glsl/convolution/conv2d/conv2d_pw_tiled.yaml +++ b/backends/vulkan/runtime/graph/ops/glsl/convolution/conv2d/conv2d_pw_tiled.yaml @@ -8,13 +8,29 @@ conv2d_pw_tiled: parameter_names_with_default_values: DTYPE: float STORAGE: texture3d + WEIGHT_STORAGE: texture2d TILE_M4: 1 TILE_K4: 1 TILE_N4: 1 TILE_M: 2 + TILE_H: 1 + PIXEL_TILES_ON_X: false generate_variant_forall: DTYPE: - VALUE: float - VALUE: half shader_variants: - NAME: conv2d_pw_tiled + - NAME: conv2d_pw_tiled_bufw + WEIGHT_STORAGE: buffer + - NAME: conv2d_pw_tiled_4x8 + TILE_M: 4 + TILE_H: 2 + TILE_N4: 2 + PIXEL_TILES_ON_X: true + - NAME: conv2d_pw_tiled_4x8_bufw + TILE_M: 4 + TILE_H: 2 + TILE_N4: 2 + WEIGHT_STORAGE: buffer + PIXEL_TILES_ON_X: true diff --git a/backends/vulkan/runtime/graph/ops/glsl/gemm/tile_utils/linear_fp_packed_weight_tile_load.glslh b/backends/vulkan/runtime/graph/ops/glsl/gemm/tile_utils/linear_fp_packed_weight_tile_load.glslh index 5592042f6f7..9011569e005 100644 --- a/backends/vulkan/runtime/graph/ops/glsl/gemm/tile_utils/linear_fp_packed_weight_tile_load.glslh +++ b/backends/vulkan/runtime/graph/ops/glsl/gemm/tile_utils/linear_fp_packed_weight_tile_load.glslh @@ -72,4 +72,26 @@ void load_packed_weight_tile_with_checks( } } +// Out-of-range blocks load the nearest in-range block instead of zeros, which +// keeps branches out of the load. Use only when the results computed from those +// blocks are discarded or multiplied by zero inputs. +void load_packed_weight_tile_clamped( + out FPWeightTile tile, + const int n4_start, + const int k4_start, + const int b, + const int N4, + const int K4) { + [[unroll]] for (int dk4 = 0; dk4 < TILE_K4; dk4++) { + const int k4 = TILE_K4 == 1 ? k4_start : min(k4_start + dk4, K4 - 1); + [[unroll]] for (int dk = 0; dk < 4; dk++) { + const int k = dk4 * 4 + dk; + [[unroll]] for (int n4 = 0; n4 < TILE_N4; n4++) { + tile.data[k][n4] = load_packed_weight_x4( + min(n4_start + n4, N4 - 1), dk, k4, b, K4, N4); + } + } + } +} + #endif // LINEAR_FP_PACKED_WEIGHT_TILE_LOAD_GLSLH diff --git a/backends/vulkan/runtime/graph/ops/impl/convolution/conv2d/Conv2dPW.cpp b/backends/vulkan/runtime/graph/ops/impl/convolution/conv2d/Conv2dPW.cpp index 05853f86a6c..448cd2e7cbb 100644 --- a/backends/vulkan/runtime/graph/ops/impl/convolution/conv2d/Conv2dPW.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/convolution/conv2d/Conv2dPW.cpp @@ -20,6 +20,34 @@ namespace vkcompute { // Shader dispatch utilities // +namespace { + +// Below this many wide tile invocations the GPU is underfilled, and Half +// convolutions are faster with the narrow tile. The threshold sits between the +// largest swept Half grid that was faster with the narrow tile on most GPUs +// (2560 wide tile invocations) and the smallest faster with the wide tile on +// every GPU (12288); grids in between were not swept. +constexpr uint64_t kMinWideTileInvocations = 4096; + +/* + * The 2x2-pixel x 8-channel tile issues half as many input and weight loads per + * FMA as the 2x1-pixel x 4-channel tile. Float uses it on every GPU except + * Mali, where the wide Float tile spills registers in the Mali offline compiler + * and was 2-31% slower than the narrow tile. + */ +bool use_wide_pw_tile(ComputeGraph* graph, const ValueRef out) { + if (graph->dtype_of(out) != vkapi::kHalf) { + return !graph->device_is_mali(); + } + const uint64_t M = uint64_t(graph->size_at(-1, out)) * + graph->size_at(-2, out); + const uint64_t N4 = utils::div_up_4(graph->size_at(-3, out)); + return utils::div_up(M, uint64_t(4)) * utils::div_up(N4, uint64_t(2)) >= + kMinWideTileInvocations; +} + +} // namespace + void resize_conv2d_pw_tiled_node( ComputeGraph* graph, const std::vector& args, @@ -60,9 +88,16 @@ vkapi::ShaderInfo pick_conv2d_pw_tiled_shader( const std::vector& resize_args) { (void)resize_args; const ValueRef out = args.at(0).refs.at(0); + const ValueRef packed_weight = args.at(1).refs.at(1); std::string kernel_name = "conv2d_pw_tiled"; kernel_name.reserve(kShaderNameReserve); + if (use_wide_pw_tile(graph, out)) { + kernel_name += "_4x8"; + } + if (graph->is_buffer_storage(packed_weight)) { + kernel_name += "_bufw"; + } add_dtype_suffix(kernel_name, graph->dtype_of(out)); return VK_KERNEL_FROM_STR(kernel_name); } @@ -78,10 +113,64 @@ GlobalWorkGrid pick_conv2d_pw_tiled_gwg( uint32_t W = graph->size_at(-1, out); uint32_t H = graph->size_at(-2, out); uint32_t C_out = graph->size_at(-3, out); - uint32_t M = H * W; uint32_t N4 = utils::div_up_4(C_out); - // TILE_N4=1, TILE_M=2. Must match conv2d_pw_tiled.yaml. - return GlobalWorkGrid({N4, utils::div_up(M, 2u), 1u}, kTiledWorkGrid); + // Tile sizes must match conv2d_pw_tiled.yaml. + const bool wide = use_wide_pw_tile(graph, out); + const uint32_t tile_n4 = wide ? 2u : 1u; + const uint32_t tile_w = 2u; + const uint32_t tile_h = wide ? 2u : 1u; + const uint32_t n_tiles = utils::div_up(N4, tile_n4); + const uint32_t m_tiles = utils::div_up(W, tile_w) * utils::div_up(H, tile_h); + // The wide tile variant puts pixel tiles on x. + if (wide) { + return GlobalWorkGrid({m_tiles, n_tiles, 1u}, kTiledWorkGrid); + } + return GlobalWorkGrid({n_tiles, m_tiles, 1u}, kTiledWorkGrid); +} + +/* + * Wide tile work groups span the channel tiles of a row of pixel tiles, so the + * input they share is read through one core's cache: y covers the channel tiles + * (rounded up to a power of two) while x keeps at least min_x pixel tiles, and x + * takes the remaining threads. + * + * The thread count and min_x per GPU family come from sweeping 12 local sizes + * (32 to 512 threads with min_x 4 to 32, and an 8x8 square) over 57 Float and + * Half 1x1 convolutions with 7x7 to 1008x756 outputs on Mali, Adreno and + * Xclipse phones. On the large convolutions, Mali was 5-8% faster with 512 + * threads than with 128, Adreno was 1-10% slower with 512 than with 128, and + * Xclipse was 30% slower with 128 than with 64. + */ +LocalWorkGroup pick_conv2d_pw_tiled_lwg( + ComputeGraph* graph, + const vkapi::ShaderInfo& shader, + const GlobalWorkGrid& gwg, + const std::vector& args, + const std::vector& resize_args) { + if (use_wide_pw_tile(graph, args.at(0).refs.at(0))) { + const bool mali = graph->device_is_mali(); + const bool adreno = graph->device_is_adreno(); + uint32_t nthreads = mali ? 512u : adreno ? 128u : 64u; + const uint32_t min_x = mali || adreno ? 16u : 4u; + const vkapi::Adapter* const adapter = graph->context()->adapter_ptr(); + while (nthreads > adapter->max_compute_workgroup_invocations() || + nthreads > adapter->max_compute_workgroup_size()[0]) { + nthreads /= 2u; + } + // y stays at or below nthreads / min_x (32 at most), within the 128 that + // Vulkan guarantees for the second work group dimension. + uint32_t y = 1u; + while (y < gwg[1] && y * min_x < nthreads) { + y *= 2u; + } + return LocalWorkGroup(nthreads / y, y, 1u); + } + // On Mali, a work group of pixel tiles that share one channel tile lets each + // warp read a single weight stream instead of one per channel tile. + if (graph->device_is_mali() && gwg[1] >= 64u) { + return LocalWorkGroup(1u, 64u, 1u); + } + return pick_xy_square_lwg(graph, shader, gwg, args, resize_args); } // @@ -107,9 +196,16 @@ ValueRef prepack_conv2d_pw_weight( const int64_t output_height = K4; const int64_t output_width = N4 * 4 * 4; + // Buffer weights are faster than texture weights on every GPU measured. + // Half buffers need 16-bit storage support. Float uses texture weights on + // Mali, where Float buffer weights produced wrong results for padded and + // strided convolutions on Mali-G715. utils::StorageType weight_storage = utils::kTexture2D; uint32_t max_extent = graph.context()->adapter_ptr()->max_texture2d_dim(); - if (output_width / 4 > max_extent || + const bool buffer_weights = graph.dtype_of(weight_data) == vkapi::kHalf + ? graph.float16_buffers_enabled() + : !graph.device_is_mali(); + if (buffer_weights || output_width / 4 > max_extent || static_cast(output_height) > max_extent) { weight_storage = utils::kBuffer; } @@ -191,7 +287,7 @@ void add_conv2d_pw_tiled_node( graph, pick_conv2d_pw_tiled_shader, pick_conv2d_pw_tiled_gwg, - pick_xy_square_lwg, + pick_conv2d_pw_tiled_lwg, // Inputs and Outputs {{out, vkapi::kWrite}, {{in, packed_weight, packed_bias}, vkapi::kRead}}, // Shader params buffers diff --git a/backends/vulkan/test/ops/benchmark/cases/conv2d/conv2d.cpp b/backends/vulkan/test/ops/benchmark/cases/conv2d/conv2d.cpp index 20b945e9346..3d40d484833 100644 --- a/backends/vulkan/test/ops/benchmark/cases/conv2d/conv2d.cpp +++ b/backends/vulkan/test/ops/benchmark/cases/conv2d/conv2d.cpp @@ -645,8 +645,8 @@ void conv2d_pw_reference_impl(TestCase& test_case) { const ValueSpec& bias_spec = test_case.inputs()[2]; ValueSpec& output = test_case.outputs()[0]; - if (input.dtype != vkapi::kFloat) { - throw std::invalid_argument("Reference only supports float"); + if (input.dtype != vkapi::kFloat && input.dtype != vkapi::kHalf) { + throw std::invalid_argument("Reference only supports float and half"); } auto input_sizes = input.get_tensor_sizes(); @@ -658,8 +658,18 @@ void conv2d_pw_reference_impl(TestCase& test_case) { int64_t W = input_sizes[3]; int64_t C_out = weight_sizes[0]; - auto& input_data = input.get_float_data(); - auto& weight_data = weight.get_float_data(); + // Mirrors create_conv2d_pw_test_case's PERF classification: large-K half + // shapes stay timing-only, since FP16 accumulation error would exceed the + // half tolerance. + if (input.dtype == vkapi::kHalf && + (C_in > kRefDimSizeLimit || C_out > kRefDimSizeLimit || + H > kRefDimSizeLimit || W > kRefDimSizeLimit)) { + throw std::invalid_argument( + "Half reference skipped for PERF shape (timing-only)"); + } + + // get_element() materializes a float regardless of dtype, so the same loop + // body serves both FP32 and FP16. auto& ref_data = output.get_ref_float_data(); ref_data.resize(N * C_out * H * W, 0.0f); @@ -669,15 +679,14 @@ void conv2d_pw_reference_impl(TestCase& test_case) { for (int64_t w = 0; w < W; ++w) { float sum = 0.0f; for (int64_t ci = 0; ci < C_in; ++ci) { - float in_val = - input_data[n * (C_in * H * W) + ci * (H * W) + h * W + w]; + float in_val = input.get_element( + n * (C_in * H * W) + ci * (H * W) + h * W + w); // weight is [C_out, C_in, 1, 1] - float w_val = weight_data[co * C_in + ci]; + float w_val = weight.get_element(co * C_in + ci); sum += in_val * w_val; } if (!bias_spec.is_none()) { - auto& bias_data = bias_spec.get_float_data(); - sum += bias_data[co]; + sum += bias_spec.get_element(co); } ref_data[n * (C_out * H * W) + co * (H * W) + h * W + w] = sum; } diff --git a/backends/vulkan/test/ops/benchmark/cases/conv2d/test_sets/pw_sets.cpp b/backends/vulkan/test/ops/benchmark/cases/conv2d/test_sets/pw_sets.cpp index afb206ea6e2..d71ef21ba96 100644 --- a/backends/vulkan/test/ops/benchmark/cases/conv2d/test_sets/pw_sets.cpp +++ b/backends/vulkan/test/ops/benchmark/cases/conv2d/test_sets/pw_sets.cpp @@ -47,11 +47,14 @@ REGISTER_TEST_CASE_SET("conv2d", "pw_correctness") { // Non-multiple-of-4 channels {1, 13, 27, 8, 8, false}, {1, 33, 17, 8, 8, false}, + // Partial pixel and channel tiles, below and above the invocation count + // where Half switches to the wide tile + {1, 33, 52, 41, 47, true}, + {1, 33, 58, 63, 61, true}, }; const std::vector groups = { - // Accuracy test cases (float only) - {accuracy_configs, {vkapi::kFloat}}, + {accuracy_configs, {vkapi::kFloat, vkapi::kHalf}}, }; return { generate_pw_test_cases(groups), @@ -98,6 +101,28 @@ REGISTER_TEST_CASE_SET("conv2d", "pw_performance") { conv2d_pw_flop_calculator}; } +REGISTER_TEST_CASE_SET("conv2d", "dsr_pw_set") { + // DSR dim=64 1x1 convs at the 1512x2016 input resolution. + const std::vector configs = { + {1, 64, 64, 756, 1008, true}, + {1, 32, 128, 756, 1008, true}, + {1, 128, 32, 756, 1008, true}, + {1, 128, 128, 378, 504, true}, + {1, 64, 256, 378, 504, true}, + {1, 256, 64, 378, 504, true}, + {1, 256, 256, 189, 252, true}, + {1, 128, 512, 189, 252, true}, + {1, 512, 128, 189, 252, true}, + }; + // Full-resolution FP16 configurations are timing-only. + return { + generate_pw_test_cases({{configs, {vkapi::kHalf}}}), + nullptr, + conv2d_pw_flop_calculator, + 5, + 20}; +} + } // namespace conv2d } // namespace prototyping } // namespace vulkan From aeb7e9ed4eef7d280a376438e3f35d0138251090 Mon Sep 17 00:00:00 2001 From: Stephen Jia Date: Mon, 2 Nov 2026 09:05:17 -0800 Subject: [PATCH 2/2] Update [ghstack-poisoned] --- .../runtime/graph/ops/impl/convolution/conv2d/Conv2dPW.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/backends/vulkan/runtime/graph/ops/impl/convolution/conv2d/Conv2dPW.cpp b/backends/vulkan/runtime/graph/ops/impl/convolution/conv2d/Conv2dPW.cpp index 448cd2e7cbb..9f7e4bfc4db 100644 --- a/backends/vulkan/runtime/graph/ops/impl/convolution/conv2d/Conv2dPW.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/convolution/conv2d/Conv2dPW.cpp @@ -131,8 +131,8 @@ GlobalWorkGrid pick_conv2d_pw_tiled_gwg( /* * Wide tile work groups span the channel tiles of a row of pixel tiles, so the * input they share is read through one core's cache: y covers the channel tiles - * (rounded up to a power of two) while x keeps at least min_x pixel tiles, and x - * takes the remaining threads. + * (rounded up to a power of two) while x keeps at least min_x pixel tiles, and + * x takes the remaining threads. * * The thread count and min_x per GPU family come from sweeping 12 local sizes * (32 to 512 threads with min_x 4 to 32, and an 8x8 square) over 57 Float and