Address-Space Qualification

Address space restriction specifiers for pointers. More...

Classes
struct	sp::get_addrspace< T >
	Get the address space of a type. More...
struct	sp::add_addrspace< T, AS >
	Add an address space qualifier to a type. More...
struct	sp::remove_addrspace< T >
	Remove address space qualifiers from a type. More...
struct	sp::remove_cva< T >
	Remove const, volatile, and address space from a type. More...
struct	sp::remove_cvaref< T >
	Remove reference const, volatile, and address space from a type. More...

Macros
#define	__AS(X)   __attribute__((address_space(X)))
	Handy macro for adding an address-space specifier to a type. More...
#define	__flat
	Address space qualifier for a pointer that is compatible with any address space. More...
#define	__device
	Address space qualifier for a pointer to GPU main memory (AKA "global" memory). More...
#define	__shared
	Address space qualifier for a pointer to shared memory. More...
#define	__constant
	Address space qualifier for a pointer to constant memory. More...
#define	__local
	Address space qualifier for a pointer to local (AKA "stack" or "private") memory. More...

Typedefs
template<typename T , AddressSpace AS>
using	sp::add_addrspace_t = typename add_addrspace< T, AS >::type
	Add an address space qualifier to a type. More...
template<typename T >
using	sp::remove_addrspace_t = typename remove_addrspace< T >::type
	Remove address space qualifiers from a type. More...
template<typename T >
using	sp::remove_cva_t = typename remove_cva< T >::type
	Remove const, volatile, and address space from a type. More...
template<typename T >
using	sp::remove_cvaref_t = typename remove_cvaref< T >::type
	Remove reference const, volatile, and address space from a type. More...

Enumerations
enum	sp::AddressSpace {   sp::FLAT = __scale_address_space_flat , sp::GENERIC = __scale_address_space_generic , sp::DEVICE = __scale_address_space_device , sp::SHARED = __scale_address_space_shared ,   sp::CONSTANT = __scale_address_space_constant , sp::LOCAL = __scale_address_space_local }
	Enum providing the address-space numbers for GPUs. More...

Variables
template<typename T >
constexpr sp::AddressSpace	sp::get_addrspace_v = get_addrspace<T>::value
	Get the address space of a type. More...

Detailed Description

Address space restriction specifiers for pointers.

You'll definitely want to read the corresponding portion of the manual

These allow you to record which memory a pointer is allowed to point to as part of its type. For instance, a __shared int* is a pointer to an int in shared memory.

__device__ functionThatExpectsSharedMemoryPointer(__shared float* x) {
    // Do something
}
 
__device__ void foo(__device float* input) {
    functionThatExpectsSharedMemoryPointer(input);  // <- Compile error.
}

This mechanism means the compiler can stop you from making common mistakes like using a host pointer when a device pointer is required. There are also modest performance wins to be had:

• Your code will become less reliant on address-space inference to avoid expensive device-side generic memory instructions
• LibSCALE APIs (such as host<->device copies) can use overload resolution to statically determine what they should do, instead of requiring you to pass an enum they switch on at runtime (or perform a surprisingly expensive pointer range check).

Macro Definition Documentation

__AS

#define __AS

(

X

)

__attribute__((address_space(X)))

Handy macro for adding an address-space specifier to a type.

__constant

#define __constant

Address space qualifier for a pointer to constant memory.

__device

#define __device

Address space qualifier for a pointer to GPU main memory (AKA "global" memory).

Pointers present in kernel arguments are always implicitly __device (regardless of how deeply nested in structs they might be).

__flat

#define __flat

Address space qualifier for a pointer that is compatible with any address space.

Can convert to any other address space, whether or not it can be dereferenced on the current target (i.e: host or device). On the host, this allows a distinction between a host T * and an "anywhere" __flat T *.

__local

#define __local

Address space qualifier for a pointer to local (AKA "stack" or "private") memory.

Note

Local variables do not automatically get this address space, but can be qualified with __local.

__shared

#define __shared

Address space qualifier for a pointer to shared memory.

Typedef Documentation

add_addrspace_t

template<typename T , AddressSpace AS>

