Skip to content

Commit

Permalink
Fix Windows compilation for on-chip memory sample
Browse files Browse the repository at this point in the history
Had to change `dimension_size` to `static size_t`, not `(automatic) auto`.
Windows requires that the variable be static, and also warns of a narrowing
conversion (`signed int` to `size_t`).
  • Loading branch information
ProGTX authored and DuncanMcBain committed Dec 17, 2018
1 parent 266b9e0 commit 2e61239
Showing 1 changed file with 16 additions and 32 deletions.
48 changes: 16 additions & 32 deletions samples/use-onchip-memory/use-onchip-memory.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,6 @@ namespace access = sycl::access;
namespace sycl_kernel {
// This kernel performs the same operation as std::iota, but also scales
// the result by two.
//
template <class>
class scaled_iota;
} // namespace sycl_kernel
Expand All @@ -49,39 +48,28 @@ void use_with_policy(Policy policy, sycl::queue& queue) {
{
auto taskContext = queue.get_context();

// 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<int, 1>{
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)}};

queue.submit([&](sycl::handler& cgh) {
constexpr auto dimension_size = 2;
static constexpr size_t dims = 2;

auto r = sycl::nd_range<dimension_size>{
sycl::range<dimension_size>{
hostData.size() / dimension_size,
dimension_size
},
sycl::range<dimension_size>{dimension_size, 1}
};
auto r =
sycl::nd_range<dims>{sycl::range<dims>{hostData.size() / dims, dims},
sycl::range<dims>{dims, 1}};
const auto access =
deviceData.get_access<access::mode::discard_write>(cgh);
cgh.parallel_for<sycl_kernel::scaled_iota<Policy>>(
r,
[access = deviceData.get_access<access::mode::discard_write>(cgh)]
(sycl::nd_item<2> id) noexcept {
const auto linearId = id.get_global_linear_id();
access[linearId] = linearId * 2;
});
// clang-format on
r, [=](sycl::nd_item<2> id) noexcept {
const auto linearId = id.get_global_linear_id();
access[linearId] = linearId * 2;
});
});
queue.wait_and_throw();
}
Expand All @@ -93,7 +81,6 @@ void use_with_policy(Policy policy, sycl::queue& queue) {
// 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(sycl::queue& queue) {
::use_with_policy(codeplay::property::prefer, queue);
}
Expand All @@ -105,18 +92,15 @@ void how_to_use_with_prefer(sycl::queue& queue) {
//
// In the event that the property isn't supported, a SYCL exception will be
// thrown.
//
void how_to_use_with_require(sycl::queue& queue) {
try {
::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 implementation "
"will raise an error.\n";
<< "\n\nThis 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 implementation will raise an error.\n";
}
}

Expand Down

0 comments on commit 2e61239

Please sign in to comment.