blob: eea33dcd8a04d7ff69715fad9de6426b4f696d0e [file] [log] [blame]
<?xml version="1.0" encoding="UTF-8"?>
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "xhtml1-transitional.dtd">
<!-- saved from url=(0013)about:internet -->
<?xml-stylesheet type="text/xsl" href="mathml.xsl"?><html xmlns="http://www.w3.org/1999/xhtml" xmlns:pref="http://www.w3.org/2002/Math/preference" xmlns:xlink="http://www.w3.org/1999/xlink" pref:renderer="mathplayer-dl">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=UTF-8" />
<style xmlns="" type="text/css">
/* This style sets a margin around the entire page */
html, body {
margin: 10px;
}
p {
font: normal 16px verdana, sans-serif;
margin: 0;
padding-bottom:12px;
}
h1 {
font: bold 25px verdana, sans-serif;
margin-top: 0;
margin-bottom: 3px;
padding-top: 0;
padding-bottom: 0;
}
h2 {
font: bold 19px verdana, sans-serif;
margin-top: 28px;
margin-bottom: 3px;
padding-top: 0;
padding-bottom: 0;
}
h3 {
font: bold 19px verdana, sans-serif !important;
margin-top: 28px;
margin-bottom: 3px;
padding-top: 0;
padding-bottom: 0;
}
li {
font: normal 16px verdana, sans-serif;
margin-top: 0;
margin-bottom: 18px;
padding-top: 0;
padding-bottom: 0;
}
.pdparam {
font: italic 16px verdana, sans-serif;
}
.term {
font: italic 16px verdana, sans-serif;
font-weight: normal;
}
.type {
font: normal 16px verdana, sans-serif !important;
}
.parameter {
font-style: italic;
}
a:link, a:visited {
color: blue;
text-decoration: none;
font: normal 16px;
}
a:hover {
background-color: #FFFF99;
font: normal 16px;
}
div.funcsynopsis {
text-align: left;
background-color: #e6e6e6;
font: normal 16px verdana, sans-serif;
padding-top: 10px;
padding-bottom: 10px;
}
div.funcsynopsis table {
border-collapse: separate;
font: normal 16px verdana, sans-serif;
}
div.funcsynopsis td {
background-color: #e6e6e6;
border: 0 solid #000;
padding: 1px;
font: normal 16px verdana, sans-serif;
}
div.refsect1 {
font-family: verdana, sans-serif;
font-size: 16px;
}
code.constant {
font: normal 16px courier new, monospace !important;
}
span.errorname {
font: normal 16px verdana, sans-serif !important;
}
code.function {
font: bold 16px verdana, sans-serif !important;
}
b.fsfunc {
font: bold 16px verdana, sans-serif !important;
}
code.varname {
font: italic 16px verdana, sans-serif;
}
code.replaceable {
font: italic 16px courier new, monospace;
}
code.funcdef {
font: normal 16px verdana, sans-serif !important;
}
.citerefentry {
font: normal 16px verdana, sans-serif !important;
}
.parameter {
font-style: italic;
}
code.fsfunc {
font: normal 16px verdana, sans-serif !important;
}
/* PARAMETER: This style controls spacing between the terms in Parameter section */
dt {
margin-top: 15px;
}
/* TABLES: These styles apply to all tables OTHER than the Synopsis and Example tables */
div.refsect1 table {
width: 100%;
margin-top: 10px;
background-color: #FFF;
border-collapse: collapse;
border-color: #000;
border-width: 1px;
font: normal 16px verdana, sans-serif;
}
div.refsect1 th {
border-collapse: collapse;
border-color: #000;
border-width: 1px;
font: bold 16px verdana, sans-serif;
}
div.refsect1 td {
background-color: #FFF;
padding: 5px;
vertical-align: text-top;
border-collapse: collapse;
border-color: #000;
border-width: 1px;
font: normal 16px verdana, sans-serif;
}
div.refsect1 p{
font: normal 16px verdana, sans-serif;
margin-top: 8px;
margin-bottom: 8px;
padding-top: 0;
padding-bottom: 0;
}
/* EXAMPLE: These styles apply only to the Example section */
div.refsect2 {
font: normal 16px courier new, monospace !important;
}
div.refsect2 table {
margin-top: 0;
background-color: #e6e6e6;
width: 100%;
border: 0 solid #000;
padding: 2px;
font: normal 16px courier new, monospace !important;
}
div.refsect2 td {
background-color: #e6e6e6;
font: normal 16px courier new, monospace !important;
white-space:pre;
}
/* COPYRIGHT: This style formats the text of the copyright statement at the bottom of the page */
div.refsect3 {
font: normal 11px verdana, sans-serif;
margin-top: 50px;
margin-bottom: 20px;
padding-top: 0;
padding-bottom: 0;
}
</style>
<title>__local</title>
<meta name="generator" content="DocBook XSL Stylesheets V1.78.1" />
<meta name="keywords" content="__local" />
</head>
<body>
<div class="refentry">
<a id="id-1"></a>
<div class="titlepage"></div>
<div xmlns="" class="refnamediv">
<a xmlns="http://www.w3.org/1999/xhtml" id="__local"></a>
<h1>__local</h1>
<p>
Address Space Qualifier.
</p>
</div>
<div class="refsect2">
<a id="synopsis"></a>
<h3>
</h3>
<div class="informaltable">
<table border="0">
<colgroup>
<col align="left" class="col1" />
</colgroup>
<tbody>
<tr>
<td align="left">
__local
local
</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="refsect1">
<a id="description"></a>
<h2>Description</h2>
<p>
OpenCL implements the following disjoint address spaces: <a class="citerefentry" href="global.html"><span class="citerefentry"><span class="refentrytitle">__global</span></span></a>,
<code class="function">__local</code>,
<a class="citerefentry" href="constant.html"><span class="citerefentry"><span class="refentrytitle">__constant</span></span></a> and <a class="citerefentry" href="private.html"><span class="citerefentry"><span class="refentrytitle">__private</span></span></a>.
</p>
<p>
</p>
<p>
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.
</p>
<p>
The address space names without the __ prefix i.e. <code class="function">global</code>, <code class="function">local</code>, <code class="function">constant</code> and <code class="function">private</code> may be substituted for the corresponding address space names with the __ prefix.
</p>
<p>
The generic address space name for arguments to a function in a program, or local variables of a function is <code class="function">__private</code>. All function arguments shall be in the <code class="function">__private</code> address space.
</p>
<p>
<a class="citerefentry" href="functionQualifiers.html"><span class="citerefentry"><span class="refentrytitle">__kernel</span></span></a> function arguments declared to be a pointer of a type can point to one of the following address spaces only: <code class="function">__global</code>, <code class="function">__local</code> or <code class="function">__constant</code>. 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.
</p>
<p>
The <code class="function">__kernel</code> function arguments declared to be of type <span class="type"><a class="citerefentry" href="otherDataTypes.html"><span class="citerefentry"><span class="refentrytitle">image2d_t</span></span></a></span> or <span class="type"><a class="citerefentry" href="otherDataTypes.html"><span class="citerefentry"><span class="refentrytitle">image3d_t</span></span></a></span>
always point to the <code class="function">__global</code> address space.
</p>
<p>
There is no generic address space name for program scope variables. All program scope variables must be declared in the <code class="function">__constant</code> address space.
</p>
<p>
The <code class="function">__constant</code> or <code class="function">constant</code>
address space name is used to describe variables allocated in
global memory and which are accessed inside a kernel(s) as read-only variables. These read-
only variables can be accessed by all (global) work-items of the kernel during its execution. Pointers to the <code class="function">__constant</code> address space are allowed arguments to functions (including <code class="function">__kernel</code> functions) and for variables declared inside functions. Variables allocated in the <code class="function">__constant</code> address space can only be defined as program scope variables and are required to be initialized.
</p>
<p>
Writes to variables declared with the <code class="function">__constant</code> address space qualifier in the OpenCL program source should be a compile-time error.
</p>
<p>
All string literal storage shall be in the <code class="function">__constant</code> address space.
</p>
<p>
NOTE: Any argument to a kernel that is declared with the <code class="function">__constant</code> address space qualifier counts as a separate constant argument. The relevant device capabilities CL_DEVICE_MAX_CONSTANT_ARGS and CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE are described in the table for <a class="citerefentry" href="clGetDeviceInfo.html"><span class="citerefentry"><span class="refentrytitle">clGetDeviceInfo</span></span></a>.
</p>
<p>
Variables inside a function or in program scope can also be declared with the <code class="function">__constant</code> address qualifier. Implementations are not required to aggregate these declarations into the fewest number of constant arguments. This behavior is implementation defined.
</p>
<p>
Thus portable code must conservatively assume that each variable declared inside a function or in program scope with the <code class="function">__constant</code> qualifier counts as a separate constant argument.
</p>
</div>
<div class="refsect1">
<a id="description"></a>
<h2>Description</h2>
<p>
The <code class="function">__local</code> or <code class="function">local</code> 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 <code class="function">__local</code> address
space are allowed as arguments to functions (including <code class="function">__kernel</code>
functions). Variables allocated in the <code class="function">__local</code> address space can
also be defined inside a <code class="function">__kernel</code> function only.
</p>
</div>
<div class="refsect2">
<a id="example1"></a>
<h3>
Example
</h3>
<div class="informaltable">
<table border="0">
<colgroup>
<col align="left" class="col1" />
</colgroup>
<tbody>
<tr>
<td align="left">
// 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];
// 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.
}
</td>
</tr>
</tbody>
</table>
</div>
<p>
Variables allocated in the <code class="function">__local</code> address space inside a <code class="function">__kernel</code> function cannot be initialized.
</p>
<div class="literallayout">
<p><br />
     __kernel void my_func(...)<br />
    {<br />
       local float a = 1;    // not allowed<br />
       __local float b;<br />
       b = 1;                // allowed<br />
    }<br />
