Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

gpu: sycl: pooling: Reduce argument size #1951

Merged
merged 1 commit into from
Jul 12, 2024
Merged

gpu: sycl: pooling: Reduce argument size #1951

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
15 changes: 10 additions & 5 deletions src/gpu/generic/sycl/pooling_kernels.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,7 +35,7 @@ namespace sycl {

using namespace nstl;
struct pooling_fwd_kernel_vec_t {
pooling_fwd_kernel_vec_t(const sycl_pooling_conf_t &conf,
pooling_fwd_kernel_vec_t(const sycl_pooling_fwd_conf_t &conf,
xpu::sycl::in_memory_arg_t &src, xpu::sycl::out_memory_arg_t &dst,
xpu::sycl::out_memory_arg_t &ws, xpu::sycl::in_memory_arg_t &src_1,
xpu::sycl::in_memory_arg_t &src_2,
Expand Down Expand Up @@ -148,13 +148,18 @@ struct pooling_fwd_kernel_vec_t {
float dst_Value(xpu::sycl::in_memory_arg_t arr, int idx, int offset) const {
auto src1_desc = conf_.src1_md[idx];
dim_t src_dim[DNNL_MAX_NDIMS];
dim_t dst_dim[DNNL_MAX_NDIMS];
auto src_dim_ = src1_desc.dims();
auto dst_dim_ = dst_md().dims();

for (int j = 0; j < src1_desc.ndims(); j++) {
src_dim[j] = src_dim_[j];
}
for (int j = 0; j < dst_md().ndims(); j++) {
dst_dim[j] = dst_dim_[j];
}
const auto off = get_binary_src1_off(
src1_desc, src_dim, offset, conf_.dst_dims, conf_.dst_ndims);
src1_desc, src_dim, offset, dst_dim, conf_.dst_md.ndims());
auto dst = load_float_value(src1_desc.data_type(), gen_ptr(arr), off);
return dst;
}
Expand Down Expand Up @@ -289,7 +294,7 @@ struct pooling_fwd_kernel_vec_t {
d /= num_summands;
}

sycl_pooling_conf_t conf_;
sycl_pooling_fwd_conf_t conf_;

xpu::sycl::in_memory_arg_t src_;
xpu::sycl::out_memory_arg_t dst_;
Expand All @@ -302,7 +307,7 @@ struct pooling_fwd_kernel_vec_t {
};

struct pooling_bwd_kernel_vec_t {
pooling_bwd_kernel_vec_t(const sycl_pooling_conf_t &conf,
pooling_bwd_kernel_vec_t(const sycl_pooling_bwd_conf_t &conf,
xpu::sycl::in_memory_arg_t &diff_dst,
xpu::sycl::out_memory_arg_t &diff_src,
xpu::sycl::in_memory_arg_t &ws)
Expand Down Expand Up @@ -479,7 +484,7 @@ struct pooling_bwd_kernel_vec_t {
}
}

sycl_pooling_conf_t conf_;
sycl_pooling_bwd_conf_t conf_;
xpu::sycl::in_memory_arg_t diff_dst_;
xpu::sycl::out_memory_arg_t diff_src_;
xpu::sycl::in_memory_arg_t ws_;
Expand Down
11 changes: 5 additions & 6 deletions src/gpu/generic/sycl/ref_pooling.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,7 @@ namespace sycl {

status_t ref_pooling_fwd_t::pd_t::init_conf() {
using namespace primitive_kind;
conf_ = sycl_pooling_conf_t();
conf_ = sycl_pooling_fwd_conf_t();
conf_.ndims = ndims();
conf_.block_size = 16;
conf_.wg_size = 32;
Expand All @@ -38,10 +38,6 @@ status_t ref_pooling_fwd_t::pd_t::init_conf() {
? xpu::sycl::md_t(workspace_md(0))
: xpu::sycl::md_t {};
conf_.zero_dims = has_zero_dim_memory();
for (int i = 0; i < DNNL_MAX_NDIMS; i++) {
conf_.dst_dims[i] = dst_md()->dims[i];
}
conf_.dst_ndims = dst_md()->ndims;
auto nelems_A = memory_desc_wrapper(src_md(0)).nelems();
int work_per_wg = conf_.wg_size * conf_.block_size;
int n_wgs = (nelems_A + work_per_wg - 1) / work_per_wg;
Expand Down Expand Up @@ -70,6 +66,9 @@ status_t ref_pooling_fwd_t::pd_t::init_conf() {

const auto *att = attr();
const auto &attr_po = att->post_ops_;
if (attr_po.len() > sycl_post_ops_t::max_post_ops) {
return dnnl_unimplemented;
}
conf_.po_len = attr_po.len();
for (auto i = 0; i < attr_po.len(); ++i) {
if (attr_po.contain(binary, i)) {
Expand Down Expand Up @@ -117,7 +116,7 @@ status_t ref_pooling_fwd_t::execute_forward(const exec_ctx_t &ctx) const {
}

status_t ref_pooling_bwd_t::pd_t::init_conf() {
conf_ = sycl_pooling_conf_t();
conf_ = sycl_pooling_bwd_conf_t();
conf_.ndims = ndims();
conf_.block_size = 16;
conf_.wg_size = 32;
Expand Down
4 changes: 2 additions & 2 deletions src/gpu/generic/sycl/ref_pooling.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -75,7 +75,7 @@ struct ref_pooling_fwd_t : public gpu::generic::sycl::primitive_t {
}

status_t init_conf();
sycl_pooling_conf_t conf_;
sycl_pooling_fwd_conf_t conf_;
};

status_t init(impl::engine_t *engine) override;
Expand Down Expand Up @@ -125,7 +125,7 @@ struct ref_pooling_bwd_t : public gpu::generic::sycl::primitive_t {
}

status_t init_conf();
sycl_pooling_conf_t conf_;
sycl_pooling_bwd_conf_t conf_;
};

status_t init(impl::engine_t *engine) override;
Expand Down
25 changes: 15 additions & 10 deletions src/gpu/generic/sycl/sycl_primitive_conf.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -301,14 +301,8 @@ struct sycl_lrn_conf_t {
int wk_size;
};

struct sycl_pooling_conf_t {
xpu::sycl::md_t src_md;
// The size "5" is lower than DNNL_MAX_NDIMS because only 5 dimension formats are supported.
xpu::sycl::md_t src1_md[5];
xpu::sycl::md_t dst_md;
struct sycl_pooling_base_conf_t {
xpu::sycl::md_t ws_md;
xpu::sycl::md_t diff_src_md;
xpu::sycl::md_t diff_dst_md;
int ndims;
int po_len;
bool zero_dims;
Expand Down Expand Up @@ -336,11 +330,20 @@ struct sycl_pooling_conf_t {
dim_t DD;
dim_t DH;
dim_t DW;
dims_t dst_dims;
int dst_ndims;
};

struct sycl_pooling_fwd_conf_t : public sycl_pooling_base_conf_t {
xpu::sycl::md_t src_md;
xpu::sycl::md_t src1_md[sycl_post_ops_t::max_post_ops];
xpu::sycl::md_t dst_md;
sycl_post_ops_t post_ops;
};

struct sycl_pooling_bwd_conf_t : public sycl_pooling_base_conf_t {
xpu::sycl::md_t diff_src_md;
xpu::sycl::md_t diff_dst_md;
};

CHECK_SYCL_KERNEL_ARG_TYPE(sycl_binary_conf_t);
CHECK_SYCL_KERNEL_ARG_TYPE(sycl_prelu_conf_t);
CHECK_SYCL_KERNEL_ARG_TYPE(sycl_shuffle_conf_t);
Expand All @@ -350,7 +353,9 @@ CHECK_SYCL_KERNEL_ARG_TYPE(sycl_softmax_conf_t);
CHECK_SYCL_KERNEL_ARG_TYPE(sycl_layer_normalization_conf_t);
CHECK_SYCL_KERNEL_ARG_TYPE(sycl_eltwise_conf_t);
CHECK_SYCL_KERNEL_ARG_TYPE(sycl_lrn_conf_t);
CHECK_SYCL_KERNEL_ARG_TYPE(sycl_pooling_conf_t);
CHECK_SYCL_KERNEL_ARG_TYPE(sycl_pooling_base_conf_t);
CHECK_SYCL_KERNEL_ARG_TYPE(sycl_pooling_fwd_conf_t);
CHECK_SYCL_KERNEL_ARG_TYPE(sycl_pooling_bwd_conf_t);

} // namespace sycl
} // namespace generic
Expand Down
Loading