generic: gpu: convolution/deconvolution/softmax: add missing checks

Author: t4c1

src/gpu/generic/sycl/ref_convolution.hpp

@@ -93,7 +93,8 @@ struct ref_convolution_fwd_t : public gpu::generic::sycl::primitive_t {
                             | sm::sum_dt)
                     && IMPLICATION(!attr()->scales_.has_default_values(),
                             attr_scales_ok())
-                    && sycl_post_ops_t::post_ops_ok(attr(), false);
+                    && sycl_post_ops_t::post_ops_ok(attr(), false)
+                    && set_default_alg_kind(alg_kind::convolution_direct);
             if (!ok) return status::unimplemented;
 
             return init_conf();
@@ -148,7 +149,8 @@ struct ref_convolution_bwd_data_t : public gpu::generic::sycl::primitive_t {
                     && attr()->has_default_values(sm::scales_runtime
                             | sm::zero_points_runtime | sm::sum_dt)
                     && IMPLICATION(!attr()->scales_.has_default_values(),
-                            attr_scales_ok());
+                            attr_scales_ok())
+                    && set_default_alg_kind(alg_kind::convolution_direct);
             if (!ok) return status::unimplemented;
 
             return init_conf();
@@ -200,10 +202,8 @@ struct ref_convolution_bwd_weights_t : public gpu::generic::sycl::primitive_t {
                             data_d, diff_weights_d, diff_dst_d)
                     && check_convolution_formats(
                             data_d, diff_weights_d, diff_dst_d)
-                    && attr()->has_default_values(sm::scales_runtime
-                            | sm::zero_points_runtime | sm::sum_dt)
-                    && IMPLICATION(!attr()->scales_.has_default_values(),
-                            attr_scales_ok());
+                    && attr()->has_default_values()
+                    && set_default_alg_kind(alg_kind::convolution_direct);
             if (!ok) return status::unimplemented;
 
             return init_conf();

src/gpu/generic/sycl/ref_deconvolution.hpp

@@ -57,9 +57,8 @@ struct ref_deconvolution_bwd_weights_t
                             data_d, diff_weights_d, diff_dst_d)
                     && check_convolution_formats(
                             data_d, diff_weights_d, diff_dst_d)
-                    && attr()->has_default_values(sm::scales_runtime
-                            | sm::zero_points_runtime | sm::post_ops
-                            | sm::sum_dt);
+                    && attr()->has_default_values()
+                    && desc()->alg_kind == alg_kind::deconvolution_direct;
             if (!ok) return status::unimplemented;
 
             return init_conf();

src/gpu/generic/sycl/ref_softmax.hpp

@@ -48,6 +48,7 @@ struct ref_sycl_softmax_fwd_t : public gpu::generic::sycl::primitive_t {
                     && sycl_post_ops_t::post_ops_ok(attr(), true, false)
                     && set_default_formats() == status::success
                     && attr_.set_default_formats(dst_md()) == status::success
+                    && check_formats(diff_src_md(), diff_dst_md())
                     && md_dims_in_range(src_md());
 
             if (!ok) return status::unimplemented;
@@ -70,6 +71,15 @@ struct ref_sycl_softmax_fwd_t : public gpu::generic::sycl::primitive_t {
             return utils::one_of(src, data_type::f32, data_type::bf16,
                     data_type::f16, data_type::s8, data_type::u8);
         }
+
+        static bool check_formats(const memory_desc_wrapper &src,
+                const memory_desc_wrapper &dst) {
+            for (const auto &mdw : {src, dst}) {
+                if (!mdw.is_plain()) return false;
+            }
+
+            return true;
+        }
     };
 
     status_t init(impl::engine_t *engine) override;
@@ -101,12 +111,22 @@ struct ref_sycl_softmax_bwd_t : public gpu::generic::sycl::primitive_t {
                     && dst_md()->data_type == diff_dst_md()->data_type
                     && attr()->has_default_values()
                     && set_default_formats() == status::success
+                    && check_formats(src_md(), dst_md())
                     && md_dims_in_range(diff_dst_md());
 
             if (!ok) return status::unimplemented;
             return init_conf();
         }
 
+        static bool check_formats(const memory_desc_wrapper &src,
+                const memory_desc_wrapper &dst) {
+            for (const auto &mdw : {src, dst}) {
+                if (!mdw.is_plain()) return false;
+            }
+
+            return true;
+        }
+
         sycl_softmax_conf_t conf_;
         status_t init_conf();
     };