</p>
</div>
<p>NOTE: Variables allocated in the <code class="function">__local</code> address space inside a <code class="function">__kernel</code> function are allocated for each work-group executing the kernel.</p>
</div>
<div class="refsect1">
<a id="specification"></a>
<h2>Specification</h2>
<p>
<img src="pdficon_small1.gif" />
<a href="http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf#page=182" target="OpenCL Spec">OpenCL Specification</a>
</p>
</div>
<div class="refsect1">
<a id="seealso"></a>
<h2>Also see</h2>
<p>
<a class="citerefentry" href="global.html"><span class="citerefentry"><span class="refentrytitle">__global</span></span></a>,
<a class="citerefentry" href="constant.html"><span class="citerefentry"><span class="refentrytitle">__constant</span></span></a>,
<a class="citerefentry" href="private.html"><span class="citerefentry"><span class="refentrytitle">__private</span></span></a>
</p>
</div>
<div xmlns="" class="refsect3" lang="en" xml:lang="en"><a xmlns="http://www.w3.org/1999/xhtml" id="Copyright"></a><h4 xmlns="http://www.w3.org/1999/xhtml"></h4><img xmlns="http://www.w3.org/1999/xhtml" src="KhronosLogo.jpg" /><p xmlns="http://www.w3.org/1999/xhtml"></p>Copyright © 2007-2010 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.
</div>
</div>
</body>
</html>