using sp::add_addrspace_t = typedef typename add_addrspace<T, AS>::type

Add an address space qualifier to a type.

Example

static_assert(std::is_same_v<__shared float, sp::add_addrspace_t<float, sp::AddressSpace::SHARED>>);
static_assert(std::is_same_v<__shared const float, sp::add_addrspace_t<const float, sp::AddressSpace::SHARED>>);
static_assert(std::is_same_v<float &, sp::add_addrspace_t<float &, sp::AddressSpace::DEVICE>>);
static_assert(std::is_same_v<const float &, sp::add_addrspace_t<const float &, sp::AddressSpace::DEVICE>>);
 
static_assert(std::is_same_v<float * __device, sp::add_addrspace_t<float *, sp::AddressSpace::DEVICE>>);
static_assert(std::is_same_v<const float * __device, sp::add_addrspace_t<const float *, sp::AddressSpace::DEVICE>>);
static_assert(std::is_same_v<float *, sp::add_addrspace_t<float *, sp::AddressSpace::GENERIC>>);
static_assert(std::is_same_v<const float *, sp::add_addrspace_t<const float *, sp::AddressSpace::GENERIC>>);
 
static_assert(std::is_same_v<float, sp::add_addrspace_t<float, sp::AddressSpace::GENERIC>>);
static_assert(std::is_same_v<const float, sp::add_addrspace_t<const float, sp::AddressSpace::GENERIC>>);
static_assert(std::is_same_v<float &, sp::add_addrspace_t<float &, sp::AddressSpace::GENERIC>>);
static_assert(std::is_same_v<const float &, sp::add_addrspace_t<const float &, sp::AddressSpace::GENERIC>>);

remove_addrspace_t

template<typename T >

using sp::remove_addrspace_t = typedef typename remove_addrspace<T>::type

Remove address space qualifiers from a type.

Example

All of the following are true.

Example

static_assert(std::is_same_v<float, sp::remove_addrspace_t<__shared float>>);
static_assert(std::is_same_v<const float, sp::remove_addrspace_t<__shared const float>>);
static_assert(std::is_same_v<__shared float &, sp::remove_addrspace_t<__shared float &>>);
static_assert(std::is_same_v<__shared const float &, sp::remove_addrspace_t<__shared const float &>>);
static_assert(std::is_same_v<float *, sp::remove_addrspace_t<float * __shared>>);
static_assert(std::is_same_v<const float *, sp::remove_addrspace_t<const float * __shared>>);
static_assert(std::is_same_v<__shared float *, sp::remove_addrspace_t<__shared float *>>);
static_assert(std::is_same_v<__shared const float *, sp::remove_addrspace_t<__shared const float *>>);
 
static_assert(std::is_same_v<float, sp::remove_addrspace_t<float>>);
static_assert(std::is_same_v<const float, sp::remove_addrspace_t<const float>>);
static_assert(std::is_same_v<float &, sp::remove_addrspace_t<float &>>);
static_assert(std::is_same_v<const float &, sp::remove_addrspace_t<const float &>>);
static_assert(std::is_same_v<float *, sp::remove_addrspace_t<float *>>);
static_assert(std::is_same_v<const float *, sp::remove_addrspace_t<const float *>>);

remove_cva_t

template<typename T >

using sp::remove_cva_t = typedef typename remove_cva<T>::type

Remove const, volatile, and address space from a type.

Example

static_assert(std::is_same_v<float, sp::remove_cva_t<__shared float>>);
static_assert(std::is_same_v<float, sp::remove_cva_t<__shared const float>>);
static_assert(std::is_same_v<__shared float &, sp::remove_cva_t<__shared float &>>);
static_assert(std::is_same_v<__shared const float &, sp::remove_cva_t<__shared const float &>>);
static_assert(std::is_same_v<float *, sp::remove_cva_t<float * __shared>>);
static_assert(std::is_same_v<const float *, sp::remove_cva_t<const float * __shared>>);
static_assert(std::is_same_v<__shared float *, sp::remove_cva_t<__shared float *>>);
static_assert(std::is_same_v<__shared const float *, sp::remove_cva_t<__shared const float *>>);
 
