OpenCL implements the following disjoint address spaces: __global
,
__local
, __constant
, and
__private
The address space qualifier may be used in variable
declarations to specify the region of memory that is used to allocate the object. The
C syntax for type qualifiers is extended in OpenCL to include an address space name as
a valid type qualifier. If the type of an object is qualified by an address space name,
the object is allocated in the specified address name; otherwise, the object is allocated
in the generic address space.
The address space names without the __prefix i.e. global
,
local
, constant
and private
may be substituted for the corresponding address space names with the __prefix.
The generic address space name for arguments to a function in a program, or local variables
of a function is __private
. All function arguments shall be in the
__private
address space.
__kernel
function arguments declared to be a pointer of a type can point to one of the following
address spaces only: __global
, __local
or
__constant
. A pointer to address space A can only be assigned to
a pointer to the same address space A. Casting a pointer to address space A to a pointer
to address space B is illegal.
Function arguments of type
image2d_t,
image3d_t,
image2d_array_t,
image1d_t,
image1d_buffer_t, and
image1d_array_t
refer to image memory objects allocated in the __global
address space.
The __local
or local
address space name is
used to describe variables that need to be allocated in local memory and are shared by
all work-items of a work-group. Pointers to the __local
address
space are allowed as arguments to functions (including __kernel
functions). Variables allocated in the __local
address space can
also be defined inside a __kernel
function scope.
There is no generic address space name for program scope variables. All program scope
variables must be declared in the __constant
address space. For example:
// declares a pointer p in the __private address space that // points to an int object in address space __global __global int *p; // declares an array of 4 floats in the __private address space. float x[4]; |
There is no address space for function return values. Using an address space qualifier in a function return type declaration will generate a compilation error, unless the return type is declared as a pointer type and the qualifier is used on the points-to address space.
For example:
__private int f() { ... } // should generate an error __local int *f() { ... } // allowed __local int * __private f() { ... }; // should generate an error. |
// Examples of variables allocated in the __local address space // inside a __kernel function __kernel void my_func(...) { __local float a; // A single float allocated // in local address space __local float b[10]; // An array of 10 floats // allocated in local address space. if (...) { // example of variable in __local address space but not // declared at __kernel function scope. __local float c; // not allowed. } } |
Variables allocated in the __local
address space inside a
__kernel
function cannot be initialized.
__kernel void my_func(...)
{
local float a = 1; // not allowed __local float b; b = 1; // allowed
}
NOTE: Variables allocated in the __local
address
space inside a __kernel
function are allocated for each work-group executing
the kernel.