Local accessor. More...

#include <sycl/accessor.hpp>

Inheritance diagram for sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >:

[legend]

Collaboration diagram for sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >:

[legend]

Public Types
using	value_type = DataT

using	reference = DataT &

using	const_reference = const DataT &

Public Member Functions
	local_accessor_base ()

template<int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>
	local_accessor_base (handler &, const detail::code_location CodeLoc=detail::code_location::current())

template<int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>
	local_accessor_base (handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())

template<int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
	local_accessor_base (range< Dimensions > AllocationSize, handler &, const detail::code_location CodeLoc=detail::code_location::current())

template<int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
	local_accessor_base (range< Dimensions > AllocationSize, handler &, const property_list &propList, const detail::code_location CodeLoc=detail::code_location::current())

size_t	get_size () const

	__SYCL2020_DEPRECATED ("get_count() is deprecated, please use size() instead") size_t get_count() const

size_t	size () const noexcept

template<int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
range< Dims >	get_range () const

template<int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0 && IsAccessAnyWrite>>
	operator RefType () const

template<int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0) && IsAccessAnyWrite>>
RefType	operator[] (id< Dimensions > Index) const

template<int Dims = Dimensions, typename = detail::enable_if_t<Dims == 1 && IsAccessAnyWrite>>
RefType	operator[] (size_t Index) const

template<int Dims = Dimensions>
	operator typename detail::enable_if_t< Dims==0 &&AccessMode==access::mode::atomic, atomic< DataT, AS >> () const

template<int Dims = Dimensions>
detail::enable_if_t<(Dims > 0) &&AccessMode==access::mode::atomic, atomic< DataT, AS > >	operator[] (id< Dimensions > Index) const

template<int Dims = Dimensions>
detail::enable_if_t< Dims==1 &&AccessMode==access::mode::atomic, atomic< DataT, AS > >	operator[] (size_t Index) const

template<int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
AccessorCommonT::template AccessorSubscript< Dims - 1, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > >	operator[] (size_t Index) const

bool	operator== (const local_accessor_base &Rhs) const

bool	operator!= (const local_accessor_base &Rhs) const

Public Member Functions inherited from sycl::_V1::detail::LocalAccessorBaseHost
	LocalAccessorBaseHost (sycl::range< 3 > Size, int Dims, int ElemSize, const property_list &PropertyList={})

sycl::range< 3 > &	getSize ()

const sycl::range< 3 > &	getSize () const

void *	getPtr ()

void *	getPtr () const

int	getNumOfDims ()

int	getElementSize ()

const property_list &	getPropList () const

Protected Types
using	AccessorCommonT = detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >

template<int Dims>
using	AccessorSubscript = typename AccessorCommonT::template AccessorSubscript< Dims, local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > >

using	ConcreteASPtrType = typename detail::DecoratedType< DataT, AS >::type *

using	RefType = detail::const_if_const_AS< AS, DataT > &

using	PtrType = detail::const_if_const_AS< AS, DataT > *

Protected Types inherited from sycl::_V1::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >
using	RefType = detail::const_if_const_AS< AS, DataT > &

using	ConstRefType = const DataT &

using	PtrType = detail::const_if_const_AS< AS, DataT > *

Protected Member Functions
	local_accessor_base (const detail::LocalAccessorImplPtr &Impl)

PtrType	getQualifiedPtr () const

void *	getPtr ()

void *	getPtr () const

const range< 3 > &	getSize () const

range< 3 > &	getSize ()

void	GDBMethodsAnchor ()

template<int Dims = AdjustedDim>
size_t	getLinearIndex (id< Dims > Id) const

Protected Member Functions inherited from sycl::_V1::detail::LocalAccessorBaseHost
	LocalAccessorBaseHost (const LocalAccessorImplPtr &Impl)

template<class Obj >
decltype(Obj::impl) friend	getSyclObjImpl (const Obj &SyclObject)

Protected Attributes
char	padding [sizeof(detail::LocalAccessorBaseDevice< AdjustedDim >)+sizeof(PtrType) - sizeof(detail::LocalAccessorBaseHost)]

Protected Attributes inherited from sycl::_V1::detail::LocalAccessorBaseHost
LocalAccessorImplPtr	impl

Static Protected Attributes
constexpr static int	AdjustedDim = Dimensions == 0 ? 1 : Dimensions

static constexpr bool	IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite

static constexpr bool	IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly

static constexpr bool	IsConst = AccessorCommonT::IsConst

Static Protected Attributes inherited from sycl::_V1::detail::accessor_common< DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder >
constexpr static access::address_space	AS

constexpr static bool	IsHostBuf

constexpr static bool	IsHostTask

constexpr static bool	IsPlaceH

constexpr static bool	IsGlobalBuf

constexpr static bool	IsConstantBuf

constexpr static bool	IsAccessAnyWrite

constexpr static bool	IsAccessReadOnly

static constexpr bool	IsConst

