generic: sycl: pooling: buf fixes when running on Intel GPU

Author: kala855

src/gpu/generic/sycl/pooling_kernels.hpp

@@ -124,7 +124,7 @@ struct pooling_fwd_kernel_vec_t {
         return 0;
     }
     float data_conv() const {
-        switch (src_md().data_type()) {
+        switch (dst_md().data_type()) {
             case data_type::bf16:
                 return (float)
                         std::numeric_limits<xpu::sycl::bfloat16_t>::lowest();

tests/benchdnn/pool/pool.cpp

@@ -162,15 +162,15 @@ void skip_invalid_prb(const prb_t *prb, res_t *res) {
 // Special function to handle Nvidia libraries issues showing up through the
 // timeline. Not recommended to remove instances to keep working state of any
 // cuda/cuDNN/cuBLAS versions.
-bool cuda_check_correctness(const prb_t *prb,
+bool gpu_check_correctness(const prb_t *prb,
         const compare::compare_t::driver_check_func_args_t &args) {
-    if (!is_nvidia_gpu()) return false;
+    if (!is_gpu()) return false;
 
-    if (args.dt == dnnl_f16) {
+    if (args.dt == dnnl_f16 && is_nvidia_gpu()) {
         // cuDNN bug: it spits f16 min value as -inf, not -65504.
         return args.exp == lowest_dt(args.dt) && std::isinf(args.got)
                 && std::signbit(args.got);
-    } else if (args.dt == dnnl_s8) {
+    } else if (args.dt == dnnl_s8 && is_nvidia_gpu()) {
         // cuDNN bug: ... and s8 min value as -127 (-INT8_MAX?), not -128.
         return args.exp == lowest_dt(args.dt) && args.got == -127;
     } else if (prb->alg == alg_t::max && (prb->dir & FLAG_BWD)
@@ -203,7 +203,7 @@ void setup_cmp(compare::compare_t &cmp, const prb_t *prb, data_kind_t kind,
     // and `kind` by value to avoid using dangling references.
     const auto pooling_add_check =
             [&, prb](const compare::compare_t::driver_check_func_args_t &args) {
-                return cuda_check_correctness(prb, args);
+                return gpu_check_correctness(prb, args);
             };
     cmp.set_driver_check_function(pooling_add_check);
 }

tests/benchdnn/pool/ref_pool.cpp

@@ -35,11 +35,7 @@ void compute_ref_fwd(const prb_t *prb, const args_t &args) {
         const int64_t SD = prb->sd, SH = prb->sh, SW = prb->sw;
         const int64_t DD = prb->dd, DH = prb->dh, DW = prb->dw;
 
-        // XXX: this is a hack to let tests with padded area to pass for bf16
-        // dt due to the library initialize values with -max_dt, but not -INF.
         float max_value = lowest_dt(prb->dst_dt());
-        if (is_nvidia_gpu() || is_amd_gpu())
-            max_value = lowest_dt(prb->src_dt());
         float avg_value = 0.;
         // Set initial value based on ws data type
         int ws_off = prb->kernel_size() <= UINT8_MAX ? UINT8_MAX : INT_MAX;