[SYCL][Doc] Add specialization constant-length alloca extension proposal

[victor-eds]

Document extension proposal for specialization constant length private memory allocations. Users will be able to perform dynamic memory allocations using specialization constants and new private_alloca and aligned_private_alloca functions returning a private_ptr to a memory region of automatic storage duration.

This is included as an experimental extension as implementation will shortly follow once the extension is approved.

[victor-eds]

@v-klochkov proposed an additional variant of private_alloca receiving a required alignment. This would be a great addition to the current API. I was thinking on something like:

namespace sycl::ext::oneapi::experimental {
template <typename ElementType, auto &SpecName,
          access::decorated DecorateAddress, std::size_t Alignment>
private_ptr<ElementType, DecorateAddress>
aligned_private_alloca(kernel_handler &h);
} // namespace sycl::ext::oneapi::experimental

i.e., receiving the alignment as an additional template argument. This goes in line with other functions in C++ and SYCL like sycl::aligned_alloc_device or std::aligned_alloc. Requiring alignment to be constant may greatly simplify implementation. The Clang builtin __builtin_alloca_with_align you mention also requires that.

[v-klochkov]

@v-klochkov proposed an additional variant of private_alloca receiving a required alignment. This would be a great addition to the current API. I was thinking on something like:

namespace sycl::ext::oneapi::experimental {
template <typename ElementType, auto &SpecName,
          access::decorated DecorateAddress, std::size_t Alignment>
private_ptr<ElementType, DecorateAddress>
aligned_private_alloca(kernel_handler &h);
} // namespace sycl::ext::oneapi::experimental

i.e., receiving the alignment as an additional template argument. This goes in line with other functions in C++ and SYCL like sycl::aligned_alloc_device or std::aligned_alloc. Requiring alignment to be constant may greatly simplify implementation. The Clang builtin __builtin_alloca_with_align you mention also requires that.

@victor-eds What is the default alignment for non aligned version of the function?
If have the extra template parameter Alignment, which goes last in order, then why not to re-use the original name and set the default value for Alignment: e.g.:

template <typename ElementType, access::decorated DecorateAddress, std::size_t Alignment = sizeof(ElementType)>
private_ptr<ElementType, DecorateAddress>
private_alloca(const specialization_constant<std::size_t> &size);

[victor-eds]

@v-klochkov proposed an additional variant of private_alloca receiving a required alignment. This would be a great addition to the current API. I was thinking on something like:

namespace sycl::ext::oneapi::experimental {
template <typename ElementType, auto &SpecName,
          access::decorated DecorateAddress, std::size_t Alignment>
private_ptr<ElementType, DecorateAddress>
aligned_private_alloca(kernel_handler &h);
} // namespace sycl::ext::oneapi::experimental

i.e., receiving the alignment as an additional template argument. This goes in line with other functions in C++ and SYCL like sycl::aligned_alloc_device or std::aligned_alloc. Requiring alignment to be constant may greatly simplify implementation. The Clang builtin __builtin_alloca_with_align you mention also requires that.

@victor-eds What is the default alignment for non aligned version of the function? If have the extra template parameter Alignment, which goes last in order, then why not to re-use the original name and set the default value for Alignment: e.g.:

template <typename ElementType, access::decorated DecorateAddress, std::size_t Alignment = sizeof(ElementType)>
private_ptr<ElementType, DecorateAddress>
private_alloca(const specialization_constant<std::size_t> &size);

Default alignment should be alignof(ElementType) (I'll mention in the extension, thanks for reminding).

Having a second aligned_private_alloca function might be more future-proof, especially when a default for DecorateAddress is agreed on. Something like:

template <typename ElementType, std::size_t Alignment, auto &SpecName, 
    access::decorated DecorateAddress>
private_ptr<ElementType, DecorateAdress>
aligned_private_alloca(kernel_handler &kh);

Can be more easily transformed to:

template <typename ElementType, std::size_t Alignment, auto &SpecName, 
    access::decorated DecorateAddress = /*default*/>
private_ptr<ElementType, DecorateAdress>
aligned_private_alloca(kernel_handler &kh);

WDYT?

[victor-eds]

@v-klochkov I added the second signature. Feel free to comment.

[victor-eds]

@gmlueck @v-klochkov

This PR has been approved by the required reviewers, but, as you also provided feedback, can you please confirm current status looks good to you? Please, resolve the open alignment conversation if so. Thanks!

[bader]

@victor-eds, could you change the file name from sycl_ext_private_alloca.asciidoc to sycl_ext_oneapi_private_alloca.asciidoc, please?

[victor-eds]

@victor-eds, could you change the file name from sycl_ext_private_alloca.asciidoc to sycl_ext_oneapi_private_alloca.asciidoc, please?

Done (#12881). Thanks for pointing out!