@@ -71,6 +71,52 @@ void dpnp_all_c(const void* array1_in, void* result1, const size_t size)
7171 event.wait ();
7272}
7373
74+ template <typename _DataType1, typename _DataType2, typename _ResultType>
75+ class dpnp_allclose_c_kernel ;
76+
77+ template <typename _DataType1, typename _DataType2, typename _ResultType>
78+ void dpnp_allclose_c (const void * array1_in, const void * array2_in, void * result1, const size_t size, double rtol_val, double atol_val)
79+ {
80+ if (!array1_in || !result1)
81+ {
82+ return ;
83+ }
84+
85+ const _DataType1* array1 = reinterpret_cast <const _DataType1*>(array1_in);
86+ const _DataType2* array2 = reinterpret_cast <const _DataType2*>(array2_in);
87+ _ResultType* result = reinterpret_cast <_ResultType*>(result1);
88+
89+ result[0 ] = true ;
90+
91+ if (!size)
92+ {
93+ return ;
94+ }
95+
96+ cl::sycl::event event;
97+
98+ cl::sycl::range<1 > gws (size);
99+ auto kernel_parallel_for_func = [=](cl::sycl::id<1 > global_id) {
100+ size_t i = global_id[0 ];
101+
102+ if (std::abs (array1[i] - array2[i]) > (atol_val + rtol_val * std::abs (array2[i]))){
103+
104+ result[0 ]= false ;
105+
106+ }
107+
108+ };
109+
110+ auto kernel_func = [&](cl::sycl::handler& cgh) {
111+ cgh.parallel_for <class dpnp_allclose_c_kernel <_DataType1, _DataType2, _ResultType>>(gws, kernel_parallel_for_func);
112+ };
113+
114+ event = DPNP_QUEUE.submit (kernel_func);
115+
116+ event.wait ();
117+
118+ }
119+
74120template <typename _DataType, typename _ResultType>
75121class dpnp_any_c_kernel ;
76122
@@ -121,6 +167,23 @@ void func_map_init_logic(func_map_t& fmap)
121167 fmap[DPNPFuncName::DPNP_FN_ALL][eft_FLT][eft_FLT] = {eft_FLT, (void *)dpnp_all_c<float , bool >};
122168 fmap[DPNPFuncName::DPNP_FN_ALL][eft_DBL][eft_DBL] = {eft_DBL, (void *)dpnp_all_c<double , bool >};
123169
170+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_INT][eft_INT] = {eft_BLN, (void *)dpnp_allclose_c<int , int , bool >};
171+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_LNG][eft_INT] = {eft_BLN, (void *)dpnp_allclose_c<long , int , bool >};
172+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_FLT][eft_INT] = {eft_BLN, (void *)dpnp_allclose_c<float , int , bool >};
173+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_DBL][eft_INT] = {eft_BLN, (void *)dpnp_allclose_c<double , int , bool >};
174+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_INT][eft_LNG] = {eft_BLN, (void *)dpnp_allclose_c<int , long , bool >};
175+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_LNG][eft_LNG] = {eft_BLN, (void *)dpnp_allclose_c<long , long , bool >};
176+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_FLT][eft_LNG] = {eft_BLN, (void *)dpnp_allclose_c<float , long , bool >};
177+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_DBL][eft_LNG] = {eft_BLN, (void *)dpnp_allclose_c<double , long , bool >};
178+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_INT][eft_FLT] = {eft_BLN, (void *)dpnp_allclose_c<int , float , bool >};
179+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_LNG][eft_FLT] = {eft_BLN, (void *)dpnp_allclose_c<long , float , bool >};
180+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_FLT][eft_FLT] = {eft_BLN, (void *)dpnp_allclose_c<float , float , bool >};
181+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_DBL][eft_FLT] = {eft_BLN, (void *)dpnp_allclose_c<double , float , bool >};
182+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_INT][eft_DBL] = {eft_BLN, (void *)dpnp_allclose_c<int , double , bool >};
183+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_LNG][eft_DBL] = {eft_BLN, (void *)dpnp_allclose_c<long , double , bool >};
184+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_FLT][eft_DBL] = {eft_BLN, (void *)dpnp_allclose_c<float , double , bool >};
185+ fmap[DPNPFuncName::DPNP_FN_ALLCLOSE][eft_DBL][eft_DBL] = {eft_BLN, (void *)dpnp_allclose_c<double , double , bool >};
186+
124187 fmap[DPNPFuncName::DPNP_FN_ANY][eft_BLN][eft_BLN] = {eft_BLN, (void *)dpnp_any_c<bool , bool >};
125188 fmap[DPNPFuncName::DPNP_FN_ANY][eft_INT][eft_INT] = {eft_INT, (void *)dpnp_any_c<int , bool >};
126189 fmap[DPNPFuncName::DPNP_FN_ANY][eft_LNG][eft_LNG] = {eft_LNG, (void *)dpnp_any_c<long , bool >};
0 commit comments