diff --git a/dpnp/backend/kernels/dpnp_krnl_bitwise.cpp b/dpnp/backend/kernels/dpnp_krnl_bitwise.cpp index d1e3d41235a1..0134e94ef39d 100644 --- a/dpnp/backend/kernels/dpnp_krnl_bitwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_bitwise.cpp @@ -68,12 +68,14 @@ DPCTLSyclEventRef dpnp_invert_c(DPCTLSyclQueueRef q_ref, sg.get_group_id()[0] * max_sg_size); if (start + static_cast(vec_sz) * max_sg_size < size) { - using multi_ptrT = - sycl::multi_ptr<_DataType, - sycl::access::address_space::global_space>; + auto input_multi_ptr = sycl::address_space_cast< + sycl::access::address_space::global_space, + sycl::access::decorated::yes>(&input_data[start]); + auto result_multi_ptr = sycl::address_space_cast< + sycl::access::address_space::global_space, + sycl::access::decorated::yes>(&result[start]); - sycl::vec<_DataType, vec_sz> x = - sg.load(multi_ptrT(&input_data[start])); + sycl::vec<_DataType, vec_sz> x = sg.load(input_multi_ptr); sycl::vec<_DataType, vec_sz> res_vec; if constexpr (std::is_same_v<_DataType, bool>) { @@ -86,7 +88,7 @@ DPCTLSyclEventRef dpnp_invert_c(DPCTLSyclQueueRef q_ref, res_vec = ~x; } - sg.store(multi_ptrT(&result[start]), res_vec); + sg.store(result_multi_ptr, res_vec); } else { for (size_t k = start + sg.get_local_id()[0]; k < size; diff --git a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp index 8fb510e9fab1..6fc494b43849 100644 --- a/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_elemwise.cpp @@ -1344,12 +1344,17 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) \ if (start + static_cast(vec_sz) * max_sg_size < \ result_size) { \ - using input1_ptrT = \ - sycl::multi_ptr<_DataType_input1, global_space>; \ - using input2_ptrT = \ - sycl::multi_ptr<_DataType_input2, global_space>; \ - using result_ptrT = \ - sycl::multi_ptr<_DataType_output, global_space>; \ + auto input1_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>( \ + &input1_data[start]); \ + auto input2_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>( \ + &input2_data[start]); \ + auto result_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>(&result[start]); \ \ sycl::vec<_DataType_output, vec_sz> res_vec; \ \ @@ -1363,11 +1368,9 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) _DataType_output>) \ { \ sycl::vec<_DataType_input1, vec_sz> x1 = \ - sg.load( \ - input1_ptrT(&input1_data[start])); \ + sg.load(input1_multi_ptr); \ sycl::vec<_DataType_input2, vec_sz> x2 = \ - sg.load( \ - input2_ptrT(&input2_data[start])); \ + sg.load(input2_multi_ptr); \ \ res_vec = __vec_operation__; \ } \ @@ -1377,24 +1380,20 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) sycl::vec<_DataType_output, vec_sz> x1 = \ dpnp_vec_cast<_DataType_output, \ _DataType_input1, vec_sz>( \ - sg.load(input1_ptrT( \ - &input1_data[start]))); \ + sg.load(input1_multi_ptr)); \ sycl::vec<_DataType_output, vec_sz> x2 = \ dpnp_vec_cast<_DataType_output, \ _DataType_input2, vec_sz>( \ - sg.load(input2_ptrT( \ - &input2_data[start]))); \ + sg.load(input2_multi_ptr)); \ \ res_vec = __vec_operation__; \ } \ } \ else { \ sycl::vec<_DataType_input1, vec_sz> x1 = \ - sg.load( \ - input1_ptrT(&input1_data[start])); \ + sg.load(input1_multi_ptr); \ sycl::vec<_DataType_input2, vec_sz> x2 = \ - sg.load( \ - input2_ptrT(&input2_data[start])); \ + sg.load(input2_multi_ptr); \ \ for (size_t k = 0; k < vec_sz; ++k) { \ const _DataType_output input1_elem = x1[k]; \ @@ -1402,8 +1401,7 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap) res_vec[k] = __operation__; \ } \ } \ - sg.store(result_ptrT(&result[start]), \ - res_vec); \ + sg.store(result_multi_ptr, res_vec); \ } \ else { \ for (size_t k = start + sg.get_local_id()[0]; \ diff --git a/dpnp/backend/kernels/dpnp_krnl_logic.cpp b/dpnp/backend/kernels/dpnp_krnl_logic.cpp index 1757d053416a..ac8f7ca4560b 100644 --- a/dpnp/backend/kernels/dpnp_krnl_logic.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_logic.cpp @@ -537,12 +537,20 @@ DPCTLSyclEventRef (*dpnp_any_ext_c)(DPCTLSyclQueueRef, \ if (start + static_cast(vec_sz) * max_sg_size < \ result_size) { \ - sycl::vec<_DataType_input1, vec_sz> x1 = sg.load( \ - sycl::multi_ptr<_DataType_input1, global_space>( \ - &input1_data[start])); \ - sycl::vec<_DataType_input2, vec_sz> x2 = sg.load( \ - sycl::multi_ptr<_DataType_input2, global_space>( \ - &input2_data[start])); \ + auto input1_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>(&input1_data[start]); \ + auto input2_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>(&input2_data[start]); \ + auto result_multi_ptr = sycl::address_space_cast< \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes>(&result[start]); \ + \ + sycl::vec<_DataType_input1, vec_sz> x1 = \ + sg.load(input1_multi_ptr); \ + sycl::vec<_DataType_input2, vec_sz> x2 = \ + sg.load(input2_multi_ptr); \ sycl::vec res_vec; \ \ for (size_t k = 0; k < vec_sz; ++k) { \ @@ -550,9 +558,7 @@ DPCTLSyclEventRef (*dpnp_any_ext_c)(DPCTLSyclQueueRef, const _DataType_input2 input2_elem = x2[k]; \ res_vec[k] = __operation__; \ } \ - sg.store( \ - sycl::multi_ptr(&result[start]), \ - res_vec); \ + sg.store(result_multi_ptr, res_vec); \ } \ else { \ for (size_t k = start; k < result_size; ++k) { \ diff --git a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp index 339924240a32..b4f5cd96b4d0 100644 --- a/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_mathematical.cpp @@ -151,8 +151,6 @@ DPCTLSyclEventRef constexpr size_t lws = 64; constexpr unsigned int vec_sz = 8; - constexpr sycl::access::address_space global_space = - sycl::access::address_space::global_space; auto gws_range = sycl::range<1>(((size + lws * vec_sz - 1) / (lws * vec_sz)) * lws); @@ -166,18 +164,20 @@ DPCTLSyclEventRef sg.get_group_id()[0] * max_sg_size); if (start + static_cast(vec_sz) * max_sg_size < size) { - using input_ptrT = - sycl::multi_ptr<_DataType_input, global_space>; - using result_ptrT = - sycl::multi_ptr<_DataType_output, global_space>; + auto array_multi_ptr = sycl::address_space_cast< + sycl::access::address_space::global_space, + sycl::access::decorated::yes>(&array1[start]); + auto result_multi_ptr = sycl::address_space_cast< + sycl::access::address_space::global_space, + sycl::access::decorated::yes>(&result[start]); sycl::vec<_DataType_input, vec_sz> data_vec = - sg.load(input_ptrT(&array1[start])); + sg.load(array_multi_ptr); sycl::vec<_DataType_output, vec_sz> res_vec = sycl::abs(data_vec); - sg.store(result_ptrT(&result[start]), res_vec); + sg.store(result_multi_ptr, res_vec); } else { for (size_t k = start + sg.get_local_id()[0]; k < size;