constexpr static bool	IsAccessReadWrite

constexpr static bool	IsAccessAtomic

Friends
template<class T >
T	detail::createSyclObjFromImpl (decltype(T::impl) ImplObj)

Detailed Description

template<typename DataT, int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>
class sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >

Local accessor.

Definition at line 2430 of file accessor.hpp.

Member Typedef Documentation

◆ AccessorCommonT

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

using sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::AccessorCommonT = detail::accessor_common<DataT, Dimensions, AccessMode, access::target::local, IsPlaceholder>

protected

Definition at line 2441 of file accessor.hpp.

◆ AccessorSubscript

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims>

using sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::AccessorSubscript = typename AccessorCommonT::template AccessorSubscript< Dims, local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder> >

protected

Definition at line 2455 of file accessor.hpp.

◆ ConcreteASPtrType

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

using sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::ConcreteASPtrType = typename detail::DecoratedType<DataT, AS>::type *

protected

Definition at line 2457 of file accessor.hpp.

◆ const_reference

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

using sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::const_reference = const DataT &

Definition at line 2537 of file accessor.hpp.

◆ PtrType

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

using sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::PtrType = detail::const_if_const_AS<AS, DataT> *

protected

Definition at line 2460 of file accessor.hpp.

◆ reference

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

using sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::reference = DataT &

Definition at line 2536 of file accessor.hpp.

◆ RefType

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

using sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::RefType = detail::const_if_const_AS<AS, DataT> &

protected

Definition at line 2459 of file accessor.hpp.

◆ value_type

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

using sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::value_type = DataT

Definition at line 2535 of file accessor.hpp.

Constructor & Destructor Documentation

◆ local_accessor_base() [1/6]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::local_accessor_base ( )

inline

Definition at line 2487 of file accessor.hpp.

◆ local_accessor_base() [2/6]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::local_accessor_base ( const detail::LocalAccessorImplPtr & Impl )

inlineprotected

Definition at line 2492 of file accessor.hpp.

◆ local_accessor_base() [3/6]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>

sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::local_accessor_base	(	handler &	,
		const detail::code_location	CodeLoc = `detail::code_location::current()`
	)

inline

Definition at line 2540 of file accessor.hpp.

◆ local_accessor_base() [4/6]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>

sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::local_accessor_base	(	handler &	,
		const property_list &	propList,
		const detail::code_location	CodeLoc = `detail::code_location::current()`
	)

inline

Definition at line 2554 of file accessor.hpp.

◆ local_accessor_base() [5/6]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>

sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::local_accessor_base	(	range< Dimensions >	AllocationSize,
		handler &	,
		const detail::code_location	CodeLoc = `detail::code_location::current()`
	)

inline

Definition at line 2571 of file accessor.hpp.

◆ local_accessor_base() [6/6]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>

sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::local_accessor_base	(	range< Dimensions >	AllocationSize,
		handler &	,
		const property_list &	propList,
		const detail::code_location	CodeLoc = `detail::code_location::current()`
	)

inline

Definition at line 2587 of file accessor.hpp.

Member Function Documentation

◆ __SYCL2020_DEPRECATED()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::__SYCL2020_DEPRECATED	(	"get_count() is	deprecated,
		please use size() instead"
	)		const

inline

Definition at line 2606 of file accessor.hpp.

◆ GDBMethodsAnchor()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

void sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::GDBMethodsAnchor ( )

inlineprotected

Definition at line 2511 of file accessor.hpp.

◆ get_range()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>

range<Dims> sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::get_range ( ) const

inline

Definition at line 2611 of file accessor.hpp.

◆ get_size()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

size_t sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::get_size ( ) const

inline

Definition at line 2604 of file accessor.hpp.

◆ getLinearIndex()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = AdjustedDim>

size_t sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::getLinearIndex ( id< Dims > Id ) const

inlineprotected

Definition at line 2524 of file accessor.hpp.

◆ getPtr() [1/2]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

void* sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::getPtr ( )

inlineprotected

Definition at line 2503 of file accessor.hpp.

◆ getPtr() [2/2]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

void* sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::getPtr ( ) const

inlineprotected

Definition at line 2504 of file accessor.hpp.

◆ getQualifiedPtr()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

PtrType sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::getQualifiedPtr ( ) const

inlineprotected

Definition at line 2499 of file accessor.hpp.

◆ getSize() [1/2]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

range<3>& sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::getSize ( )

inlineprotected

Definition at line 2508 of file accessor.hpp.

◆ getSize() [2/2]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

const range<3>& sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::getSize ( ) const

inlineprotected

Definition at line 2505 of file accessor.hpp.

◆ operator RefType()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0 && IsAccessAnyWrite>>

sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::operator RefType ( ) const

inline

Definition at line 2617 of file accessor.hpp.

◆ operator typename detail::enable_if_t< Dims==0 &&AccessMode==access::mode::atomic, atomic< DataT, AS >>()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions>

sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::operator typename detail::enable_if_t< Dims==0 &&AccessMode==access::mode::atomic, atomic< DataT, AS >> ( ) const

inline

Definition at line 2636 of file accessor.hpp.

◆ operator!=()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

bool sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::operator!= ( const local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > & Rhs ) const

inline

Definition at line 2670 of file accessor.hpp.

◆ operator==()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

bool sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::operator== ( const local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder > & Rhs ) const

inline

Definition at line 2667 of file accessor.hpp.

References sycl::_V1::detail::LocalAccessorBaseHost::impl.

◆ operator[]() [1/5]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0) && IsAccessAnyWrite>>

RefType sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::operator[] ( id< Dimensions > Index ) const

inline

Definition at line 2623 of file accessor.hpp.

References sycl::_V1::detail::getLinearIndex().

◆ operator[]() [2/5]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions>

detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic, atomic<DataT, AS> > sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::operator[] ( id< Dimensions > Index ) const

inline

Definition at line 2645 of file accessor.hpp.

References sycl::_V1::detail::getLinearIndex().

◆ operator[]() [3/5]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions, typename = detail::enable_if_t<Dims == 1 && IsAccessAnyWrite>>

RefType sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::operator[] ( size_t Index ) const

inline

Definition at line 2630 of file accessor.hpp.

◆ operator[]() [4/5]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions>

detail::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic, atomic<DataT, AS> > sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::operator[] ( size_t Index ) const

inline

Definition at line 2654 of file accessor.hpp.

◆ operator[]() [5/5]

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>

AccessorCommonT::template AccessorSubscript< Dims - 1, local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder> > sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::operator[] ( size_t Index ) const

inline

Definition at line 2663 of file accessor.hpp.

◆ size()

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

size_t sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::size ( ) const

inlinenoexcept

Definition at line 2608 of file accessor.hpp.

Friends And Related Function Documentation

◆ detail::createSyclObjFromImpl

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

template<class T >

T detail::createSyclObjFromImpl ( decltype(T::impl) ImplObj )

friend

Member Data Documentation

◆ AdjustedDim

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

constexpr static int sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::AdjustedDim = Dimensions == 0 ? 1 : Dimensions

staticconstexprprotected

Definition at line 2437 of file accessor.hpp.

◆ IsAccessAnyWrite

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

constexpr bool sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite

staticconstexprprotected

Definition at line 2447 of file accessor.hpp.

◆ IsAccessReadOnly

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

constexpr bool sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly

staticconstexprprotected

Definition at line 2448 of file accessor.hpp.

◆ IsConst

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

constexpr bool sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::IsConst = AccessorCommonT::IsConst

staticconstexprprotected

Definition at line 2449 of file accessor.hpp.

◆ padding

template<typename DataT , int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>

char sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >::padding[sizeof(detail::LocalAccessorBaseDevice< AdjustedDim >)+ sizeof(PtrType) - sizeof(detail::LocalAccessorBaseHost)]

protected

Definition at line 2496 of file accessor.hpp.

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

include/sycl/accessor.hpp

Public Types

Public Member Functions

Protected Types

Protected Member Functions

Protected Attributes

Static Protected Attributes

Friends

Detailed Description

template<typename DataT, int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder> class sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >

Member Typedef Documentation

◆ AccessorCommonT

◆ AccessorSubscript

◆ ConcreteASPtrType

◆ const_reference

◆ PtrType

◆ reference

◆ RefType

◆ value_type

Constructor & Destructor Documentation

◆ local_accessor_base() [1/6]

◆ local_accessor_base() [2/6]

◆ local_accessor_base() [3/6]

◆ local_accessor_base() [4/6]

◆ local_accessor_base() [5/6]

◆ local_accessor_base() [6/6]

Member Function Documentation

◆ __SYCL2020_DEPRECATED()

◆ GDBMethodsAnchor()

◆ get_range()

◆ get_size()

◆ getLinearIndex()

◆ getPtr() [1/2]

◆ getPtr() [2/2]

◆ getQualifiedPtr()

◆ getSize() [1/2]

◆ getSize() [2/2]

◆ operator RefType()

◆ operator typename detail::enable_if_t< Dims==0 &&AccessMode==access::mode::atomic, atomic< DataT, AS >>()

◆ operator!=()

◆ operator==()

◆ operator[]() [1/5]

◆ operator[]() [2/5]

◆ operator[]() [3/5]

◆ operator[]() [4/5]

◆ operator[]() [5/5]

◆ size()

Friends And Related Function Documentation

◆ detail::createSyclObjFromImpl

Member Data Documentation

◆ AdjustedDim

◆ IsAccessAnyWrite

◆ IsAccessReadOnly

◆ IsConst

◆ padding

template<typename DataT, int Dimensions, access::mode AccessMode, access::placeholder IsPlaceholder>
class sycl::_V1::local_accessor_base< DataT, Dimensions, AccessMode, IsPlaceholder >