COMPMID-1269: (Nightly) Fix CL/Winograd/OutputTransform mismatches

Check if the depth is multiple of tile size for NHWC if not write to
dummy padding.

Change-Id: Ie854dcbc75aa94bd1686f7769a009dd2654fdfed
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/135055
Reviewed-by: Pablo Tello <pablo.tello@arm.com>
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>

3 files changed, 49 insertions, 28 deletions

diff --git a/src/core/CL/cl_kernels/winograd.cl b/src/core/CL/cl_kernels/winograd.cl
index 6a570277ab..c7ca8f6752 100644
--- a/src/core/CL/cl_kernels/winograd.cl
+++ b/src/core/CL/cl_kernels/winograd.cl
@@ -1586,15 +1586,15 @@ __kernel void winograd_output_transform_4x4_3x3_nchw(
  * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
  * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in]  dst_size                          Size of the destination tensor, minus the last padding
  */
 __kernel void winograd_output_transform_4x4_3x3_nhwc(
     TENSOR3D_DECLARATION(src),
-    TENSOR3D_DECLARATION(dst)
+    TENSOR3D_DECLARATION(dst),
 #if defined(HAS_BIAS)
-    ,
-    VECTOR_DECLARATION(bias)
+    VECTOR_DECLARATION(bias),
 #endif // defined(HAS_BIAS)
-)
+    int dst_size)
 {
     // Each thread stores a 4x4 tile
     Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
@@ -1734,25 +1734,27 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc(
 #endif // defined(HAS_BIAS)
 
     // Get output address
-    __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_out * dst_stride_x + y_out * dst_stride_y + z_out * dst_stride_z;
+    int4 offset = (int4)(dst_offset_first_element_in_bytes + x_out * sizeof(float) + y_out * dst_stride_y + z_out * dst_stride_z);
+    offset      = min(offset + (int4)(0, 1, 2, 3) * (int4)dst_stride_z, dst_size); // If address is beyond the last plane, clamp it to dst_size (which points to the last padding).
+    int4 mult_y = min(dst_size - offset, 1);                                       // If out of bound, we don't want to increase dst_stride_y, so we set the multiplier to 0. It will be 1 otherwise.
 
     // Store the 4x4 output tile
-    *((__global float *)(dst_addr + 0 * dst_stride_y + 0 * dst_stride_z)) = out00;
-    *((__global float *)(dst_addr + 1 * dst_stride_y + 0 * dst_stride_z)) = out01;
-    *((__global float *)(dst_addr + 2 * dst_stride_y + 0 * dst_stride_z)) = out02;
-    *((__global float *)(dst_addr + 3 * dst_stride_y + 0 * dst_stride_z)) = out03;
-    *((__global float *)(dst_addr + 0 * dst_stride_y + 1 * dst_stride_z)) = out10;
-    *((__global float *)(dst_addr + 1 * dst_stride_y + 1 * dst_stride_z)) = out11;
-    *((__global float *)(dst_addr + 2 * dst_stride_y + 1 * dst_stride_z)) = out12;
-    *((__global float *)(dst_addr + 3 * dst_stride_y + 1 * dst_stride_z)) = out13;
-    *((__global float *)(dst_addr + 0 * dst_stride_y + 2 * dst_stride_z)) = out20;
-    *((__global float *)(dst_addr + 1 * dst_stride_y + 2 * dst_stride_z)) = out21;
-    *((__global float *)(dst_addr + 2 * dst_stride_y + 2 * dst_stride_z)) = out22;
-    *((__global float *)(dst_addr + 3 * dst_stride_y + 2 * dst_stride_z)) = out23;
-    *((__global float *)(dst_addr + 0 * dst_stride_y + 3 * dst_stride_z)) = out30;
-    *((__global float *)(dst_addr + 1 * dst_stride_y + 3 * dst_stride_z)) = out31;
-    *((__global float *)(dst_addr + 2 * dst_stride_y + 3 * dst_stride_z)) = out32;
-    *((__global float *)(dst_addr + 3 * dst_stride_y + 3 * dst_stride_z)) = out33;
+    *((__global float *)(dst_ptr + mult_y.s0 * 0 * dst_stride_y + offset.s0)) = out00;
+    *((__global float *)(dst_ptr + mult_y.s0 * 1 * dst_stride_y + offset.s0)) = out01;
+    *((__global float *)(dst_ptr + mult_y.s0 * 2 * dst_stride_y + offset.s0)) = out02;
+    *((__global float *)(dst_ptr + mult_y.s0 * 3 * dst_stride_y + offset.s0)) = out03;
+    *((__global float *)(dst_ptr + mult_y.s1 * 0 * dst_stride_y + offset.s1)) = out10;
+    *((__global float *)(dst_ptr + mult_y.s1 * 1 * dst_stride_y + offset.s1)) = out11;
+    *((__global float *)(dst_ptr + mult_y.s1 * 2 * dst_stride_y + offset.s1)) = out12;
+    *((__global float *)(dst_ptr + mult_y.s1 * 3 * dst_stride_y + offset.s1)) = out13;
+    *((__global float *)(dst_ptr + mult_y.s2 * 0 * dst_stride_y + offset.s2)) = out20;
+    *((__global float *)(dst_ptr + mult_y.s2 * 1 * dst_stride_y + offset.s2)) = out21;
+    *((__global float *)(dst_ptr + mult_y.s2 * 2 * dst_stride_y + offset.s2)) = out22;
+    *((__global float *)(dst_ptr + mult_y.s2 * 3 * dst_stride_y + offset.s2)) = out23;
+    *((__global float *)(dst_ptr + mult_y.s3 * 0 * dst_stride_y + offset.s3)) = out30;
+    *((__global float *)(dst_ptr + mult_y.s3 * 1 * dst_stride_y + offset.s3)) = out31;
+    *((__global float *)(dst_ptr + mult_y.s3 * 2 * dst_stride_y + offset.s3)) = out32;
+    *((__global float *)(dst_ptr + mult_y.s3 * 3 * dst_stride_y + offset.s3)) = out33;
 }
 
 #define COMPUTE_TMP_COL(col, d0, d1, d2, d3, d4, d5, d6, d7, comm_fact)  \
diff --git a/src/core/CL/kernels/CLWinogradOutputTransformKernel.cpp b/src/core/CL/kernels/CLWinogradOutputTransformKernel.cpp
index 416d8e8d5f..5377bd33d2 100644
--- a/src/core/CL/kernels/CLWinogradOutputTransformKernel.cpp
+++ b/src/core/CL/kernels/CLWinogradOutputTransformKernel.cpp
@@ -49,6 +49,8 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *bias, con
 {
     ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32);
 
+    ARM_COMPUTE_RETURN_ERROR_ON(output->data_layout() != winograd_info.output_data_layout);
+
     const PadStrideInfo conv_info        = winograd_info.convolution_info;
     const Size2D        output_tile_size = winograd_info.output_tile_size;
     const Size2D        kernel_size      = winograd_info.kernel_size;
@@ -94,20 +96,31 @@ std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITen
     Window win            = calculate_max_window(*input, Steps(num_elems_processed_per_iteration));
     bool   window_changed = false;
 
-    AccessWindowRectangle input_access(input, 0, 0, num_elems_processed_per_iteration, num_elems_processed_per_iteration);
-    AccessWindowStatic    output_access(output, 0, 0, ceil_to_multiple(output->dimension(0), output_tile_size.width), ceil_to_multiple(output->dimension(1), output_tile_size.height));
+    int output_static_window_end_x = 0;
+    int output_static_window_end_y = 0;
 
-    if(bias != nullptr)
+    if(output->data_layout() == DataLayout::NCHW)
     {
-        AccessWindowStatic bias_access(bias, 0, 0, bias->dimension(0), bias->dimension(1));
-        window_changed = update_window_and_padding(win, input_access, bias_access, output_access);
+        output_static_window_end_x = ceil_to_multiple(output->dimension(0), output_tile_size.width);
+        output_static_window_end_y = ceil_to_multiple(output->dimension(1), output_tile_size.height);
     }
     else
     {
-        window_changed = update_window_and_padding(win, input_access, output_access);
+        output_static_window_end_x = output->dimension(0);
+        output_static_window_end_y = std::max(ceil_to_multiple(output->dimension(1), output_tile_size.width), output->dimension(1) + 1 /* For out of bound reads towards the z axis */);
     }
+
+    AccessWindowRectangle input_access(input, 0, 0, num_elems_processed_per_iteration, num_elems_processed_per_iteration);
+    AccessWindowStatic    output_access(output, 0, 0, output_static_window_end_x, output_static_window_end_y);
+    window_changed = update_window_and_padding(win, input_access, output_access);
     output->set_valid_region(ValidRegion(Coordinates(), output->tensor_shape()));
 
+    if(bias != nullptr)
+    {
+        AccessWindowStatic bias_access(bias, 0, 0, bias->dimension(0), bias->dimension(1));
+        window_changed = window_changed || update_window_and_padding(win, bias_access);
+    }
+
     Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
     return std::make_pair(err, win);
 }
