Skip to content

[SYCL] removing std::memcpy from sycl headers #11436

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

Merged
merged 3 commits into from
Oct 16, 2023
Merged

[SYCL] removing std::memcpy from sycl headers #11436

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
951997d
removing std::memcpy from sycl headers, continuation of work original…
cperkinsintel Oct 4, 2023
e328db3
reviewer feedback
cperkinsintel Oct 13, 2023
8dc54a5
reviewer feedback
cperkinsintel Oct 16, 2023
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
26 changes: 13 additions & 13 deletions sycl/include/sycl/detail/spirv.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@

#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp> // for IdToMaskPosition

#include <cstring> // for std::memcpy
#include <sycl/detail/memcpy.hpp> // sycl::detail::memcpy

namespace sycl {
inline namespace _V1 {
Expand Down Expand Up @@ -370,9 +370,9 @@ EnableIfGenericBroadcast<T, IdT> GroupBroadcast(Group g, T x, IdT local_id) {
char *ResultBytes = reinterpret_cast<char *>(&Result);
auto BroadcastBytes = [=](size_t Offset, size_t Size) {
uint64_t BroadcastX, BroadcastResult;
std::memcpy(&BroadcastX, XBytes + Offset, Size);
detail::memcpy(&BroadcastX, XBytes + Offset, Size);
BroadcastResult = GroupBroadcast(g, BroadcastX, local_id);
std::memcpy(ResultBytes + Offset, &BroadcastResult, Size);
detail::memcpy(ResultBytes + Offset, &BroadcastResult, Size);
};
GenericCall<T>(BroadcastBytes);
return Result;
Expand Down Expand Up @@ -424,9 +424,9 @@ EnableIfGenericBroadcast<T> GroupBroadcast(Group g, T x,
char *ResultBytes = reinterpret_cast<char *>(&Result);
auto BroadcastBytes = [=](size_t Offset, size_t Size) {
uint64_t BroadcastX, BroadcastResult;
std::memcpy(&BroadcastX, XBytes + Offset, Size);
detail::memcpy(&BroadcastX, XBytes + Offset, Size);
BroadcastResult = GroupBroadcast(g, BroadcastX, local_id);
std::memcpy(ResultBytes + Offset, &BroadcastResult, Size);
detail::memcpy(ResultBytes + Offset, &BroadcastResult, Size);
};
GenericCall<T>(BroadcastBytes);
return Result;
Expand Down Expand Up @@ -957,9 +957,9 @@ EnableIfGenericShuffle<T> SubgroupShuffle(T x, id<1> local_id) {
char *ResultBytes = reinterpret_cast<char *>(&Result);
auto ShuffleBytes = [=](size_t Offset, size_t Size) {
ShuffleChunkT ShuffleX, ShuffleResult;
std::memcpy(&ShuffleX, XBytes + Offset, Size);
detail::memcpy(&ShuffleX, XBytes + Offset, Size);
ShuffleResult = SubgroupShuffle(ShuffleX, local_id);
std::memcpy(ResultBytes + Offset, &ShuffleResult, Size);
detail::memcpy(ResultBytes + Offset, &ShuffleResult, Size);
};
GenericCall<T>(ShuffleBytes);
return Result;
Expand All @@ -972,9 +972,9 @@ EnableIfGenericShuffle<T> SubgroupShuffleXor(T x, id<1> local_id) {
char *ResultBytes = reinterpret_cast<char *>(&Result);
auto ShuffleBytes = [=](size_t Offset, size_t Size) {
ShuffleChunkT ShuffleX, ShuffleResult;
std::memcpy(&ShuffleX, XBytes + Offset, Size);
detail::memcpy(&ShuffleX, XBytes + Offset, Size);
ShuffleResult = SubgroupShuffleXor(ShuffleX, local_id);
std::memcpy(ResultBytes + Offset, &ShuffleResult, Size);
detail::memcpy(ResultBytes + Offset, &ShuffleResult, Size);
};
GenericCall<T>(ShuffleBytes);
return Result;
Expand All @@ -987,9 +987,9 @@ EnableIfGenericShuffle<T> SubgroupShuffleDown(T x, uint32_t delta) {
char *ResultBytes = reinterpret_cast<char *>(&Result);
auto ShuffleBytes = [=](size_t Offset, size_t Size) {
ShuffleChunkT ShuffleX, ShuffleResult;
std::memcpy(&ShuffleX, XBytes + Offset, Size);
detail::memcpy(&ShuffleX, XBytes + Offset, Size);
ShuffleResult = SubgroupShuffleDown(ShuffleX, delta);
std::memcpy(ResultBytes + Offset, &ShuffleResult, Size);
detail::memcpy(ResultBytes + Offset, &ShuffleResult, Size);
};
GenericCall<T>(ShuffleBytes);
return Result;
Expand All @@ -1002,9 +1002,9 @@ EnableIfGenericShuffle<T> SubgroupShuffleUp(T x, uint32_t delta) {
char *ResultBytes = reinterpret_cast<char *>(&Result);
auto ShuffleBytes = [=](size_t Offset, size_t Size) {
ShuffleChunkT ShuffleX, ShuffleResult;
std::memcpy(&ShuffleX, XBytes + Offset, Size);
detail::memcpy(&ShuffleX, XBytes + Offset, Size);
ShuffleResult = SubgroupShuffleUp(ShuffleX, delta);
std::memcpy(ResultBytes + Offset, &ShuffleResult, Size);
detail::memcpy(ResultBytes + Offset, &ShuffleResult, Size);
};
GenericCall<T>(ShuffleBytes);
return Result;
Expand Down
13 changes: 7 additions & 6 deletions sycl/include/sycl/ext/oneapi/experimental/bfloat16_math.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,10 +9,11 @@
#pragma once

#include <sycl/builtins.hpp> // for ceil, cos, exp, exp10, exp2
#include <sycl/detail/memcpy.hpp> // sycl::detail::memcpy
#include <sycl/ext/oneapi/bfloat16.hpp> // for bfloat16, bfloat16ToBits
#include <sycl/marray.hpp> // for marray

#include <cstring> // for size_t, memcpy
#include <cstring> // for size_t
#include <stdint.h> // for uint32_t
#include <type_traits> // for enable_if_t, is_same

Expand All @@ -24,7 +25,7 @@ namespace detail {
template <size_t N>
uint32_t to_uint32_t(sycl::marray<bfloat16, N> x, size_t start) {
uint32_t res;
std::memcpy(&res, &x[start], sizeof(uint32_t));
sycl::detail::memcpy(&res, &x[start], sizeof(uint32_t));
return res;
}
} // namespace detail
Expand Down Expand Up @@ -71,7 +72,7 @@ sycl::marray<bfloat16, N> fabs(sycl::marray<bfloat16, N> x) {
(__SYCL_CUDA_ARCH__ >= 800)
for (size_t i = 0; i < N / 2; i++) {
auto partial_res = __clc_fabs(detail::to_uint32_t(x, i * 2));
std::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
sycl::detail::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
}

if (N % 2) {
Expand Down Expand Up @@ -126,7 +127,7 @@ sycl::marray<bfloat16, N> fmin(sycl::marray<bfloat16, N> x,
for (size_t i = 0; i < N / 2; i++) {
auto partial_res = __clc_fmin(detail::to_uint32_t(x, i * 2),
detail::to_uint32_t(y, i * 2));
std::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
sycl::detail::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
}

if (N % 2) {
Expand Down Expand Up @@ -182,7 +183,7 @@ sycl::marray<bfloat16, N> fmax(sycl::marray<bfloat16, N> x,
for (size_t i = 0; i < N / 2; i++) {
auto partial_res = __clc_fmax(detail::to_uint32_t(x, i * 2),
detail::to_uint32_t(y, i * 2));
std::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
sycl::detail::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
}

if (N % 2) {
Expand Down Expand Up @@ -226,7 +227,7 @@ sycl::marray<bfloat16, N> fma(sycl::marray<bfloat16, N> x,
auto partial_res =
__clc_fma(detail::to_uint32_t(x, i * 2), detail::to_uint32_t(y, i * 2),
detail::to_uint32_t(z, i * 2));
std::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
sycl::detail::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
}

if (N % 2) {
Expand Down
7 changes: 4 additions & 3 deletions sycl/include/sycl/ext/oneapi/experimental/builtins.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,10 +13,11 @@
#include <sycl/detail/builtins.hpp> // for __invoke_exp2, __invo...
#include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE
#include <sycl/detail/generic_type_traits.hpp> // for is_svgenfloath, is_sv...
#include <sycl/detail/memcpy.hpp> // detail::memcpy
#include <sycl/marray.hpp> // for marray
#include <sycl/types.hpp> // for vec

#include <cstring> // for memcpy, size_t
#include <cstring> // for size_t
#include <stdio.h> // for printf
#include <type_traits> // for enable_if_t

Expand Down Expand Up @@ -124,7 +125,7 @@ inline __SYCL_ALWAYS_INLINE
auto partial_res = __sycl_std::__invoke_tanh<sycl::vec<T, 2>>(
sycl::detail::to_vec2(x, i * 2));
#endif
std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>));
sycl::detail::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>));
}
if (N % 2) {
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
Expand Down Expand Up @@ -167,7 +168,7 @@ exp2(sycl::marray<half, N> x) __NOEXC {
auto partial_res = __sycl_std::__invoke_exp2<sycl::vec<half, 2>>(
sycl::detail::to_vec2(x, i * 2));
#endif
std::memcpy(&res[i * 2], &partial_res, sizeof(vec<half, 2>));
sycl::detail::memcpy(&res[i * 2], &partial_res, sizeof(vec<half, 2>));
}
if (N % 2) {
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
Expand Down
7 changes: 3 additions & 4 deletions sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@
#pragma once

#include <sycl/detail/helpers.hpp> // for Builder
#include <sycl/detail/memcpy.hpp> // detail::memcpy
#include <sycl/exception.hpp> // for errc, exception
#include <sycl/feature_test.hpp> // for SYCL_EXT_ONEAPI_SUB_GROUP_MASK
#include <sycl/id.hpp> // for id
Expand All @@ -16,7 +17,6 @@

#include <assert.h> // for assert
#include <climits> // for CHAR_BIT
#include <cstring> // for memcpy
#include <stddef.h> // for size_t
#include <stdint.h> // for uint32_t
#include <system_error> // for error_code
Expand Down Expand Up @@ -110,9 +110,8 @@ struct sub_group_mask {
size_t RemainingBytes = sizeof(Bits) - BytesCopied;
size_t BytesToCopy =
RemainingBytes < sizeof(T) ? RemainingBytes : sizeof(T);
// TODO: memcpy is not guaranteed to work in kernels. Find alternative.
std::memcpy(reinterpret_cast<char *>(&Bits) + BytesCopied, &val[I],
BytesToCopy);
sycl::detail::memcpy(reinterpret_cast<char *>(&Bits) + BytesCopied,
&val[I], BytesToCopy);
BytesCopied += BytesToCopy;
}
}
Expand Down