Skip to content

Commit b9fb939

Browse files
sgeor255mgouicem
sgeor255
authored and
mgouicem
committed
generic: fix several gtest issues
1 parent 2bf5ffc commit b9fb939

40 files changed

+416
-13
lines changed

‎src/gpu/generic/ref_concat.hpp

+3
Original file line numberDiff line numberDiff line change
@@ -129,6 +129,9 @@ struct ref_concat_t : public gpu::primitive_t {
129129
}
130130

131131
status_t execute(const exec_ctx_t &ctx) const override {
132+
if (memory_desc_wrapper(pd()->dst_md()).size() == 0)
133+
return status::success;
134+
132135
using namespace memory_tracking::names;
133136
impl::engine_t *engine = ctx.stream()->engine();
134137
const auto n = pd()->n_inputs();

‎src/gpu/generic/sycl/ref_binary.cpp

+3
Original file line numberDiff line numberDiff line change
@@ -55,12 +55,15 @@ status_t ref_binary_t::pd_t::init_conf() {
5555
}
5656

5757
status_t ref_binary_t::init(impl::engine_t *engine) {
58+
if (memory_desc_wrapper(pd()->dst_md()).size() == 0) return status::success;
59+
5860
const auto kid = ::sycl::get_kernel_id<binary_kernel_vec_t>();
5961
CHECK(create_kernel(engine, kid, &kernel_));
6062
return status::success;
6163
}
6264

6365
status_t ref_binary_t::execute(const exec_ctx_t &ctx) const {
66+
if (memory_desc_wrapper(pd()->dst_md()).size() == 0) return status::success;
6467

6568
ctx.zero_pad_output(DNNL_ARG_TO);
6669

‎src/gpu/generic/sycl/ref_convolution.cpp

+6-1
Original file line numberDiff line numberDiff line change
@@ -74,6 +74,8 @@ status_t ref_convolution_fwd_t::init(impl::engine_t *engine) {
7474
}
7575

7676
status_t ref_convolution_fwd_t::execute(const exec_ctx_t &ctx) const {
77+
if (memory_desc_wrapper(pd()->dst_md()).size() == 0) return status::success;
78+
7779
parallel_for(ctx, kernel_, [&](::sycl::handler &cgh) {
7880
convolution_kernel_fwd_t convolution_kernel(pd()->conf_, cgh, ctx);
7981

@@ -111,7 +113,7 @@ status_t ref_convolution_bwd_data_t::pd_t::init_conf() {
111113
conf_.single_data_zeropoint = attr()->zero_points_.common(DNNL_ARG_SRC_0);
112114
conf_.single_dst_zeropoint = attr()->zero_points_.common(DNNL_ARG_DST);
113115

114-
conf_.post_ops = sycl_post_ops_t(attr(), dst_md());
116+
conf_.post_ops = sycl_post_ops_t(attr(), diff_src_md());
115117

116118
conf_.padding[0] = static_cast<int>(desc()->padding[0][0]);
117119
conf_.padding[1] = static_cast<int>(desc()->padding[0][1]);
@@ -134,6 +136,9 @@ status_t ref_convolution_bwd_data_t::init(impl::engine_t *engine) {
134136
}
135137

136138
status_t ref_convolution_bwd_data_t::execute(const exec_ctx_t &ctx) const {
139+
if (memory_desc_wrapper(pd()->diff_src_md()).size() == 0)
140+
return status::success;
141+
137142
parallel_for(ctx, kernel_, [&](::sycl::handler &cgh) {
138143
convolution_kernel_bwd_data_t convolution_kernel(pd()->conf_, cgh, ctx);
139144

‎src/gpu/generic/sycl/ref_layer_normalizations.cpp

+5
Original file line numberDiff line numberDiff line change
@@ -91,6 +91,8 @@ status_t ref_layer_normalization_fwd_t::init(impl::engine_t *engine) {
9191

9292
status_t ref_layer_normalization_fwd_t::execute_forward(
9393
const exec_ctx_t &ctx) const {
94+
if (memory_desc_wrapper(pd()->dst_md()).size() == 0) return status::success;
95+
9496
if (pd()->stats_are_src()) {
9597
return parallel_for(ctx, kernel_, [&](::sycl::handler &cgh) {
9698
layer_normalization_fwd_kernel_vec_t layer_normalization_fwd_kernel(
@@ -163,6 +165,9 @@ status_t ref_layer_normalization_bwd_t::init(impl::engine_t *engine) {
163165

164166
status_t ref_layer_normalization_bwd_t::execute_backward(
165167
const exec_ctx_t &ctx) const {
168+
if (memory_desc_wrapper(pd()->diff_src_md()).size() == 0)
169+
return status::success;
170+
166171
if (pd()->conf_.use_scale || pd()->conf_.use_shift) {
167172
auto status = parallel_for(ctx, kernel_, [&](::sycl::handler &cgh) {
168173
auto nelems_A = memory_desc_wrapper(pd()->src_md(0)).nelems();

‎src/gpu/generic/sycl/ref_matmul.cpp

+2
Original file line numberDiff line numberDiff line change
@@ -142,6 +142,8 @@ status_t ref_matmul_t::init(impl::engine_t *engine) {
142142
}
143143

144144
status_t ref_matmul_t::execute(const exec_ctx_t &ctx) const {
145+
if (memory_desc_wrapper(pd()->dst_md()).size() == 0) return status::success;
146+
145147
sycl_matmul_conf_t conf = pd()->conf_;
146148
if (pd()->any_runtime_params_) {
147149
const auto src_d = ctx.memory_mdw(DNNL_ARG_SRC, pd()->src_md());

‎src/gpu/generic/sycl/ref_pooling.cpp

+9
Original file line numberDiff line numberDiff line change
@@ -74,6 +74,11 @@ status_t ref_pooling_fwd_t::init(impl::engine_t *engine) {
7474
}
7575

7676
status_t ref_pooling_fwd_t::execute_forward(const exec_ctx_t &ctx) const {
77+
// XXX: Add support for 0-dim src
78+
for (auto &md : {pd()->src_md(), pd()->dst_md()}) {
79+
if (memory_desc_wrapper(md).size() == 0) return status::success;
80+
}
81+
7782
return parallel_for(ctx, kernel_, [&](::sycl::handler &cgh) {
7883
auto nelems_A = memory_desc_wrapper(pd()->src_md(0)).nelems();
7984
pooling_fwd_kernel_vec_t pooling_fwd_kernel(pd()->conf_, cgh, ctx);
@@ -135,6 +140,10 @@ status_t ref_pooling_bwd_t::init(impl::engine_t *engine) {
135140
}
136141

137142
status_t ref_pooling_bwd_t::execute_backward(const exec_ctx_t &ctx) const {
143+
for (auto &md : {pd()->diff_src_md(), pd()->diff_dst_md()}) {
144+
if (memory_desc_wrapper(md).size() == 0) return status::success;
145+
}
146+
138147
return parallel_for(ctx, kernel_, [&](::sycl::handler &cgh) {
139148
auto nelems_A = memory_desc_wrapper(pd()->diff_src_md(0)).nelems();
140149
pooling_bwd_kernel_vec_t pooling_bwd_kernel(pd()->conf_, cgh, ctx);

‎src/gpu/generic/sycl/ref_prelu.hpp

+4-2
Original file line numberDiff line numberDiff line change
@@ -56,7 +56,8 @@ struct ref_prelu_fwd_t : public gpu::generic::sycl::primitive_t {
5656
&& (weights_md(0)->format_desc.blocking.inner_nblks == 0)
5757
&& check_data_types(data_d, weights_d, dst_d)
5858
&& md_dims_in_range(src_md())
59-
&& md_dims_in_range(weights_md());
59+
&& md_dims_in_range(weights_md())
60+
&& attr()->has_default_values();
6061

6162
if (!ok) return status::unimplemented;
6263
return init_conf();
@@ -109,7 +110,8 @@ struct ref_prelu_bwd_t : public gpu::generic::sycl::primitive_t {
109110
&& diff_weights_md(0)->data_type == weights_md(0)->data_type
110111
&& check_data_types(data_d, weights_d, diff_dst_d)
111112
&& md_dims_in_range(diff_src_md())
112-
&& md_dims_in_range(weights_md());
113+
&& md_dims_in_range(weights_md())
114+
&& attr()->has_default_values();
113115

114116
if (!ok) return status::unimplemented;
115117

‎src/gpu/generic/sycl/ref_shuffle.hpp

+2
Original file line numberDiff line numberDiff line change
@@ -55,6 +55,8 @@ struct ref_shuffle_t : public gpu::generic::sycl::primitive_t {
5555
&& IMPLICATION(is_fwd(),
5656
src_md(0)->format_desc.blocking.inner_nblks == 0)
5757
&& attr()->has_default_values()
58+
&& memory_desc_wrapper(src_data_md)
59+
== memory_desc_wrapper(dst_data_md)
5860
&& md_dims_in_range(src_md());
5961
if (!ok) return status::unimplemented;
6062
return init_conf();

‎src/gpu/generic/sycl/ref_softmax.cpp

+4
Original file line numberDiff line numberDiff line change
@@ -52,6 +52,8 @@ status_t ref_sycl_softmax_fwd_t::init(impl::engine_t *engine) {
5252
}
5353

5454
status_t ref_sycl_softmax_fwd_t::execute_forward(const exec_ctx_t &ctx) const {
55+
if (pd()->has_zero_dim_memory()) return status::success;
56+
5557
return parallel_for(ctx, kernel_, [&](::sycl::handler &cgh) {
5658
softmax_fwd_kernel_vec_t softmax_fwd_kernel_(pd()->conf_, cgh, ctx);
5759

@@ -82,6 +84,8 @@ status_t ref_sycl_softmax_bwd_t::init(impl::engine_t *engine) {
8284
}
8385

8486
status_t ref_sycl_softmax_bwd_t::execute_backward(const exec_ctx_t &ctx) const {
87+
if (pd()->has_zero_dim_memory()) return status::success;
88+
8589
return parallel_for(ctx, kernel_, [&](::sycl::handler &cgh) {
8690
softmax_bwd_kernel_vec_t softmax_bwd_kernel(pd()->conf_, cgh, ctx);
8791

‎src/gpu/generic/sycl/ref_softmax.hpp

+3-10
Original file line numberDiff line numberDiff line change
@@ -111,22 +111,15 @@ struct ref_sycl_softmax_bwd_t : public gpu::generic::sycl::primitive_t {
111111
&& dst_md()->data_type == diff_dst_md()->data_type
112112
&& attr()->has_default_values()
113113
&& set_default_formats() == status::success
114-
&& check_formats(diff_src_md(), diff_dst_md())
114+
&& memory_desc_wrapper(diff_src_md()).is_plain()
115+
&& memory_desc_wrapper(diff_dst_md()).is_plain()
116+
&& memory_desc_wrapper(dst_md()).is_plain()
115117
&& md_dims_in_range(diff_dst_md());
116118

117119
if (!ok) return status::unimplemented;
118120
return init_conf();
119121
}
120122

121-
static bool check_formats(const memory_desc_wrapper &src,
122-
const memory_desc_wrapper &dst) {
123-
for (const auto &mdw : {src, dst}) {
124-
if (!mdw.is_plain()) return false;
125-
}
126-
127-
return true;
128-
}
129-
130123
sycl_softmax_conf_t conf_;
131124
status_t init_conf();
132125
};

‎tests/gtests/dnnl_test_macros.hpp

+15
Original file line numberDiff line numberDiff line change
@@ -69,6 +69,21 @@
6969
#define SKIP_FOR_LOOP_HIP(cond, message)
7070
#endif
7171

72+
#ifdef DNNL_SYCL_GENERIC
73+
#define SKIP_IF_GENERIC(cond, message) \
74+
do { \
75+
SKIP_IF(get_test_engine_kind() == engine::kind::gpu && (cond), \
76+
(message)); \
77+
} while (0)
78+
79+
#define SKIP_FOR_LOOP_GENERIC(cond, message) \
80+
SKIP_FOR_LOOP( \
81+
get_test_engine_kind() == engine::kind::gpu && (cond), (message));
82+
#else
83+
#define SKIP_IF_GENERIC(cond, message)
84+
#define SKIP_FOR_LOOP_GENERIC(cond, message)
85+
#endif
86+
7287
#define TEST_F_(test_fixture, test_name) TEST_F(test_fixture, test_name)
7388

7489
#define CPU_TEST_F(test_fixture, test_name) \

‎tests/gtests/sycl/api/CMakeLists.txt

+17
Original file line numberDiff line numberDiff line change
@@ -20,3 +20,20 @@ file(GLOB TEST_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/test_*.cpp)
2020
list(APPEND TEST_SOURCES ${MAIN_SRC_GTEST})
2121

2222
register_exe(${TEST_EXE} "${TEST_SOURCES}" "test" "dnnl_gtest")
23+
# InteropReorderAndUserKernel & EltwiseWithUserKernel tests run SYCL kernels
24+
# so they need to be compiled with the correct device triple
25+
if(DNNL_WITH_SYCL)
26+
if(DNNL_SYCL_GENERIC)
27+
CHECK_CXX_COMPILER_FLAG("-fsycl -fsycl-targets=nvptx64-nvidia-cuda" NVIDIA_TARGET_SUPPORTED)
28+
endif()
29+
30+
# Enable linking SYCL kernels.
31+
if(DNNL_SYCL_CUDA OR (DNNL_SYCL_GENERIC AND NVIDIA_TARGET_SUPPORTED))
32+
append(CMAKE_CXX_FLAGS "-fsycl-targets=nvptx64-nvidia-cuda")
33+
append(CMAKE_CXX_FLAGS "-Wno-linker-warnings")
34+
endif()
35+
36+
if(DNNL_AMD_ENABLE_SYCL_KERNELS)
37+
append(CMAKE_CXX_FLAGS "-fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=${DNNL_AMD_SYCL_KERNELS_TARGET_ARCH}")
38+
endif()
39+
endif()

‎tests/gtests/test_batch_normalization.cpp

+10
Original file line numberDiff line numberDiff line change
@@ -59,6 +59,11 @@ bool hip_check_format_tag(tag first_tag, Rest... rest_tags) {
5959
return hip_check_format_tag(rest_tags...);
6060
}
6161

62+
bool generic_check_format_tag(tag atag) {
63+
return impl::utils::one_of(atag, tag::ncw, tag::nchw, tag::ncdhw, tag::nwc,
64+
tag::nhwc, tag::ndhwc, tag::any);
65+
}
66+
6267
class batch_normalization_test_t
6368
: public ::testing::TestWithParam<batch_normalization_test_params_t> {
6469
private:
@@ -80,6 +85,11 @@ class batch_normalization_test_t
8085
SKIP_IF_HIP(!hip_check_format_tag(p.src_tag, p.dst_tag),
8186
"Unsupported format tag");
8287

88+
SKIP_IF_GENERIC(
89+
!generic_check_format_tag(p.src_tag), "Unsupported format tag");
90+
SKIP_IF_GENERIC(
91+
!generic_check_format_tag(p.dst_tag), "Unsupported format tag");
92+
8393
SKIP_IF_CUDA(p.src_dt != p.dst_dt && p.src_dt != dt::undef
8494
&& p.dst_dt != dt::undef,
8595
"Unsupported different data types for source and "

‎tests/gtests/test_binary.cpp

+10
Original file line numberDiff line numberDiff line change
@@ -69,6 +69,8 @@ class binary_test_t : public ::testing::TestWithParam<binary_test_params_t> {
6969
"Unsupported source format tag");
7070
SKIP_IF_HIP(!hip_check_format_tag(tag),
7171
"Unsupported source format tag");
72+
SKIP_IF_GENERIC(!generic_check_format_tag(tag),
73+
"Unsupported source format tag");
7274
}
7375
SKIP_IF_CUDA(!cuda_check_format_tag(p.dst_format),
7476
"Unsupported destination format tag");
@@ -101,6 +103,14 @@ class binary_test_t : public ::testing::TestWithParam<binary_test_params_t> {
101103
return atag == tag::abcd || atag == tag::acdb;
102104
}
103105
bool hip_check_format_tag(tag atag) { return atag == tag::abcd; }
106+
bool generic_check_format_tag(tag atag) {
107+
return impl::utils::one_of(atag, tag::a, tag::ab, tag::abc, tag::abcd,
108+
tag::abcde, tag::abcdef, tag::abdec, tag::acb, tag::acbde,
109+
tag::acbdef, tag::acdb, tag::acdeb, tag::ba, tag::bac,
110+
tag::bacd, tag::bca, tag::bcda, tag::bcdea, tag::cba, tag::cdba,
111+
tag::cdeba, tag::decab, tag::defcab, tag::Ab32a, tag::aBc32b,
112+
tag::any);
113+
}
104114

105115
void Test() {
106116
auto eng = get_test_engine();

‎tests/gtests/test_concat.cpp

+12
Original file line numberDiff line numberDiff line change
@@ -99,6 +99,14 @@ class concat_test_t : public ::testing::TestWithParam<concat_test_params_t> {
9999
dnnl_aBcde4b);
100100
}
101101

102+
bool generic_supported_format_tag(memory::format_tag tag) {
103+
return impl::utils::one_of(tag, dnnl_a, dnnl_ab, dnnl_abc, dnnl_abcd,
104+
dnnl_abcde, dnnl_abcdef, dnnl_abdec, dnnl_acb, dnnl_acbde,
105+
dnnl_acbdef, dnnl_acdb, dnnl_acdeb, dnnl_ba, dnnl_bac,
106+
dnnl_bacd, dnnl_bca, dnnl_bcda, dnnl_bcdea, dnnl_cba, dnnl_cdba,
107+
dnnl_cdeba, dnnl_decab, dnnl_defcab);
108+
}
109+
102110
void SetUp() override {
103111
auto data_type = data_traits<data_t>::data_type;
104112
SKIP_IF_HIP(true, "Concat operator is not supported");
@@ -109,10 +117,14 @@ class concat_test_t : public ::testing::TestWithParam<concat_test_params_t> {
109117
for (size_t i = 0; i < p.srcs_cds.size(); i++) {
110118
SKIP_IF_CUDA(!cuda_supported_format_tag(p.srcs_format[i]),
111119
"Unsupported format tag");
120+
SKIP_IF_GENERIC(!generic_supported_format_tag(p.srcs_format[i]),
121+
"Unsupported format tag");
112122
}
113123

114124
SKIP_IF_CUDA(!cuda_supported_format_tag(p.dst_format),
115125
"Unsupported format tag");
126+
SKIP_IF_GENERIC(!generic_supported_format_tag(p.dst_format),
127+
"Unsupported format tag");
116128
catch_expected_failures(
117129
[&]() { Test(); }, p.expect_to_fail, p.expected_status, false);
118130
}

‎tests/gtests/test_convolution_backward_data_common.hpp

+33
Original file line numberDiff line numberDiff line change
@@ -138,6 +138,23 @@ class convolution_backward_data_test
138138
p.formats.weights_format, p.aalgorithm)),
139139
"Format is not supported.");
140140

141+
SKIP_IF_GENERIC(
142+
!(generic_check_format_tags(p.formats.src_format)
143+
&& generic_check_format_tags(p.formats.dst_format)
144+
&& (generic_check_format_tags(p.formats.weights_format)
145+
|| (impl::utils::one_of(
146+
p.formats.weights_format,
147+
memory::format_tag::goiw,
148+
memory::format_tag::goihw,
149+
memory::format_tag::goidhw,
150+
memory::format_tag::oiw,
151+
memory::format_tag::oihw,
152+
memory::format_tag::oidhw)))
153+
&& check_generic_dt<data_t_diff_src>()
154+
&& check_generic_dt<data_t_diff_dst>()
155+
&& check_generic_dt<data_t_wei>()),
156+
"Format is not supported.");
157+
141158
catch_expected_failures(
142159
[&]() { Test(); }, p.expect_to_fail, p.expected_status);
143160
}
@@ -156,6 +173,14 @@ class convolution_backward_data_test
156173
memory::format_tag::acdeb);
157174
}
158175

176+
bool generic_check_format_tags(memory::format_tag tag) {
177+
return impl::utils::one_of(tag, memory::format_tag::ab,
178+
memory::format_tag::abc, memory::format_tag::abcd,
179+
memory::format_tag::abcde, memory::format_tag::abcdef,
180+
memory::format_tag::acb, memory::format_tag::acdb,
181+
memory::format_tag::acdeb, memory::format_tag::any);
182+
}
183+
159184
bool check_cuda_alg_format(memory::format_tag dst_fmt,
160185
memory::format_tag wei_fmt, algorithm alg) {
161186
bool res = dst_fmt == wei_fmt;
@@ -182,6 +207,14 @@ class convolution_backward_data_test
182207
return res;
183208
}
184209

210+
template <typename dt>
211+
bool check_generic_dt() {
212+
return impl::utils::one_of(data_traits<dt>::data_type,
213+
memory::data_type::f32, memory::data_type::bf16,
214+
memory::data_type::f16, memory::data_type::s32,
215+
memory::data_type::s8, memory::data_type::u8);
216+
}
217+
185218
void Test() {
186219
auto p = ::testing::TestWithParam<
187220
test_convolution_params_t>::GetParam();

0 commit comments

Comments
 (0)