From 4959401ee77e05a647ba93537baa74d60396a6ba Mon Sep 17 00:00:00 2001 From: Christopher Di Bella Date: Tue, 11 Dec 2018 17:13:03 +0000 Subject: [PATCH 1/4] Added use-onchip-memory sample. --- samples/use-onchip-memory/CMakeLists.txt | 21 +++ .../use-onchip-memory/use-onchip-memory.cpp | 142 ++++++++++++++++++ 2 files changed, 163 insertions(+) create mode 100644 samples/use-onchip-memory/CMakeLists.txt create mode 100644 samples/use-onchip-memory/use-onchip-memory.cpp diff --git a/samples/use-onchip-memory/CMakeLists.txt b/samples/use-onchip-memory/CMakeLists.txt new file mode 100644 index 0000000..9dfd53e --- /dev/null +++ b/samples/use-onchip-memory/CMakeLists.txt @@ -0,0 +1,21 @@ +set(SOURCE_NAME "use-onchip-memory") + +add_executable( + ${SOURCE_NAME} + ${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_NAME}.cpp +) +include_directories( + ${PROJECT_SOURCE_DIR}/include +) +add_sycl_to_target( + TARGET ${SOURCE_NAME} + SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_NAME}.cpp +) +add_test( + NAME ${SOURCE_NAME} + COMMAND ${SOURCE_NAME} +) +install( + TARGETS ${SOURCE_NAME} + RUNTIME DESTINATION bin +) diff --git a/samples/use-onchip-memory/use-onchip-memory.cpp b/samples/use-onchip-memory/use-onchip-memory.cpp new file mode 100644 index 0000000..6c41b83 --- /dev/null +++ b/samples/use-onchip-memory/use-onchip-memory.cpp @@ -0,0 +1,142 @@ +/*************************************************************************** + * + * Copyright Codeplay Software Limited + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * For your convenience, a copy of the License has been included in this + * repository. + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * Codeplay's ComputeCpp SDK + * + * example_vptr.cpp + * + * Description: + * Sample code that demonstrates the use of the use_onchip_memory extension + * to SYCL provided by ComputeCpp. + * + **************************************************************************/ + +#include +#include +#include + +namespace sycl = cl::sycl; +namespace access = sycl::access; + +// Sample kernel +// ------------- +// Since kernels are no longer allowed to be declared inline, I dislike +// declaring a random name in the global namespace that can potentially be +// repurposed. +// +// This kernel performs the same operation as std::iota, but also scales +// the result by two. +// +class scaled_iota { + public: + using accessor = sycl::accessor; + + explicit scaled_iota(accessor deviceAccess) noexcept + : m_deviceAccess{std::move(deviceAccess)} {} + + void operator()(sycl::nd_item<2> itemId) const noexcept { + const auto linearId = itemId.get_global_linear_id(); + m_deviceAccess[linearId] = linearId * 2; + } + + private: + accessor m_deviceAccess; +}; + +namespace codeplay = sycl::codeplay; + +template +void use_with_policy(Policy policy, sycl::queue& queue) { + auto hostData = sycl::vector_class(1024); + { + auto taskContext = queue.get_context(); + + // Notice that the Codeplay property takes a policy argument: this is used + // to indicate whether the property is advantageous or genuinely necessary. + // + auto deviceData = sycl::buffer{ + hostData.data(), // + sycl::range<1>(hostData.size()), // + sycl::property_list{ + sycl::property::buffer::context_bound(taskContext), + codeplay::property::buffer::use_onchip_memory(policy) // <-------- + } // ^--------------------- + }; + + deviceData.set_final_data(hostData.data()); + + queue.submit([&](sycl::handler & cgh) { + constexpr auto dimension_size = 2; + auto r = sycl::nd_range{ + sycl::range{ + hostData.size() / dimension_size, // + dimension_size // + }, + sycl::range{dimension_size, 1} // + }; + cgh.parallel_for( + r, + scaled_iota(deviceData.get_access(cgh))); + }); + + queue.wait_and_throw(); + } +} + +// Codeplay policy extensions have two different enabling mechanisms: the first +// is to indicate that a policy is preferred. Using this policy means that if +// the system supports the feature, then the feature will be enabled. If the +// feature is not present on the system, then it will not be enabled. +// +// Puns aside, this is the preferred default. +// +void how_to_use_with_prefer() { + auto queue = sycl::queue{}; + ::use_with_policy(codeplay::property::prefer, queue); +} + +// Alternatively, if you can guarantee that your system will support this +// policy, or if it is expected any system using your software must support the +// policy, then you can use the require tag to indicate that the feature is +// required by your software. +// +// In the event that the property isn't supported, an exception will be thrown. +// +void how_to_use_with_require() { + try { + auto queue = sycl::queue{}; + ::use_with_policy(codeplay::property::require, queue); + } catch (const sycl::exception& e) { + std::cerr << "An error occurred: " << e.what() + << "\n" + "\n" + "This particular error has occurred because you are requiring " + "the policy use_onchip_memory be available, and your hardware " + "doesn't support the use_onchip_memory, so the SYCL ecosystem " + "will raise an error.\n"; + } +} + +int main() { + // Using the on-chip memory policy with the require tag. + how_to_use_with_require(); + + // Using the on-chip memory policy with the prefer tag. + how_to_use_with_prefer(); +} From b0834a42c98fbef5202f6a35dd8eebdd682a2228 Mon Sep 17 00:00:00 2001 From: Christopher Di Bella Date: Wed, 12 Dec 2018 13:33:17 +0000 Subject: [PATCH 2/4] Applied review changes --- .../use-onchip-memory/use-onchip-memory.cpp | 97 +++++++++---------- 1 file changed, 45 insertions(+), 52 deletions(-) diff --git a/samples/use-onchip-memory/use-onchip-memory.cpp b/samples/use-onchip-memory/use-onchip-memory.cpp index 6c41b83..c6b00aa 100644 --- a/samples/use-onchip-memory/use-onchip-memory.cpp +++ b/samples/use-onchip-memory/use-onchip-memory.cpp @@ -18,7 +18,7 @@ * * Codeplay's ComputeCpp SDK * - * example_vptr.cpp + * use-onchip-memory.cpp * * Description: * Sample code that demonstrates the use of the use_onchip_memory extension @@ -33,31 +33,13 @@ namespace sycl = cl::sycl; namespace access = sycl::access; -// Sample kernel -// ------------- -// Since kernels are no longer allowed to be declared inline, I dislike -// declaring a random name in the global namespace that can potentially be -// repurposed. -// +namespace sycl_kernel { // This kernel performs the same operation as std::iota, but also scales // the result by two. // -class scaled_iota { - public: - using accessor = sycl::accessor; - - explicit scaled_iota(accessor deviceAccess) noexcept - : m_deviceAccess{std::move(deviceAccess)} {} - - void operator()(sycl::nd_item<2> itemId) const noexcept { - const auto linearId = itemId.get_global_linear_id(); - m_deviceAccess[linearId] = linearId * 2; - } - - private: - accessor m_deviceAccess; -}; +template +class scaled_iota; +} // namespace sycl_kernel namespace codeplay = sycl::codeplay; @@ -67,60 +49,69 @@ void use_with_policy(Policy policy, sycl::queue& queue) { { auto taskContext = queue.get_context(); - // Notice that the Codeplay property takes a policy argument: this is used - // to indicate whether the property is advantageous or genuinely necessary. + // clang-format off + // + // Notice that the on_chip_memory property takes a policy argument: this is + // used to indicate whether the property is advantageous or genuinely + // necessary. // auto deviceData = sycl::buffer{ - hostData.data(), // - sycl::range<1>(hostData.size()), // + hostData.data(), + sycl::range<1>(hostData.size()), sycl::property_list{ sycl::property::buffer::context_bound(taskContext), - codeplay::property::buffer::use_onchip_memory(policy) // <-------- - } // ^--------------------- + codeplay::property::buffer::use_onchip_memory(policy) + } }; + // clang-format on deviceData.set_final_data(hostData.data()); - queue.submit([&](sycl::handler & cgh) { + queue.submit([&](sycl::handler& cgh) { constexpr auto dimension_size = 2; + + // clang-format off auto r = sycl::nd_range{ - sycl::range{ - hostData.size() / dimension_size, // - dimension_size // - }, - sycl::range{dimension_size, 1} // + sycl::range{ + hostData.size() / dimension_size, + dimension_size + }, + sycl::range{dimension_size, 1} }; - cgh.parallel_for( - r, - scaled_iota(deviceData.get_access(cgh))); + cgh.parallel_for>( + r, + [access = deviceData.get_access(cgh)] + (sycl::nd_item<2> id) noexcept { + const auto linearId = id.get_global_linear_id(); + access[linearId] = linearId * 2; + }); + // clang-format on }); - queue.wait_and_throw(); } } -// Codeplay policy extensions have two different enabling mechanisms: the first -// is to indicate that a policy is preferred. Using this policy means that if -// the system supports the feature, then the feature will be enabled. If the +// use_onchip_memory has two different enabling mechanisms: the first is to +// indicate that a policy is preferred. Using this policy means that if the +// system supports the feature, then the feature will be enabled. If the // feature is not present on the system, then it will not be enabled. // // Puns aside, this is the preferred default. // -void how_to_use_with_prefer() { - auto queue = sycl::queue{}; +void how_to_use_with_prefer(sycl::queue& queue) { ::use_with_policy(codeplay::property::prefer, queue); } // Alternatively, if you can guarantee that your system will support this -// policy, or if it is expected any system using your software must support the -// policy, then you can use the require tag to indicate that the feature is -// required by your software. +// policy, or if it is expected any system using your software must support +// the policy, then you can use the require tag to indicate that the feature +// is required by your software. // -// In the event that the property isn't supported, an exception will be thrown. +// In the event that the property isn't supported, a SYCL exception will be +// thrown. // -void how_to_use_with_require() { +void how_to_use_with_require(sycl::queue& queue) { try { - auto queue = sycl::queue{}; ::use_with_policy(codeplay::property::require, queue); } catch (const sycl::exception& e) { std::cerr << "An error occurred: " << e.what() @@ -134,9 +125,11 @@ void how_to_use_with_require() { } int main() { + auto queue = sycl::queue{}; + // Using the on-chip memory policy with the require tag. - how_to_use_with_require(); + how_to_use_with_require(queue); // Using the on-chip memory policy with the prefer tag. - how_to_use_with_prefer(); + how_to_use_with_prefer(queue); } From a199dd5eabbf9e141238df68cd1e9a5ef395bbe7 Mon Sep 17 00:00:00 2001 From: Christopher Di Bella Date: Thu, 13 Dec 2018 13:30:17 +0000 Subject: [PATCH 3/4] Applied feedback --- samples/use-onchip-memory/CMakeLists.txt | 3 --- samples/use-onchip-memory/use-onchip-memory.cpp | 4 ---- 2 files changed, 7 deletions(-) diff --git a/samples/use-onchip-memory/CMakeLists.txt b/samples/use-onchip-memory/CMakeLists.txt index 9dfd53e..8caa510 100644 --- a/samples/use-onchip-memory/CMakeLists.txt +++ b/samples/use-onchip-memory/CMakeLists.txt @@ -4,9 +4,6 @@ add_executable( ${SOURCE_NAME} ${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_NAME}.cpp ) -include_directories( - ${PROJECT_SOURCE_DIR}/include -) add_sycl_to_target( TARGET ${SOURCE_NAME} SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_NAME}.cpp diff --git a/samples/use-onchip-memory/use-onchip-memory.cpp b/samples/use-onchip-memory/use-onchip-memory.cpp index c6b00aa..62f761f 100644 --- a/samples/use-onchip-memory/use-onchip-memory.cpp +++ b/samples/use-onchip-memory/use-onchip-memory.cpp @@ -63,14 +63,10 @@ void use_with_policy(Policy policy, sycl::queue& queue) { codeplay::property::buffer::use_onchip_memory(policy) } }; - // clang-format on - - deviceData.set_final_data(hostData.data()); queue.submit([&](sycl::handler& cgh) { constexpr auto dimension_size = 2; - // clang-format off auto r = sycl::nd_range{ sycl::range{ hostData.size() / dimension_size, From 11eaab0f57ee339e4311868073c0971b5ac58a52 Mon Sep 17 00:00:00 2001 From: Christopher Di Bella Date: Fri, 14 Dec 2018 16:35:24 +0000 Subject: [PATCH 4/4] Update use-onchip-memory.cpp --- samples/use-onchip-memory/use-onchip-memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/samples/use-onchip-memory/use-onchip-memory.cpp b/samples/use-onchip-memory/use-onchip-memory.cpp index 62f761f..a6f1cd9 100644 --- a/samples/use-onchip-memory/use-onchip-memory.cpp +++ b/samples/use-onchip-memory/use-onchip-memory.cpp @@ -115,7 +115,7 @@ void how_to_use_with_require(sycl::queue& queue) { "\n" "This particular error has occurred because you are requiring " "the policy use_onchip_memory be available, and your hardware " - "doesn't support the use_onchip_memory, so the SYCL ecosystem " + "doesn't support the use_onchip_memory, so the SYCL implementation " "will raise an error.\n"; } }