[SYCL] Implement SYCL part of sycl_ext_oneapi_prefetch #11458
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
KornevNikita marked this conversation as resolved.
Show resolved
Hide resolved
|
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,301 @@ | ||
| //==--------------- prefetch.hpp --- SYCL prefetch extension ---------------==// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #pragma once | ||
|
|
||
| #include <CL/__spirv/spirv_ops.hpp> | ||
| #include <sycl/ext/oneapi/properties/properties.hpp> | ||
|
|
||
| namespace sycl { | ||
| inline namespace _V1 { | ||
| namespace ext::oneapi::experimental { | ||
|
|
||
| enum class cache_level { L1 = 0, L2 = 1, L3 = 2, L4 = 3 }; | ||
|
|
||
| struct nontemporal; | ||
|
|
||
| struct prefetch_hint_key { | ||
| template <cache_level Level, typename Hint> | ||
| using value_t = | ||
| property_value<prefetch_hint_key, | ||
| std::integral_constant<cache_level, Level>, Hint>; | ||
| }; | ||
|
|
||
| template <cache_level Level, typename Hint> | ||
| inline constexpr prefetch_hint_key::value_t<Level, Hint> prefetch_hint; | ||
|
|
||
| inline constexpr prefetch_hint_key::value_t<cache_level::L1, void> | ||
| prefetch_hint_L1; | ||
| inline constexpr prefetch_hint_key::value_t<cache_level::L2, void> | ||
| prefetch_hint_L2; | ||
| inline constexpr prefetch_hint_key::value_t<cache_level::L3, void> | ||
| prefetch_hint_L3; | ||
| inline constexpr prefetch_hint_key::value_t<cache_level::L4, void> | ||
| prefetch_hint_L4; | ||
|
|
||
| inline constexpr prefetch_hint_key::value_t<cache_level::L1, nontemporal> | ||
| prefetch_hint_L1_nt; | ||
| inline constexpr prefetch_hint_key::value_t<cache_level::L2, nontemporal> | ||
| prefetch_hint_L2_nt; | ||
| inline constexpr prefetch_hint_key::value_t<cache_level::L3, nontemporal> | ||
| prefetch_hint_L3_nt; | ||
| inline constexpr prefetch_hint_key::value_t<cache_level::L4, nontemporal> | ||
| prefetch_hint_L4_nt; | ||
|
|
||
| namespace detail { | ||
| template <> struct IsCompileTimeProperty<prefetch_hint_key> : std::true_type {}; | ||
|
|
||
| template <cache_level Level, typename Hint> | ||
| struct PropertyMetaInfo<prefetch_hint_key::value_t<Level, Hint>> { | ||
| static constexpr const char *name = std::is_same_v<Hint, nontemporal> | ||
| ? "sycl-prefetch-hint-nt" | ||
| : "sycl-prefetch-hint"; | ||
| static constexpr int value = static_cast<int>(Level); | ||
KornevNikita marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| }; | ||
|
|
||
| template <typename Properties> | ||
| void prefetch_impl(void *ptr, size_t bytes, Properties properties) { | ||
| #ifdef __SYCL_DEVICE_ONLY__ | ||
KornevNikita marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| auto *ptrGlobalAS = __SYCL_GenericCastToPtrExplicit_ToGlobal<char>(ptr); | ||
| __attribute__((opencl_global)) char *ptrAnnotated = nullptr; | ||
| if constexpr (!properties.template has_property<prefetch_hint_key>()) { | ||
| ptrAnnotated = __builtin_intel_sycl_ptr_annotation( | ||
| ptrGlobalAS, "sycl-prefetch-hint", static_cast<int>(cache_level::L1)); | ||
| } else { | ||
| auto prop = properties.template get_property<prefetch_hint_key>(); | ||
| ptrAnnotated = __builtin_intel_sycl_ptr_annotation( | ||
| ptrGlobalAS, PropertyMetaInfo<decltype(prop)>::name, | ||
| PropertyMetaInfo<decltype(prop)>::value); | ||
| } | ||
| __spirv_ocl_prefetch(ptrAnnotated, bytes); | ||
| #endif | ||
| std::ignore = ptr; | ||
| std::ignore = bytes; | ||
| std::ignore = properties; | ||
| } | ||
| } // namespace detail | ||
|
|
||
| template <typename Properties = empty_properties_t> | ||
| void prefetch(void *ptr, Properties properties = {}) { | ||
| detail::prefetch_impl(ptr, 1, properties); | ||
| } | ||
|
|
||
| template <typename Properties = empty_properties_t> | ||
| void prefetch(void *ptr, size_t bytes, Properties properties = {}) { | ||
| detail::prefetch_impl(ptr, bytes, properties); | ||
| } | ||
|
|
||
| template <typename T, typename Properties = empty_properties_t> | ||
| void prefetch(T *ptr, Properties properties = {}) { | ||
| prefetch((void *)ptr, sizeof(T), properties); | ||
| } | ||
|
|
||
| template <typename T, typename Properties = empty_properties_t> | ||
| void prefetch(T *ptr, size_t count, Properties properties = {}) { | ||
| prefetch((void *)ptr, count * sizeof(T), properties); | ||
| } | ||
|
|
||
| // Only available if AddressSpace == global_space || AddressSpace == | ||
| // generic_space | ||
KornevNikita marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| template <access::address_space AddressSpace, access::decorated IsDecorated, | ||
| typename Properties = empty_properties_t> | ||
| std::enable_if_t<AddressSpace == access::address_space::global_space || | ||
| AddressSpace == access::address_space::generic_space, | ||
| void> | ||
KornevNikita marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, | ||
| Properties properties = {}) { | ||
| prefetch(ptr.get(), properties); | ||
| } | ||
|
|
||
| // Only available if AddressSpace == global_space || AddressSpace == | ||
| // generic_space | ||
| template <access::address_space AddressSpace, access::decorated IsDecorated, | ||
| typename Properties = empty_properties_t> | ||
| std::enable_if_t<AddressSpace == access::address_space::global_space || | ||
| AddressSpace == access::address_space::generic_space, | ||
| void> | ||
| prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes, | ||
| Properties properties = {}) { | ||
| prefetch(ptr.get(), bytes, properties); | ||
| } | ||
|
|
||
| // Only available if AddressSpace == global_space || AddressSpace == | ||
| // generic_space | ||
| template <typename T, access::address_space AddressSpace, | ||
| access::decorated IsDecorated, | ||
| typename Properties = empty_properties_t> | ||
| std::enable_if_t<AddressSpace == access::address_space::global_space || | ||
| AddressSpace == access::address_space::generic_space, | ||
| void> | ||
| prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, | ||
| Properties properties = {}) { | ||
| prefetch(ptr.get(), properties); | ||
| } | ||
|
|
||
| // Only available if AddressSpace == global_space || AddressSpace == | ||
| // generic_space | ||
| template <typename T, access::address_space AddressSpace, | ||
| access::decorated IsDecorated, | ||
| typename Properties = empty_properties_t> | ||
| std::enable_if_t<AddressSpace == access::address_space::global_space || | ||
| AddressSpace == access::address_space::generic_space, | ||
| void> | ||
| prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count, | ||
| Properties properties = {}) { | ||
| prefetch(ptr.get(), count, properties); | ||
| } | ||
|
|
||
| // Only available if Dimensions > 0 && (AccessMode == read || AccessMode == | ||
| // read_write) | ||
| template <typename DataT, int Dimensions, access_mode AccessMode, | ||
| access::placeholder IsPlaceholder, | ||
| typename Properties = empty_properties_t> | ||
| std::enable_if_t<(Dimensions > 0) && (AccessMode == access_mode::read || | ||
| AccessMode == access_mode::read_write), | ||
| void> | ||
| prefetch( | ||
| accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc, | ||
| id<Dimensions> offset, Properties properties = {}) { | ||
| prefetch((void *)&acc[offset], sizeof(DataT), properties); | ||
KornevNikita marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| } | ||
|
|
||
| // Only available if Dimensions > 0 && (AccessMode == read || AccessMode == | ||
| // read_write) | ||
| template <typename DataT, int Dimensions, access_mode AccessMode, | ||
| access::placeholder IsPlaceholder, | ||
| typename Properties = empty_properties_t> | ||
| std::enable_if_t<(Dimensions > 0) && (AccessMode == access_mode::read || | ||
| AccessMode == access_mode::read_write), | ||
| void> | ||
| prefetch( | ||
| accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc, | ||
| size_t offset, size_t count, Properties properties = {}) { | ||
| prefetch((void *)&acc[offset], count * sizeof(DataT), properties); | ||
|
There was a problem hiding this comment. Do we know for sure that There was a problem hiding this comment. Good point. @Pennycook I'm not sure how it's intended to work in case of N-dim offset. Should we call the __spirv_ocl_prefetch spirv instruction several times for different memory segments in such case or there should be some constraint? There was a problem hiding this comment. We struggled a bit with multi-dimensional prefetches (see the issues). Where we landed (for now) is that the block being prefetched is assumed to be contiguous, and only the offset itself is multi-dimensional. It's effectively a shorthand to avoid computing the linear offset from the start of the buffer. Note that the specification says for the multi-dimensional cases:
If somebody requests a multi-dimensional prefetch later, we can describe it with a |
||
| } | ||
|
|
||
| template <typename Group, typename Properties = empty_properties_t> | ||
| typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void> | ||
| joint_prefetch(Group g, void *ptr, Properties properties = {}) { | ||
| detail::prefetch_impl(ptr, 1, properties); | ||
| } | ||
|
|
||
| template <typename Group, typename Properties = empty_properties_t> | ||
| typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void> | ||
| joint_prefetch(Group g, void *ptr, size_t bytes, Properties properties = {}) { | ||
| detail::prefetch_impl(ptr, bytes, properties); | ||
| } | ||
|
|
||
| template <typename Group, typename T, typename Properties = empty_properties_t> | ||
| typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void> | ||
| joint_prefetch(Group g, T *ptr, Properties properties = {}) { | ||
| joint_prefetch((void *)ptr, sizeof(T), properties); | ||
| } | ||
|
|
||
| template <typename Group, typename T, typename Properties = empty_properties_t> | ||
| typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void> | ||
| joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) { | ||
| joint_prefetch((void *)ptr, count * sizeof(T), properties); | ||
| } | ||
|
|
||
| // Only available if AddressSpace == global_space || AddressSpace == | ||
| // generic_space | ||
| template <typename Group, access::address_space AddressSpace, | ||
| access::decorated IsDecorated, | ||
| typename Properties = ext::oneapi::experimental::empty_properties_t> | ||
| typename std::enable_if_t< | ||
| sycl::is_group_v<std::decay_t<Group>> && | ||
| (AddressSpace == access::address_space::global_space || | ||
| AddressSpace == access::address_space::generic_space), | ||
| void> | ||
| joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr, | ||
| Properties properties = {}) { | ||
| joint_prefetch(g, ptr.get(), properties); | ||
| } | ||
|
|
||
| // Only available if AddressSpace == global_space || AddressSpace == | ||
| // generic_space | ||
| template <typename Group, access::address_space AddressSpace, | ||
| access::decorated IsDecorated, | ||
| typename Properties = ext::oneapi::experimental::empty_properties_t> | ||
| typename std::enable_if_t< | ||
| sycl::is_group_v<std::decay_t<Group>> && | ||
| (AddressSpace == access::address_space::global_space || | ||
| AddressSpace == access::address_space::generic_space), | ||
| void> | ||
| joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr, | ||
| size_t bytes, Properties properties = {}) { | ||
| joint_prefetch(g, ptr.get(), bytes, properties); | ||
| } | ||
|
|
||
| // Only available if AddressSpace == global_space || AddressSpace == | ||
| // generic_space | ||
| template <typename Group, typename T, access::address_space AddressSpace, | ||
| access::decorated IsDecorated, | ||
| typename Properties = ext::oneapi::experimental::empty_properties_t> | ||
| typename std::enable_if_t< | ||
| sycl::is_group_v<std::decay_t<Group>> && | ||
| (AddressSpace == access::address_space::global_space || | ||
| AddressSpace == access::address_space::generic_space), | ||
| void> | ||
| joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr, | ||
| Properties properties = {}) { | ||
| joint_prefetch(g, ptr.get(), properties); | ||
| } | ||
|
|
||
| // Only available if AddressSpace == global_space || AddressSpace == | ||
| // generic_space | ||
| template <typename Group, typename T, access::address_space AddressSpace, | ||
| access::decorated IsDecorated, | ||
| typename Properties = ext::oneapi::experimental::empty_properties_t> | ||
| typename std::enable_if_t< | ||
| sycl::is_group_v<std::decay_t<Group>> && | ||
| (AddressSpace == access::address_space::global_space || | ||
| AddressSpace == access::address_space::generic_space), | ||
| void> | ||
| joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr, | ||
| size_t count, Properties properties = {}) { | ||
| joint_prefetch(g, ptr.get(), count, properties); | ||
| } | ||
|
|
||
| // Only available if Dimensions > 0 && (AccessMode == read || AccessMode == | ||
| // read_write) | ||
| template <typename Group, typename DataT, int Dimensions, | ||
| access_mode AccessMode, access::placeholder IsPlaceholder, | ||
| typename Properties = ext::oneapi::experimental::empty_properties_t> | ||
| typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> && | ||
| (Dimensions > 0) && | ||
| (AccessMode == access_mode::read || | ||
| AccessMode == access_mode::read_write), | ||
| void> | ||
| joint_prefetch( | ||
| Group g, | ||
| accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc, | ||
| size_t offset, Properties properties = {}) { | ||
| joint_prefetch(g, (void *)&acc[offset], sizeof(DataT), properties); | ||
| } | ||
|
|
||
| // Only available if Dimensions > 0 && (AccessMode == read || AccessMode == | ||
| // read_write) | ||
| template <typename Group, typename DataT, int Dimensions, | ||
| access_mode AccessMode, access::placeholder IsPlaceholder, | ||
| typename Properties = ext::oneapi::experimental::empty_properties_t> | ||
| typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> && | ||
| (Dimensions > 0) && | ||
| (AccessMode == access_mode::read || | ||
| AccessMode == access_mode::read_write), | ||
| void> | ||
| joint_prefetch( | ||
| Group g, | ||
| accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc, | ||
| size_t offset, size_t count, Properties properties = {}) { | ||
| joint_prefetch(g, (void *)&acc[offset], count * sizeof(DataT), properties); | ||
| } | ||
|
|
||
| } // namespace ext::oneapi::experimental | ||
| } // namespace _V1 | ||
| } // namespace sycl | ||
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,39 @@ | ||
| // RUN: %clangxx -fsycl-device-only -S %s -o - | FileCheck %s | ||
|
|
||
| #include <sycl/sycl.hpp> | ||
|
|
||
| char data[] = {0, 1, 2, 3}; | ||
|
|
||
| // CHECK: [[PREFETCH_STR:@.*]] = private unnamed_addr addrspace(1) constant [19 x i8] c"sycl-prefetch-hint\00", section "llvm.metadata" | ||
|
There was a problem hiding this comment. Are these manually written? If so, do we have an utility script to generate them automatically? If yes, one might need to disable instrumentation/use |
||
| // CHECK: [[PREFETCH_LVL0:@.*]] = private unnamed_addr addrspace(1) constant [2 x i8] c"0\00", section "llvm.metadata" | ||
| // CHECK: [[ANNOTATION1:@.*]] = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) [[PREFETCH_STR]], ptr addrspace(1) [[PREFETCH_LVL0]] }, section "llvm.metadata" | ||
| // CHECK: [[PREFETCH_LVL1:@.*]] = private unnamed_addr addrspace(1) constant [2 x i8] c"1\00", section "llvm.metadata" | ||
| // CHECK: [[ANNOTATION2:@.*]] = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) [[PREFETCH_STR]], ptr addrspace(1) [[PREFETCH_LVL1]] }, section "llvm.metadata" | ||
| // CHECK: [[PREFETCH_STR_NT:@.*]] = private unnamed_addr addrspace(1) constant [22 x i8] c"sycl-prefetch-hint-nt\00", section "llvm.metadata" | ||
| // CHECK: [[PREFETCH_LVL2:@.*]] = private unnamed_addr addrspace(1) constant [2 x i8] c"2\00", section "llvm.metadata" | ||
| // CHECK: [[ANNOTATION3:@.*]] = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) [[PREFETCH_STR_NT]], ptr addrspace(1) [[PREFETCH_LVL2]] }, section "llvm.metadata" | ||
|
|
||
| int main() { | ||
| namespace syclex = sycl::ext::oneapi::experimental; | ||
| sycl::queue q; | ||
| void *dataPtr = &data; | ||
| q.single_task([=]() { | ||
| // CHECK: [[CASTED:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobal{{.*}} | ||
|
|
||
| // CHECK: [[ANNOTATED1:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], {{.*}}, ptr addrspace(1) [[ANNOTATION1]]) | ||
| // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED1]], i64 noundef 1) | ||
| syclex::prefetch(dataPtr); | ||
|
|
||
| // CHECK: [[ANNOTATED2:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], {{.*}}, ptr addrspace(1) [[ANNOTATION2]]) | ||
| // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED2]], i64 noundef 1) | ||
| syclex::prefetch(dataPtr, syclex::properties{syclex::prefetch_hint_L2}); | ||
|
|
||
| // CHECK: [[ANNOTATED3:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], {{.*}}, ptr addrspace(1) [[ANNOTATION3]]) | ||
| // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED3]], i64 noundef 4) | ||
| syclex::prefetch(dataPtr, 4, | ||
| syclex::properties{syclex::prefetch_hint_L3_nt}); | ||
| }); | ||
| q.wait(); | ||
|
|
||
| return 0; | ||
| } | ||
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This doesn't feel right to me, but it's outside of SYCL RT...
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Looks like currently it's the only way to transform these properties into spirv decorations for this pointer. I tried to create
spirv.Decorationsmetadata instead, but the compiler eliminates them with optimization flags.Annotation also may be eliminated, but it's much less likely.