xe: ref_conv: align BWD_D post-ops with FWD

Author: echeresh

src/gpu/intel/ocl/ref_convolution.cl

@@ -191,26 +191,27 @@ __kernel void ref_convolution_bwd_data(__global SRC_DATA_T *diff_src,
         }
     }
 
-    ACC_DATA_T sum_src;
-#if WITH_SUM
-    sum_src = TO_ACC(SRC_TO_REF(diff_src[SRC_OFF(n, g * IC + ic, id, ih, iw)]));
-#endif
-
-    ACC_DATA_T accumulator = TO_ACC(d);
+    POST_OP_DATA_T tmp = d;
 
 #if WITH_SRC_SCALES
-    accumulator *= src_scales[0];
+    tmp *= src_scales[0];
 #endif
 #if WITH_WEI_SCALES
 #if WEI_SCALES_MASK == 0
-    accumulator *= wei_scales[0];
+    tmp *= wei_scales[0];
 #else
-    accumulator *= wei_scales[g * IC + ic];
+    tmp *= wei_scales[g * IC + ic];
 #endif
 #endif
 
 #if WITH_BIAS
-    accumulator += BIA_TO_REF(bias[g * IC + ic]);
+    tmp += (POST_OP_DATA_T)BIA_TO_REF(bias[g * IC + ic]);
+#endif
+
+    POST_OP_DATA_T sum_src;
+#if WITH_SUM
+    sum_src = (POST_OP_DATA_T)SUM_TO_REF(
+            AS_SUM_DATA_T(diff_src[SRC_OFF(n, g * IC + ic, id, ih, iw)]));
 #endif
 
 #if NDIMS == 3
@@ -230,19 +231,19 @@ __kernel void ref_convolution_bwd_data(__global SRC_DATA_T *diff_src,
     const unsigned po_d3 = 0;
     const unsigned po_d4 = 0;
 #endif
-    APPLY_POST_OPS_SERIAL(accumulator, ACC_DATA_T, sum_src, float, n, 1,
-            g *IC + ic, 1, po_d2, 1, po_d3, 1, po_d4, 1, 0, 1);
+    APPLY_POST_OPS_SERIAL(tmp, POST_OP_DATA_T, sum_src, POST_OP_DATA_T, n, 1,
+            g * IC + ic, 1, po_d2, 1, po_d3, 1, po_d4, 1, 0, 1);
 
 #if WITH_DST_SCALES
-    accumulator /= dst_scales[0];
+    tmp /= dst_scales[0];
 #endif
 
 #if WITH_DST_ZPOINTS
     const int dst_zp = dst_zpoints[WITH_DST_ZPOINTS_PER_OC ? g * IC + ic : 0];
-    accumulator += dst_zp;
+    tmp += dst_zp;
 #endif // WITH_DST_ZPOINTS
 
-    diff_src[SRC_OFF(n, g * IC + ic, id, ih, iw)] = TO_SRC(accumulator);
+    diff_src[SRC_OFF(n, g * IC + ic, id, ih, iw)] = TO_SRC(tmp);
 }
 #endif