diff --git a/src/gpu/ocl/ref_eltwise.cl b/src/gpu/ocl/ref_eltwise.cl
index 48b56121c14..b4ae7e6cb2b 100644
--- a/src/gpu/ocl/ref_eltwise.cl
+++ b/src/gpu/ocl/ref_eltwise.cl
@@ -1,5 +1,5 @@
 /*******************************************************************************
-* Copyright 2019-2021 Intel Corporation
+* Copyright 2019-2023 Intel Corporation
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
@@ -23,8 +23,6 @@
 #define DIFF_DATA_OFF(x0, x1, x2, x3, x4, x5) \
     OFF_MD(DIFF_DATA, x0, x1, x2, x3, x4, x5)
 
-#define KERNEL_ATTR __attribute__((intel_reqd_sub_group_size(SUB_GROUP_SIZE)))
-
 #if IS_FWD
 KERNEL_ATTR
 __kernel void ref_eltwise_fwd(__global DATA_T *src, __global DATA_T *dst,
diff --git a/src/gpu/ocl/ref_eltwise.cpp b/src/gpu/ocl/ref_eltwise.cpp
index d8c2b7a8ab2..7baeab1c763 100644
--- a/src/gpu/ocl/ref_eltwise.cpp
+++ b/src/gpu/ocl/ref_eltwise.cpp
@@ -41,7 +41,6 @@ static status_t init_conf_common(eltwise_conf_t &conf, offsets_t &off,
     conf.alg = alg;
     conf.is_forward = is_forward;
     conf.attr_info = attr_info_t::create(pd->attr());
-    conf.sub_group_size = 32;
 
     set_offsets(src_d, off.src_off);
     set_offsets(diff_data_d, off.dst_off);
@@ -78,7 +77,6 @@ static status_t init_kernel_ctx_common(compute::kernel_ctx_t &kernel_ctx,
     kernel_ctx.define_int("GWS0", conf.dispatch.nd_range().global_range()[0]);
     kernel_ctx.define_int("GWS1", conf.dispatch.nd_range().global_range()[1]);
     kernel_ctx.define_int("GWS2", conf.dispatch.nd_range().global_range()[2]);
-    kernel_ctx.define_int("SUB_GROUP_SIZE", conf.sub_group_size);
 
     bool with_binary_post_ops
             = post_ops.find(primitive_kind_t::dnnl_binary) != -1;
diff --git a/src/gpu/ocl/ref_eltwise.hpp b/src/gpu/ocl/ref_eltwise.hpp
index 1b13eb3b271..26dd47a7191 100644
--- a/src/gpu/ocl/ref_eltwise.hpp
+++ b/src/gpu/ocl/ref_eltwise.hpp
@@ -62,8 +62,6 @@ struct ref_eltwise_fwd_t : public gpu_primitive_t {
             if (!ok) return status::unimplemented;
 
             CHECK(init_conf(engine));
-            if (!compute_engine->mayiuse_sub_group(conf.sub_group_size))
-                return status::unimplemented;
             return status::success;
         }
 
@@ -110,9 +108,6 @@ struct ref_eltwise_bwd_t : public gpu_primitive_t {
             using namespace utils;
             assert(engine->kind() == engine_kind::gpu);
 
-            auto *compute_engine
-                    = utils::downcast<compute::compute_engine_t *>(engine);
-
             using namespace alg_kind;
             const bool ok = !is_fwd()
                     && !memory_desc_ndims_ok(data_md(), diff_dst_md())
@@ -127,8 +122,6 @@ struct ref_eltwise_bwd_t : public gpu_primitive_t {
             if (!ok) return status::unimplemented;
 
             CHECK(init_conf(engine));
-            if (!compute_engine->mayiuse_sub_group(conf.sub_group_size))
-                return status::unimplemented;
             return status::success;
         }