Skip to content

[ESIMD][NFC] - removed unnecessary macro checks #2900

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 1 commit into from
Dec 18, 2020
Merged

[ESIMD][NFC] - removed unnecessary macro checks #2900

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
28 changes: 14 additions & 14 deletions sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -200,12 +200,12 @@ ESIMD_INLINE ESIMD_NODEBUG simd<T, n> block_load(AccessorTy acc,
static_assert(Sz <= 8 * __esimd::OWORD,
"block size must be at most 8 owords");

#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__)
#if defined(__SYCL_DEVICE_ONLY__)
auto surf_ind = AccessorPrivateProxy::getNativeImageObj(acc);
return __esimd_block_read<T, n>(surf_ind, offset);
#else
return __esimd_block_read<T, n>(acc, offset);
#endif // __SYCL_DEVICE_ONLY__ && __SYCL_EXPLICIT_SIMD__
#endif // __SYCL_DEVICE_ONLY__
}

/// Flat-address block-store.
Expand Down Expand Up @@ -240,12 +240,12 @@ ESIMD_INLINE ESIMD_NODEBUG void block_store(AccessorTy acc, uint32_t offset,
static_assert(Sz <= 8 * __esimd::OWORD,
"block size must be at most 8 owords");

#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__)
#if defined(__SYCL_DEVICE_ONLY__)
auto surf_ind = AccessorPrivateProxy::getNativeImageObj(acc);
__esimd_block_write<T, n>(surf_ind, offset >> 4, vals.data());
#else
__esimd_block_write<T, n>(acc, offset >> 4, vals.data());
#endif // __SYCL_DEVICE_ONLY__ && __SYCL_EXPLICIT_SIMD__
#endif // __SYCL_DEVICE_ONLY__
}

/// Accessor-based gather.
Expand Down Expand Up @@ -290,7 +290,7 @@ ESIMD_INLINE ESIMD_NODEBUG
using PromoT =
typename sycl::detail::conditional_t<std::is_signed<T>::value, int32_t,
uint32_t>;
#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__)
#if defined(__SYCL_DEVICE_ONLY__)
const auto surf_ind = AccessorPrivateProxy::getNativeImageObj(acc);
const simd<PromoT, N> promo_vals =
__esimd_surf_read<PromoT, N, decltype(surf_ind), TypeSizeLog2, L1H,
Expand All @@ -302,7 +302,7 @@ ESIMD_INLINE ESIMD_NODEBUG
#endif
return sycl::INTEL::gpu::convert<T>(promo_vals);
} else {
#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__)
#if defined(__SYCL_DEVICE_ONLY__)
const auto surf_ind = AccessorPrivateProxy::getNativeImageObj(acc);
return __esimd_surf_read<T, N, decltype(surf_ind), TypeSizeLog2, L1H, L3H>(
scale, surf_ind, glob_offset, offsets);
Expand Down Expand Up @@ -359,7 +359,7 @@ ESIMD_INLINE ESIMD_NODEBUG
typename sycl::detail::conditional_t<std::is_signed<T>::value, int32_t,
uint32_t>;
const simd<PromoT, N> promo_vals = sycl::INTEL::gpu::convert<PromoT>(vals);
#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__)
#if defined(__SYCL_DEVICE_ONLY__)
const auto surf_ind = AccessorPrivateProxy::getNativeImageObj(acc);
__esimd_surf_write<PromoT, N, decltype(surf_ind), TypeSizeLog2, L1H, L3H>(
pred, scale, surf_ind, glob_offset, offsets, promo_vals);
Expand All @@ -368,7 +368,7 @@ ESIMD_INLINE ESIMD_NODEBUG
pred, scale, acc, glob_offset, offsets, promo_vals);
#endif
} else {
#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__)
#if defined(__SYCL_DEVICE_ONLY__)
const auto surf_ind = AccessorPrivateProxy::getNativeImageObj(acc);
__esimd_surf_write<T, N, decltype(surf_ind), TypeSizeLog2, L1H, L3H>(
pred, scale, surf_ind, glob_offset, offsets, vals);
Expand Down Expand Up @@ -734,7 +734,7 @@ media_block_load(AccessorTy acc, unsigned x, unsigned y) {
static_assert(Width <= 64u, "valid block width is in range [1, 64]");
static_assert(m <= 64u, "valid block height is in range [1, 64]");
static_assert(plane <= 3u, "valid plane index is in range [0, 3]");
#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__)
#if defined(__SYCL_DEVICE_ONLY__)
constexpr unsigned int RoundedWidth =
Width < 4 ? 4 : __esimd::getNextPowerOf2<Width>();

Expand All @@ -751,7 +751,7 @@ media_block_load(AccessorTy acc, unsigned x, unsigned y) {
}
#else
return __esimd_media_block_load<T, m, n>(0, acc, plane, sizeof(T) * n, x, y);
#endif // __SYCL_DEVICE_ONLY__ && __SYCL_EXPLICIT_SIMD__
#endif // __SYCL_DEVICE_ONLY__
}

/// Media block store.
Expand All @@ -776,7 +776,7 @@ media_block_store(AccessorTy acc, unsigned x, unsigned y, simd<T, m * n> vals) {
static_assert(Width <= 64u, "valid block width is in range [1, 64]");
static_assert(m <= 64u, "valid block height is in range [1, 64]");
static_assert(plane <= 3u, "valid plane index is in range [0, 3]");
#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__)
#if defined(__SYCL_DEVICE_ONLY__)
constexpr unsigned int RoundedWidth =
Width < 4 ? 4 : __esimd::getNextPowerOf2<Width>();
constexpr unsigned int n1 = RoundedWidth / sizeof(T);
Expand All @@ -796,7 +796,7 @@ media_block_store(AccessorTy acc, unsigned x, unsigned y, simd<T, m * n> vals) {
}
#else
__esimd_media_block_store<T, m, n>(0, acc, plane, sizeof(T) * n, x, y, vals);
#endif // __SYCL_DEVICE_ONLY__ && __SYCL_EXPLICIT_SIMD__
#endif // __SYCL_DEVICE_ONLY__
}

#ifndef __SYCL_DEVICE_ONLY__
Expand All @@ -813,11 +813,11 @@ SYCL_EXTERNAL void slm_init(uint32_t size) {}
/// \ingroup sycl_esimd
template <typename AccessorTy>
ESIMD_INLINE ESIMD_NODEBUG uint32_t esimd_get_value(AccessorTy acc) {
#if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__)
#if defined(__SYCL_DEVICE_ONLY__)
return __esimd_get_value(AccessorPrivateProxy::getNativeImageObj(acc));
#else
return __esimd_get_value(acc);
#endif // __SYCL_DEVICE_ONLY__ && __SYCL_EXPLICIT_SIMD__
#endif // __SYCL_DEVICE_ONLY__
}

/// \defgroup sycl_esimd_raw_send_api Raw send APIs
Expand Down