gpu: ocl: binary: enable mixed data types for gen9 kernels

oneapi-src · Oct 9, 2020 · ac0ae4b · ac0ae4b
1 parent 2f90d44
commit ac0ae4b
Show file tree

Hide file tree

Showing 2 changed files with 4 additions and 5 deletions.
diff --git a/src/gpu/ocl/gen9_binary.cl b/src/gpu/ocl/gen9_binary.cl
@@ -340,13 +340,13 @@ __kernel void gen9_binary(__global SRC0_DATA_T *src0,
 
 #if WITH_SUM
 #if NVECT == 1
-    dst_data = CONVERT_FLOAT_T(DST_BLOCK_READ(&src0[0]));
+    dst_data = CONVERT_FLOAT_T(DST_BLOCK_READ(&dst[0]));
 #elif NVECT == 2
-    dst_data = CONVERT_FLOAT2_T(DST_BLOCK_READ2(&src0[0]));
+    dst_data = CONVERT_FLOAT2_T(DST_BLOCK_READ2(&dst[0]));
 #elif NVECT == 4
-    dst_data = CONVERT_FLOAT4_T(DST_BLOCK_READ4(&src0[0]));
+    dst_data = CONVERT_FLOAT4_T(DST_BLOCK_READ4(&dst[0]));
 #elif NVECT == 8
-    dst_data = CONVERT_FLOAT8_T(DST_BLOCK_READ8(&src0[0]));
+    dst_data = CONVERT_FLOAT8_T(DST_BLOCK_READ8(&dst[0]));
 #endif
 #endif
 

diff --git a/src/gpu/ocl/gen9_binary.hpp b/src/gpu/ocl/gen9_binary.hpp
@@ -55,7 +55,6 @@ struct gen9_binary_t : public gpu_primitive_t {
                             || utils::everyone_is(f16, src_md(0)->data_type,
                                     src_md(1)->data_type, dst_md()->data_type)
                             || utils::one_of(src_md(0)->data_type, s8, u8))
-                    && (src_md(0)->data_type == dst_md()->data_type)
                     && IMPLICATION(!attr()->scales_.has_default_values(),
                             utils::one_of(src_md(0)->data_type, s8, u8)
                                     && utils::one_of(