From c18c5c548d670dd28a2218cde9cd4056198b0a35 Mon Sep 17 00:00:00 2001
From: Victor Perez <victor.perez@codeplay.com>
Date: Thu, 15 Feb 2024 13:48:50 +0000
Subject: [PATCH] Add storage duration example

Signed-off-by: Victor Perez <victor.perez@codeplay.com>
---
 ...neapi_spec_constant_length_alloca.asciidoc | 33 +++++++++++++++++++
 1 file changed, 33 insertions(+)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_spec_constant_length_alloca.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_spec_constant_length_alloca.asciidoc
index 89642fda12dd0..37d3ddf31fce8 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_spec_constant_length_alloca.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_spec_constant_length_alloca.asciidoc
@@ -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);
@@ -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