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

Conversion from raw to multi_ptr should be done with address_space_cast #1538

Merged
merged 1 commit into from
Aug 27, 2023
Merged

Conversion from raw to multi_ptr should be done with address_space_cast #1538

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
14 changes: 8 additions & 6 deletions dpnp/backend/kernels/dpnp_krnl_bitwise.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,12 +68,14 @@ DPCTLSyclEventRef dpnp_invert_c(DPCTLSyclQueueRef q_ref,
sg.get_group_id()[0] * max_sg_size);

if (start + static_cast<size_t>(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<vec_sz>(multi_ptrT(&input_data[start]));
sycl::vec<_DataType, vec_sz> x = sg.load<vec_sz>(input_multi_ptr);
sycl::vec<_DataType, vec_sz> res_vec;

if constexpr (std::is_same_v<_DataType, bool>) {
Expand All @@ -86,7 +88,7 @@ DPCTLSyclEventRef dpnp_invert_c(DPCTLSyclQueueRef q_ref,
res_vec = ~x;
}

sg.store<vec_sz>(multi_ptrT(&result[start]), res_vec);
sg.store<vec_sz>(result_multi_ptr, res_vec);
}
else {
for (size_t k = start + sg.get_local_id()[0]; k < size;
Expand Down
38 changes: 18 additions & 20 deletions dpnp/backend/kernels/dpnp_krnl_elemwise.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1344,12 +1344,17 @@ static void func_map_init_elemwise_1arg_1type(func_map_t &fmap)
\
if (start + static_cast<size_t>(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; \
\
Expand All @@ -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<vec_sz>( \
input1_ptrT(&input1_data[start])); \
sg.load<vec_sz>(input1_multi_ptr); \
sycl::vec<_DataType_input2, vec_sz> x2 = \
sg.load<vec_sz>( \
input2_ptrT(&input2_data[start])); \
sg.load<vec_sz>(input2_multi_ptr); \
\
res_vec = __vec_operation__; \
} \
Expand All @@ -1377,33 +1380,28 @@ 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<vec_sz>(input1_ptrT( \
&input1_data[start]))); \
sg.load<vec_sz>(input1_multi_ptr)); \
sycl::vec<_DataType_output, vec_sz> x2 = \
dpnp_vec_cast<_DataType_output, \
_DataType_input2, vec_sz>( \
sg.load<vec_sz>(input2_ptrT( \
&input2_data[start]))); \
sg.load<vec_sz>(input2_multi_ptr)); \
\
res_vec = __vec_operation__; \
} \
} \
else { \
sycl::vec<_DataType_input1, vec_sz> x1 = \
sg.load<vec_sz>( \
input1_ptrT(&input1_data[start])); \
sg.load<vec_sz>(input1_multi_ptr); \
sycl::vec<_DataType_input2, vec_sz> x2 = \
sg.load<vec_sz>( \
input2_ptrT(&input2_data[start])); \
sg.load<vec_sz>(input2_multi_ptr); \
\
for (size_t k = 0; k < vec_sz; ++k) { \
const _DataType_output input1_elem = x1[k]; \
const _DataType_output input2_elem = x2[k]; \
res_vec[k] = __operation__; \
} \
} \
sg.store<vec_sz>(result_ptrT(&result[start]), \
res_vec); \
sg.store<vec_sz>(result_multi_ptr, res_vec); \
} \
else { \
for (size_t k = start + sg.get_local_id()[0]; \
Expand Down
24 changes: 15 additions & 9 deletions dpnp/backend/kernels/dpnp_krnl_logic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -537,22 +537,28 @@ DPCTLSyclEventRef (*dpnp_any_ext_c)(DPCTLSyclQueueRef,
\
if (start + static_cast<size_t>(vec_sz) * max_sg_size < \
result_size) { \
sycl::vec<_DataType_input1, vec_sz> x1 = sg.load<vec_sz>( \
sycl::multi_ptr<_DataType_input1, global_space>( \
&input1_data[start])); \
sycl::vec<_DataType_input2, vec_sz> x2 = sg.load<vec_sz>( \
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<vec_sz>(input1_multi_ptr); \
sycl::vec<_DataType_input2, vec_sz> x2 = \
sg.load<vec_sz>(input2_multi_ptr); \
sycl::vec<bool, vec_sz> res_vec; \
\
for (size_t k = 0; k < vec_sz; ++k) { \
const _DataType_input1 input1_elem = x1[k]; \
const _DataType_input2 input2_elem = x2[k]; \
res_vec[k] = __operation__; \
} \
sg.store<vec_sz>( \
sycl::multi_ptr<bool, global_space>(&result[start]), \
res_vec); \
sg.store<vec_sz>(result_multi_ptr, res_vec); \
} \
else { \
for (size_t k = start; k < result_size; ++k) { \
Expand Down
16 changes: 8 additions & 8 deletions dpnp/backend/kernels/dpnp_krnl_mathematical.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -166,18 +164,20 @@ DPCTLSyclEventRef
sg.get_group_id()[0] * max_sg_size);

if (start + static_cast<size_t>(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<vec_sz>(input_ptrT(&array1[start]));
sg.load<vec_sz>(array_multi_ptr);

sycl::vec<_DataType_output, vec_sz> res_vec =
sycl::abs(data_vec);

sg.store<vec_sz>(result_ptrT(&result[start]), res_vec);
sg.store<vec_sz>(result_multi_ptr, res_vec);
}
else {
for (size_t k = start + sg.get_local_id()[0]; k < size;
Expand Down