blob: 693499430bd4b86e30ad8f08d1f53253cab3615f [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>genericAddressSpace</title>
<meta name="generator" content="DocBook XSL Stylesheets V1.79.1" />
<meta name="keywords" content="genericAddressSpace" />
</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="genericAddressSpace"></a>
<h1>generic</h1>
<p>
Generic Address Space.
</p>
</div>
<div class="refsect2">
<a id="synopsis"></a>
<h3>
</h3>
<div class="informaltable">
<table class="informaltable" border="0">
<colgroup>
<col align="left" class="col1" />
</colgroup>
<tbody>
<tr>
<td align="left">
__generic (unnamed address space)
</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="refsect1">
<a id="description"></a>
<h2>Description</h2>
<p>
Generic address space (glossary): An address space that include the private, local, and global address
spaces available to a device. The generic address space supports conversion of pointers to and
from private, local and global address spaces, and hence lets a programmer write a single
function that at compile time can take arguments from any of the three named address spaces.
</p>
</div>
<div class="refsect1">
<a id="notes"></a>
<h2>Notes</h2>
<p>
Programmers often need functions callable from kernels where the pointers manipulated by those
functions can point to multiple named address spaces. This saves a programmer from the error-
prone and wasteful practice of creating multiple copies of functions; one for each named address
space. Therefore the global, local and private address spaces belong to a single
<span class="emphasis"><em>generic address space</em></span>.
This is closely modeled after the concept of a generic address space used in the embedded
C standard (ISO/IEC 9899:1999). Since they all belong to a single generic address space, the
following properties are supported for pointers to named address spaces in device memory:
</p>
<p>
The following rules apply when using pointers that point to the generic address space:
</p>
<div class="itemizedlist">
<ul class="itemizedlist" style="list-style-type: disc; ">
<li class="listitem">
<p>
A pointer that points to the <code class="function">global</code>,
<code class="function">local</code> or <code class="function">private</code> address space can be
implicitly converted to a pointer to the unnamed
generic address space but not vice-versa.
</p>
</li>
<li class="listitem">
<p>
Pointer casts can be used to cast a pointer that points to the
<code class="function">global</code>, <code class="function">local</code>
or <code class="function">private</code> space to
the unnamed generic addresss space and vice-versa.
</p>
</li>
<li class="listitem">
<p>
A pointer that points to the <code class="function">constant</code>
address space cannot be cast or implicitly
converted to the generic address space.
</p>
</li>
</ul>
</div>
</div>
<div class="refsect2">
<a id="example1"></a>
<h3>
Example
</h3>
<p>
This is the canonical example. In this example, function
<code class="function">foo</code>() is declared with an argument that
is a pointer with no address space qualifier.
</p>
<div class="informaltable">
<table class="informaltable" border="0">
<colgroup>
<col align="left" class="col1" />
</colgroup>
<tbody>
<tr>
<td align="left">
void foo(int *a)
{
*a = *a + 2;
}
kernel void k1(local int *a)
{
foo(a);
}
kernel void k2(global int *a)
{
foo(a);
}
</td>
</tr>
</tbody>
</table>
</div>
<p>
In the example below, <code class="varname">var</code> is in the unnamed
generic address space which gets mapped to the
<code class="function">global</code> or <code class="function">local</code>
address space depending on the result of the conditional expression.
</p>
<div class="informaltable">
<table class="informaltable" border="0">
<colgroup>
<col align="left" class="col1" />
</colgroup>
<tbody>
<tr>
<td align="left">
kernel void bar(global int *g, local int *l)
{
int *var;
if (is_even(get_global_id(0))
var = g;
else
var = l;
*var = 42;
}
</td>
</tr>
</tbody>
</table>
</div>
<p>
The example below is an example with one unnamed
generic address space pointer with multiple
named address space assignments.
</p>
<div class="informaltable">
<table class="informaltable" border="0">
<colgroup>
<col align="left" class="col1" />
</colgroup>
<tbody>
<tr>
<td align="left">
int *ptr;
global int g;
ptr = &amp;g; // legal
local int l;
ptr = &amp;l; // legal
private int p;
ptr = &amp;p; // legal
constant int c;
ptr = &amp;c; // illegal
</td>
</tr>
</tbody>
</table>
</div>
<p>
The example below is an example with one unnamed
generic address space pointer being
assigned to point to several named address spaces.
</p>
<div class="informaltable">
<table class="informaltable" border="0">
<colgroup>
<col align="left" class="col1" />
</colgroup>
<tbody>
<tr>
<td align="left">
global int * gp;
local int *lp;
private int *pp;
int *p;
p = gp; // legal
p = lp; // legal
p = pp; // legal
// it is illegal to convert from a generic pointer
// to an explicit address space pointer without a cast:
gp = p; // compile-time error
lp = p; // compile-time error
pp = p; // compile-time error
</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="refsect1">
<a id="specification"></a>
<h2>Specification</h2>
<p>
<img src="pdficon_small1.gif" />
<a href="https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#namedest=genericAddressSpace" 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>,
<a class="citerefentry" href="qualifiers.html"><span class="citerefentry"><span class="refentrytitle">qualifiers</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-2015 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>