@@ -197,6 +210,12 @@ void CLWinogradOutputTransformKernel::run(const Window &window, cl::CommandQueue
         add_1D_tensor_argument(idx1, _bias, slice_biases);
     }
 
+    if(_output->info()->data_layout() == DataLayout::NHWC)
+    {
+        unsigned int idx2 = 2 * num_arguments_per_3D_tensor() + ((_bias != nullptr) ? num_arguments_per_1D_tensor() : 0);
+        _kernel.setArg(idx2, static_cast<int>(_output->info()->total_size() - _output->info()->strides_in_bytes().y()));
+    }
+
     do
     {
         unsigned int idx = 0;
diff --git a/tests/validation/CL/Winograd.cpp b/tests/validation/CL/Winograd.cpp
index 17dd58ffdc..4565c2d59d 100644
--- a/tests/validation/CL/Winograd.cpp
+++ b/tests/validation/CL/Winograd.cpp
@@ -297,7 +297,7 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(framework::da
 
     // Create tensors
     CLTensor a = create_tensor<CLTensor>(shape_a, data_type);
-    CLTensor b = create_tensor<CLTensor>(shape_b, data_type);
+    CLTensor b = create_tensor<CLTensor>(shape_b, data_type, 1, 0, QuantizationInfo(), winograd_info.output_data_layout);
 
     ARM_COMPUTE_EXPECT(a.info()->is_resizable(), framework::LogLevel::ERRORS);
     ARM_COMPUTE_EXPECT(b.info()->is_resizable(), framework::LogLevel::ERRORS);