Address space qualifier.
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.
functionQualifiers(3clc) 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 otherDataTypes(3clc), otherDataTypes(3clc), otherDataTypes(3clc), otherDataTypes(3clc), otherDataTypes(3clc), and otherDataTypes(3clc) 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.
\m[blue]OpenCL Specification\m[]\s-2\u[1]\d\s+2
global(3clc), constant(3clc), private(3clc)
The Khronos Group
Copyright © 2007-2011y The Khronos Group Inc.
Permission is hereby granted, free of charge, to any person obtaining a copy of this software and/or associated documentation files (the "Materials"), to deal in the Materials without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Materials, and to permit persons to whom the Materials are furnished to do so, subject to the condition that this copyright notice and permission notice shall be included in all copies or substantial portions of the Materials.
OpenCL Specification
page 222, section 6.5 - Address Space Qualifiers