static_assert(std::is_same_v<float, sp::remove_cva_t<float>>);
static_assert(std::is_same_v<float, sp::remove_cva_t<const float>>);
static_assert(std::is_same_v<float &, sp::remove_cva_t<float &>>);
static_assert(std::is_same_v<const float &, sp::remove_cva_t<const float &>>);
static_assert(std::is_same_v<float *, sp::remove_cva_t<float *>>);
static_assert(std::is_same_v<const float *, sp::remove_cva_t<const float *>>);

remove_cvaref_t

template<typename T >

using sp::remove_cvaref_t = typedef typename remove_cvaref<T>::type

Remove reference const, volatile, and address space from a type.

Example

static_assert(std::is_same_v<float, sp::remove_cvaref_t<__shared float>>);
static_assert(std::is_same_v<float, sp::remove_cvaref_t<__shared const float>>);
static_assert(std::is_same_v<float, sp::remove_cvaref_t<__shared float &>>);
static_assert(std::is_same_v<float, sp::remove_cvaref_t<__shared const float &>>);
static_assert(std::is_same_v<float *, sp::remove_cvaref_t<float * __shared>>);
static_assert(std::is_same_v<const float *, sp::remove_cvaref_t<const float * __shared>>);
static_assert(std::is_same_v<__shared float *, sp::remove_cvaref_t<__shared float *>>);
static_assert(std::is_same_v<__shared const float *, sp::remove_cvaref_t<__shared const float *>>);
 
static_assert(std::is_same_v<float, sp::remove_cvaref_t<float>>);
static_assert(std::is_same_v<float, sp::remove_cvaref_t<const float>>);
static_assert(std::is_same_v<float, sp::remove_cvaref_t<float &>>);
static_assert(std::is_same_v<float, sp::remove_cvaref_t<const float &>>);
static_assert(std::is_same_v<float *, sp::remove_cvaref_t<float *>>);
static_assert(std::is_same_v<const float *, sp::remove_cvaref_t<const float *>>);

Enumeration Type Documentation

AddressSpace

enum sp::AddressSpace

Enum providing the address-space numbers for GPUs.

Warning

This may become an enum class in future, so it's recommended to only access these via the sp::AddressSpace scope, and not the sp scope.

These numbers are implementation dependent. Their values and relative ordering should not be depended on.

Enumerator
FLAT	A __flat pointer.
GENERIC	Language generic address space: Pointer that is compatible with any address space on the current target. Can convert to any other address space that can be dereferenced on current target (i.e: host or device, depending on the context), but may lead to the use of generic load/store instructions. On the host, this is the only address space that can be dereferenced.
DEVICE	A __device pointer.
SHARED	A __shared pointer.
CONSTANT	A __constant pointer.
LOCAL	A __local pointer.

Variable Documentation

get_addrspace_v

template<typename T >

constexpr sp::AddressSpace sp::get_addrspace_v = get_addrspace<T>::value

constexpr

Get the address space of a type.

Example

static_assert(sp::get_addrspace_v<__device float> == sp::AddressSpace::DEVICE);
static_assert(sp::get_addrspace_v<__device const float> == sp::AddressSpace::DEVICE);
static_assert(sp::get_addrspace_v<__shared float &> == sp::AddressSpace::GENERIC);
static_assert(sp::get_addrspace_v<__shared const float &> == sp::AddressSpace::GENERIC);
 
static_assert(sp::get_addrspace_v<float * __device> == sp::AddressSpace::DEVICE);
static_assert(sp::get_addrspace_v<const float * __device> == sp::AddressSpace::DEVICE);
static_assert(sp::get_addrspace_v<__device float *> == sp::AddressSpace::GENERIC);
static_assert(sp::get_addrspace_v<__device const float *> == sp::AddressSpace::GENERIC);
 
static_assert(sp::get_addrspace_v<float> == sp::AddressSpace::GENERIC);
static_assert(sp::get_addrspace_v<const float> == sp::AddressSpace::GENERIC);
static_assert(sp::get_addrspace_v<float &> == sp::AddressSpace::GENERIC);
static_assert(sp::get_addrspace_v<const float &> == sp::AddressSpace::GENERIC);