Skip to content

Commit

Permalink
Add storage duration example
Browse files Browse the repository at this point in the history
Signed-off-by: Victor Perez <victor.perez@codeplay.com>
  • Loading branch information
victor-eds committed Feb 15, 2024
1 parent 92be499 commit c18c5c5
Showing 1 changed file with 33 additions and 0 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -158,6 +158,8 @@ conditions results in undefined behaviour.

This non-normative section shows some example usages of the extension.

=== Basic Usage

[source,c++]
----
constexpr specialization_id<int> size(1);
Expand All @@ -179,6 +181,37 @@ void run(queue q, const float *in, float *out, size_t n) {
});
----

=== Storage Duration Clarification

The following example is intended to clarify storage duration of memory
allocated by `private_alloca`.

[source,c++]
----
constexpr specialization_id<int> size(1);
class Kernel;
SYCL_EXTERNAL void impl(const float *in, float *out, size_t n,
raw_private_ptr<float> ptr);
void run(queue q, const float *in, float *out, size_t n) {
q.submit([&](handler &h) {
h.set_specialization_constant<size>(n);
h.parallel_for<Kernel>(n, [=](id<1> i, kernel_handler kh) {
raw_private_ptr<float> ptr;
{
ptr = private_alloca<float, size, access::decorated::no>(kh);
// 'private_alloca' has allocated a private memory region we can use in
// this block.
impl(in, out, kh.get_specialization_constant<size>(), ptr);
}
// Memory allocated by 'private_alloca' has been deallocated.
// Dereferencing 'ptr' at this program point is undefined behaviour.
});
});
----

== Design Constraints

The big design constraint stems from the unknown allocation size at compile
Expand Down

0 comments on commit c18c5c5

Please sign in to comment.