StaticLocalStorage< T, size, condition > Class Template Reference

#include <gromacs/gpu_utils/sycl_kernel_utils.h>

Description

template<typename T, size_t size, bool condition = true>
class StaticLocalStorage< T, size, condition >

Shim for the two implementations of local storage.

A sycl::local_accessor works as a facet of a command-group handler, so when using handler-free submission, we must also use the extension for work-group static storage, which does not depend on a command-group handler.

Because we have the above two implementation flavours for SYCL device kernels, this class permits the code to specify the base type, buffer size, and condition under which it exists in one place, and then use that correctly for both implementations of local storage.

Note that the declaration of a sycl::local_accessor must be in the host code, whereas the declaration of a work_group_static must be in the kernel code.

Public Types
using	Underlying = detail::UnderlyingType< T >
	Helper typedef.
using	WorkGroupStatic = std::nullptr_t
using	ActiveHostStorage = std::conditional_t< sc_useCommandGroupHandler, sycl::local_accessor< T, 1 >, std::nullptr_t >
	The type to declare local storage in host code, given the build configuration. More...
using	ActiveDeviceStorage = std::conditional_t< sc_useCommandGroupHandler, std::nullptr_t, WorkGroupStatic >
	The type to declare local storage in device code, given the build configuration. More...
using	HostStorage = std::conditional_t< condition, ActiveHostStorage, std::nullptr_t >
	Type of local storage in the host code, given condition.
using	DeviceStorage = std::conditional_t< condition, ActiveDeviceStorage, std::nullptr_t >
	Type of local storage in the device code, given condition.
using	LocalPtr = sycl::multi_ptr< T, sycl::access::address_space::local_space >
	Helper typedef for get_pointer() implementation.

Static Public Member Functions
static HostStorage	makeHostStorage (sycl::handler &cgh)
	Helper function to make storage for the local_accessor case. More...
static HostStorage	makeHostStorage (std::nullptr_t)
	Helper function to make storage for the work_group_static case. More...
static LocalPtr	get_pointer (const sycl::local_accessor< T, 1 > &hostStorage, std::nullptr_t &) noexcept
	Provide a consistent way to get a sycl::local_ptr when using sycl::local_accessor. More...
static LocalPtr	get_pointer (const std::nullptr_t &, std::nullptr_t &) noexcept
	Provide a consistent way to get a sycl::local_ptr when using a work_group_static. More...

Member Typedef Documentation

template<typename T , size_t size, bool condition = true>

using StaticLocalStorage< T, size, condition >::ActiveDeviceStorage = std::conditional_t<sc_useCommandGroupHandler, std::nullptr_t, WorkGroupStatic>

The type to declare local storage in device code, given the build configuration.

With !sc_useCommandGroupHandler, this type (ie. work_group_static) can only be successfully declared in the device kernel.

template<typename T , size_t size, bool condition = true>

using StaticLocalStorage< T, size, condition >::ActiveHostStorage = std::conditional_t<sc_useCommandGroupHandler, sycl::local_accessor<T, 1>, std::nullptr_t>

The type to declare local storage in host code, given the build configuration.

With sc_useCommandGroupHandler, this type (i.e. sycl::local_accessor) can only be successfully declared in the host code.

Member Function Documentation

template<typename T , size_t size, bool condition = true>

static LocalPtr StaticLocalStorage< T, size, condition >::get_pointer ( const sycl::local_accessor< T, 1 > &  hostStorage, std::nullptr_t &    )

inlinestaticnoexcept

Provide a consistent way to get a sycl::local_ptr when using sycl::local_accessor.

Note that the sycl::local_accessor is created on the host and is const in the device kernel, but the sycl::local_ptr obtained from it can be either const or non-const according to need.

template<typename T , size_t size, bool condition = true>

static LocalPtr StaticLocalStorage< T, size, condition >::get_pointer ( const std::nullptr_t &  , std::nullptr_t &    )

inlinestaticnoexcept

Provide a consistent way to get a sycl::local_ptr when using a work_group_static.

Note that the work_group_static is created on the device and can be treated as either const or non-const according to need. But since the device kernels generally need read-write access, we only need the non-const overload.

Provide a consistent way to get a sycl::local_ptr containing nullptr when the storage is empty.

template<typename T , size_t size, bool condition = true>

static HostStorage StaticLocalStorage< T, size, condition >::makeHostStorage ( sycl::handler &  cgh )

inlinestatic

Helper function to make storage for the local_accessor case.

May only be called on the host.

Note that the equivalent approach for work_group_static compiles and fails at runtime, because the address of the local storage is invalid.

template<typename T , size_t size, bool condition = true>

static HostStorage StaticLocalStorage< T, size, condition >::makeHostStorage ( std::nullptr_t  )

inlinestatic

Helper function to make storage for the work_group_static case.

Should only be called on the host, to match its other implementation.

---

The documentation for this class was generated from the following file:

• src/gromacs/gpu_utils/sycl_kernel_utils.h