From 19dba1eff66e95c318ce368299229f99a87ba1f5 Mon Sep 17 00:00:00 2001 From: Stephen Jia Date: Mon, 2 Nov 2026 13:30:23 -0800 Subject: [PATCH] [ET-VK][conv2d] Use a 2x2x8 tile and tuned local sizes for pointwise convolution Pull Request resolved: https://github.com/pytorch/executorch/pull/24433 ## Ulterior Motive Closing the latency gap between ET-VK and LiteRT (TFLite GPU) on the DSR denoising model on mobile GPUs. The model's 80 1x1 convolutions were the largest per-operator gap on Mali. ## Rationale **What**: Reworks `conv2d_pw_tiled` around a 2x2-pixel x 8-channel tile with buffer weights, branch-free bounds handling and per-GPU local sizes. The 2x1-pixel x 4-channel tile stays for small Half grids and for Float on Mali, and Float on Mali uses texture weights. **Why**: Pointwise convolutions are load/store bound on Mali. The old tile issued 6 texel loads per 8 vec4 FMAs, read weights through a texture, and guarded every load with a bounds check. The 2x2 x 8 tile issues half as many loads per FMA. ## Details - Tile - **Before**: one tile, 2 consecutive output positions x 4 output channels per invocation. - **After**: a 2x2 output pixel block x 8 output channels (`_4x8` variants) for Float on every GPU except Mali, and for Half when the grid has at least 4096 invocations. Smaller Half grids keep the 2x1 x 4 tile, which fills the GPU better there, and so does Float on Mali. - Weight storage - **Before**: texture2d, unless the packed weights exceed texture limits. - **After**: buffer (`_bufw` variants) for Float on every GPU except Mali, and for Half when the device supports 16-bit storage buffers; texture2d otherwise. On a Mali-G715 in CI, Float buffer weights produced wrong results for padded and strided convolutions with either tile. The same cases passed with the parent's texture weights. - Bounds handling - **Before**: every input and weight load is guarded. - **After**: pixels and channel tiles past the edge read the nearest in-range address and are discarded at the store, so weight loads have no branches. One redundant K check stays on the input load; without any branch in the K loop, the Adreno compiler reads the weight buffer through a much slower path. - Local size - **Before**: 64-invocation square for every dispatch. - **After**: wide tile work groups span the channel tiles of a row of pixel tiles, with 512 threads on Mali, 128 on Adreno and 64 elsewhere, halved until they fit the device limits. The narrow tile uses 1x64 along pixel tiles on Mali when there are at least 64 of them, and the 64-invocation square otherwise. - Tests - The pointwise reference supports Half. `pw_correctness` runs Half too, and adds partial-tile shapes below and above the Half tile switch. - New timing-only `dsr_pw_set` with the DSR model's 1x1 convolution shapes. The tile switch and local sizes come from a sweep over 12 local sizes and both tiles on 57 shapes (see Test Plan). Float on Mali keeps the 2x1 x 4 tile: in the Mali offline compiler the 2x2 x 8 Float tile spills 88 bytes per thread while the 2x1 x 4 Float tile does not spill, and the wide Float tile was 2-31% slower in the sweep. ghstack-source-id: 462948223 @exported-using-ghexport Differential Revision: [D128786678](https://our.internmc.facebook.com/intern/diff/D128786678/) --- .../convolution/conv2d/conv2d_pw_tiled.glsl | 112 +++++++++++------- .../convolution/conv2d/conv2d_pw_tiled.yaml | 16 +++ .../linear_fp_packed_weight_tile_load.glslh | 22 ++++ .../ops/impl/convolution/conv2d/Conv2dPW.cpp | 103 +++++++++++++++- .../ops/benchmark/cases/conv2d/conv2d.cpp | 27 +++-- .../cases/conv2d/test_sets/pw_sets.cpp | 29 ++++- 6 files changed, 249 insertions(+), 60 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..b703d418612 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,24 @@ #version 450 core #define PRECISION ${PRECISION} -#define VEC4_T ${texel_load_type(DTYPE, STORAGE)} +#define T ${buffer_scalar_type(DTYPE)} +#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 +34,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,46 +62,52 @@ ${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]; + // 1 for pixels inside the input, 0 for padding. + T in_scale[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 W_in, - const int H_in) { + const int H_out, + const ivec2 in_extents) { + 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); - coords.valid[m] = true; + coords.in_scale[m] = T(1); } else { - const int in_x = out_x * stride_w - padding_w; - const int in_y = out_y * stride_h - padding_h; - coords.pos[m] = ivec2(in_x, in_y); - coords.valid[m] = in_x >= 0 && in_x < W_in && in_y >= 0 && in_y < H_in; + // Padding pixels read the nearest input pixel and are zeroed by a + // multiply. Zeroing them with a branch or select instead dropped valid + // pixels on Adreno 830 (Float) and PowerVR DXT (Half). + const ivec2 pos = + ivec2(out_x * stride_w - padding_w, out_y * stride_h - padding_h); + coords.pos[m] = clamp(pos, ivec2(0), in_extents - 1); + coords.in_scale[m] = T(all(equal(coords.pos[m], pos)) ? 1.0 : 0.0); } } } @@ -102,9 +119,15 @@ 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) { 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)); + if (stride_1_padding_0 == 0) { + tile.data[m][k4] *= coords.in_scale[m]; + } } else { tile.data[m][k4] = VEC4_T(0.0); } @@ -116,45 +139,45 @@ 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; const int K = in_sizes.z; const int K4 = div_up_4(K); 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 +188,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, in_sizes.xy); 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 +202,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..81f624a891a 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,13 @@ 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. 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(); + if (buffer_weights || output_width / 4 > max_extent || static_cast(output_height) > max_extent) { weight_storage = utils::kBuffer; } @@ -191,7 +284,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