<!DOCTYPE html
PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" lang="en-us" xml:lang="en-us">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8"></meta>
<meta http-equiv="X-UA-Compatible" content="IE=edge"></meta>
<meta name="copyright" content="(C) Copyright 2005"></meta>
<meta name="DC.rights.owner" content="(C) Copyright 2005"></meta>
<meta name="DC.Type" content="concept"></meta>
<meta name="DC.Title" content="NVVM IR Specification 1.1"></meta>
<meta name="abstract" content="Reference guide to the NVVM compiler IR (internal representation) based on the LVVM IR."></meta>
<meta name="description" content="Reference guide to the NVVM compiler IR (internal representation) based on the LVVM IR."></meta>
<meta name="DC.Coverage" content="Compiler SDK"></meta>
<meta name="DC.subject" content="NVVM, NVVM IR, NVVM compiler, NVVM compiler IR, NVVM IR specification, NVVM specification, NVVM LLVM IR, NVVM CUDA, NVVM CUDA C compiler, NVVM GPU"></meta>
<meta name="keywords" content="NVVM, NVVM IR, NVVM compiler, NVVM compiler IR, NVVM IR specification, NVVM specification, NVVM LLVM IR, NVVM CUDA, NVVM CUDA C compiler, NVVM GPU"></meta>
<meta name="DC.Format" content="XHTML"></meta>
<meta name="DC.Identifier" content="abstract"></meta>
<link rel="stylesheet" type="text/css" href="../common/formatting/commonltr.css"></link>
<link rel="stylesheet" type="text/css" href="../common/formatting/site.css"></link>
<title>NVVM IR :: CUDA Toolkit Documentation</title>
<!--[if lt IE 9]>
<script src="../common/formatting/html5shiv-printshiv.min.js"></script>
<![endif]-->
<script type="text/javascript" charset="utf-8" src="../common/scripts/tynt/tynt.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.min.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.ba-hashchange.min.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.scrollintoview.min.js"></script>
<script type="text/javascript" src="../search/htmlFileList.js"></script>
<script type="text/javascript" src="../search/htmlFileInfoList.js"></script>
<script type="text/javascript" src="../search/nwSearchFnt.min.js"></script>
<script type="text/javascript" src="../search/stemmers/en_stemmer.min.js"></script>
<script type="text/javascript" src="../search/index-1.js"></script>
<script type="text/javascript" src="../search/index-2.js"></script>
<script type="text/javascript" src="../search/index-3.js"></script>
<link rel="canonical" href="http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html"></link>
<link rel="stylesheet" type="text/css" href="../common/formatting/qwcode.highlight.css"></link>
</head>
<body>
<header id="header"><span id="company">NVIDIA</span><span id="site-title">CUDA Toolkit Documentation</span><form id="search" method="get" action="search">
<input type="text" name="search-text"></input><fieldset id="search-location">
<legend>Search In:</legend>
<label><input type="radio" name="search-type" value="site"></input>Entire Site</label>
<label><input type="radio" name="search-type" value="document"></input>Just This Document</label></fieldset>
<button type="reset">clear search</button>
<button id="submit" type="submit">search</button></form>
</header>
<div id="site-content">
<nav id="site-nav">
<div class="category closed"><a href="../index.html" title="The root of the site.">CUDA Toolkit
v6.5</a></div>
<div class="category"><a href="index.html" title="NVVM IR">NVVM IR</a></div>
<ul>
<li>
<div class="section-link"><a href="#introduction">1. Introduction</a></div>
</li>
<li>
<div class="section-link"><a href="#identifiers">2. Identifiers</a></div>
</li>
<li>
<div class="section-link"><a href="#high-level-structure">3. High Level Structure</a></div>
<ul>
<li>
<div class="section-link"><a href="#linkage-types">3.1. Linkage Types</a></div>
</li>
<li>
<div class="section-link"><a href="#calling-conventions">3.2. Calling Conventions</a></div>
<ul>
<li>
<div class="section-link"><a href="#rules-and-restrictions">3.2.1. Rules and Restrictions</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#visibility-styles">3.3. Visibility Styles</a></div>
</li>
<li>
<div class="section-link"><a href="#named-types">3.4. Named Types</a></div>
</li>
<li>
<div class="section-link"><a href="#global-variables">3.5. Global Variables</a></div>
</li>
<li>
<div class="section-link"><a href="#functions">3.6. Functions</a></div>
</li>
<li>
<div class="section-link"><a href="#aliases">3.7. Aliases</a></div>
</li>
<li>
<div class="section-link"><a href="#named-metadata">3.8. Named Metadata</a></div>
</li>
<li>
<div class="section-link"><a href="#parameter-attributes">3.9. Parameter Attributes</a></div>
</li>
<li>
<div class="section-link"><a href="#garbage-collector-names">3.10. Garbage Collector Names</a></div>
</li>
<li>
<div class="section-link"><a href="#function-attributes">3.11. Function Attributes</a></div>
</li>
<li>
<div class="section-link"><a href="#module-level-inline-assembly">3.12. Module-Level Inline Assembly</a></div>
</li>
<li>
<div class="section-link"><a href="#data-layout">3.13. Data Layout</a></div>
</li>
<li>
<div class="section-link"><a href="#pointer-aliasing-rules">3.14. Pointer Aliasing Rules</a></div>
</li>
<li>
<div class="section-link"><a href="#volatile-memory-access">3.15. Volatile Memory Access</a></div>
</li>
<li>
<div class="section-link"><a href="#atomic-memory-ordering-constraints">3.16. Atomic Memory Ordering Constraints</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#type-system">4. Type System</a></div>
</li>
<li>
<div class="section-link"><a href="#constants">5. Constants</a></div>
</li>
<li>
<div class="section-link"><a href="#other-values">6. Other Values</a></div>
<ul>
<li>
<div class="section-link"><a href="#inline-assembler-expressions">6.1. Inline Assembler Expressions</a></div>
</li>
<li>
<div class="section-link"><a href="#metadata-nodes-and-metadata-strings">6.2. Metadata Nodes and Metadata Strings</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#intrinsic-global-variables">7. Intrinsic Global Variables</a></div>
</li>
<li>
<div class="section-link"><a href="#instructions">8. Instructions</a></div>
<ul>
<li>
<div class="section-link"><a href="#terminator-instructions">8.1. Terminator Instructions</a></div>
</li>
<li>
<div class="section-link"><a href="#binary-operations">8.2. Binary Operations</a></div>
</li>
<li>
<div class="section-link"><a href="#bitwise-binary-operations">8.3. Bitwise Binary Operations</a></div>
</li>
<li>
<div class="section-link"><a href="#vector-operations">8.4. Vector Operations</a></div>
</li>
<li>
<div class="section-link"><a href="#aggregate-operations">8.5. Aggregate Operations</a></div>
</li>
<li>
<div class="section-link"><a href="#memory-access-and-addressing-operations">8.6. Memory Access and Addressing Operations</a></div>
<ul>
<li>
<div class="section-link"><a href="#alloca-instruction">8.6.1. alloca Instruction</a></div>
</li>
<li>
<div class="section-link"><a href="#load-instruction">8.6.2. load Instruction</a></div>
</li>
<li>
<div class="section-link"><a href="#store-instruction">8.6.3. store Instruction</a></div>
</li>
<li>
<div class="section-link"><a href="#fence-instruction">8.6.4. fence Instruction</a></div>
</li>
<li>
<div class="section-link"><a href="#cmpxchg-instruction">8.6.5. cmpxchg Instruction</a></div>
</li>
<li>
<div class="section-link"><a href="#atomicrmw-instruction">8.6.6. atomicrmw Instruction</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#conversion-operations">8.7. Conversion Operations</a></div>
</li>
<li>
<div class="section-link"><a href="#other-operations">8.8. Other Operations</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#intrinsic-functions">9. Intrinsic Functions</a></div>
<ul>
<li>
<div class="section-link"><a href="#variable-argument-handling-intrinsics">9.1. Variable Argument Handling Intrinsics</a></div>
</li>
<li>
<div class="section-link"><a href="#accurate-garbage-collection-intrinics">9.2. Accurate Garbage Collection Intrinsics</a></div>
</li>
<li>
<div class="section-link"><a href="#code-generator-intrinics">9.3. Code Generator Intrinsics</a></div>
</li>
<li>
<div class="section-link"><a href="#standard-c-library-intrinics">9.4. Standard C Library Intrinsics</a></div>
</li>
<li>
<div class="section-link"><a href="#bit-manipulations-intrinics">9.5. Bit Manipulations Intrinsics</a></div>
</li>
<li>
<div class="section-link"><a href="#arithmetic-with-overflow-intrinsics">9.6. Arithmetic with Overflow Intrinsics</a></div>
</li>
<li>
<div class="section-link"><a href="#half-precision-floating-point-intrinsics">9.7. Half Precision Floating Point Intrinsics</a></div>
</li>
<li>
<div class="section-link"><a href="#debugger-intrinsics">9.8. Debugger Intrinsics</a></div>
</li>
<li>
<div class="section-link"><a href="#exception-handling-intrinsics">9.9. Exception Handling Intrinsics</a></div>
</li>
<li>
<div class="section-link"><a href="#trampoline-intrinsics">9.10. Trampoline Intrinsics</a></div>
</li>
<li>
<div class="section-link"><a href="#memory-use-markers">9.11. Memory Use Markers</a></div>
</li>
<li>
<div class="section-link"><a href="#general-intrinsics">9.12. General Intrinsics</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#address-space">10. Address Space</a></div>
<ul>
<li>
<div class="section-link"><a href="#address-spaces">10.1. Address Spaces</a></div>
</li>
<li>
<div class="section-link"><a href="#generic-pointers-and-non-generic-pointers">10.2. Generic Pointers and Non-Generic Pointers</a></div>
<ul>
<li>
<div class="section-link"><a href="#generic-pointers-vs-non-generic-pointers">10.2.1. Generic Pointers vs. Non-generic Pointers</a></div>
</li>
<li>
<div class="section-link"><a href="#conversion">10.2.2. Conversion</a></div>
</li>
<li>
<div class="section-link"><a href="#no-aliasing-between-two-different-specific-address-spaces">10.2.3. No Aliasing between Two Different Specific Address Spaces</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#the-alloca-instruction">10.3. The alloca Instruction</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#global-property-annotation-chapter-11">11. Global Property Annotation</a></div>
<ul>
<li>
<div class="section-link"><a href="#overview">11.1. Overview</a></div>
</li>
<li>
<div class="section-link"><a href="#representation-of-properties">11.2. Representation of Properties</a></div>
</li>
<li>
<div class="section-link"><a href="#supported-properties">11.3. Supported Properties</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#texture-and-surface">12. Texture and Surface</a></div>
<ul>
<li>
<div class="section-link"><a href="#texture-variable-and-surface-variable">12.1. Texture Variable and Surface Variable</a></div>
</li>
<li>
<div class="section-link"><a href="#accessing-texture-memory-or-surface-memory">12.2. Accessing Texture Memory or Surface Memory</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#nvvm-specific-intrinsic-functions">13. NVVM Specific Intrinsic Functions</a></div>
<ul>
<li>
<div class="section-link"><a href="#nvvm-intrin-atomic">13.1. Atomic</a></div>
</li>
<li>
<div class="section-link"><a href="#nvvm-intrin-barrier">13.2. Barrier and Memory Fence</a></div>
</li>
<li>
<div class="section-link"><a href="#nvvm-intrin-addrsp">13.3. Address space conversion</a></div>
</li>
<li>
<div class="section-link"><a href="#nvvm-intrin-spreg">13.4. Special Registers</a></div>
</li>
<li>
<div class="section-link"><a href="#nvvm-intrin-texture">13.5. Texture/surface Access</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#nvvm-abi-for-ptx">14. NVVM ABI for PTX</a></div>
<ul>
<li>
<div class="section-link"><a href="#linkage-types-nvvm">14.1. Linkage Types</a></div>
</li>
<li>
<div class="section-link"><a href="#argument-passing-and-return">14.2. Argument for Passing and Return</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#revision-history">A. Revision History</a></div>
</li>
</ul>
</nav>
<div id="resize-nav"></div>
<nav id="search-results">
<h2>Search Results</h2>
<ol></ol>
</nav>
<div id="contents-container">
<div id="breadcrumbs-container">
<div id="release-info">NVVM IR
(<a href="../../pdf/NVVM_IR_Specification.pdf">PDF</a>)
-
v6.5
(<a href="https://developer.nvidia.com/cuda-toolkit-archive">older</a>)
-
Last updated August 1, 2014
-
<a href="mailto:cudatools@nvidia.com?subject=CUDA Toolkit Documentation Feedback: NVVM IR">Send Feedback</a>
-
<span class="st_facebook"></span><span class="st_twitter"></span><span class="st_linkedin"></span><span class="st_reddit"></span><span class="st_slashdot"></span><span class="st_tumblr"></span><span class="st_sharethis"></span></div>
</div>
<article id="contents">
<div class="topic nested0" id="abstract"><a name="abstract" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#abstract" name="abstract" shape="rect">NVVM IR Specification 1.1</a></h2>
<div class="body conbody"></div>
</div>
<div class="topic concept nested0" id="introduction"><a name="introduction" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#introduction" name="introduction" shape="rect">1. Introduction</a></h2>
<div class="body conbody">
<p class="p">NVVM IR is a compiler IR (internal representation) based on the LLVM IR. The NVVM IR is
designed to represent GPU compute kernels (for example, CUDA kernels). High-level
language front-ends, like the CUDA C compiler front-end, can generate NVVM IR. The NVVM
compiler (which is based on LLVM) generates PTX code from NVVM IR.
</p>
<p class="p">NVVM IR and NVVM compilers are mostly agnostic about the source language being used. The
PTX codegen part of a NVVM compiler needs to know the source language because of the
difference in DCI (driver/compiler interface).
</p>
<p class="p">Technically speaking, NVVM IR is LLVM IR with a set of rules, restrictions, and
conventions, plus a set of supported intrinsic functions. A program specified in NVVM IR
is always a legal LLVM program. A legal LLVM program may not be a legal NVVM program.
</p>
<p class="p">There are three levels of support for LLVM IR.</p>
<ul class="ul">
<li class="li">Supported: The feature is fully supported. Most IR features should fall into this
category.
</li>
<li class="li">Accepted and ignored: The NVVM compiler will accept this IR feature, but will ignore
the required semantics. This applies to some IR features that do not have meaningful
semantics on GPUs and that can be ignored. Calling convention markings are an
example.
</li>
<li class="li">Illegal, not supported: The specified semantics is not supported, such as a
<samp class="ph codeph">va_arg</samp> function. Future versions of NVVM may either support or
accept and ignore IRs that are illegal in the current version.
</li>
</ul>
<p class="p">The current NVVM IR is based on LLVM 3.2. For the complete semantics of the IR, readers
of this document should check the official LLVM Language Reference Manual (<a class="xref" href="http://llvm.org/releases/3.2/docs/LangRef.html" target="_blank" shape="rect">http://llvm.org/releases/3.2/docs/LangRef.html</a>).
</p>
</div>
</div>
<div class="topic concept nested0" id="identifiers"><a name="identifiers" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#identifiers" name="identifiers" shape="rect">2. Identifiers</a></h2>
<div class="body conbody">
<p class="p">The name of a named global identifier must have the form:</p>
<p class="p"><samp class="ph codeph">@[a-zA-Z$_][a-zA-Z$_0-9]*</samp></p>
<p class="p">Note that it cannot contain the . character.</p>
<p class="p"><samp class="ph codeph">[@%]llvm.nvvm.*</samp> and <samp class="ph codeph">[@%]nvvm.*</samp> are reserved words.
</p>
</div>
</div>
<div class="topic concept nested0" id="high-level-structure"><a name="high-level-structure" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#high-level-structure" name="high-level-structure" shape="rect">3. High Level Structure</a></h2>
<div class="topic concept nested1" id="linkage-types"><a name="linkage-types" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#linkage-types" name="linkage-types" shape="rect">3.1. Linkage Types</a></h3>
<div class="body conbody">
<p class="p">All linkage types (except for <samp class="ph codeph">dllimport</samp> and <samp class="ph codeph">dllexport</samp>)
are supported. <a class="xref" href="index.html#nvvm-abi-for-ptx" shape="rect">NVVM ABI for PTX</a> provides details on how they are
translated to PTX.
</p>
</div>
</div>
<div class="topic concept nested1" id="calling-conventions"><a name="calling-conventions" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#calling-conventions" name="calling-conventions" shape="rect">3.2. Calling Conventions</a></h3>
<div class="body conbody">
<p class="p">All LLVM calling convention markings are accepted and ignored. Functions and calls are
generated according to the PTX calling convention.
</p>
</div>
<div class="topic concept nested2" id="rules-and-restrictions"><a name="rules-and-restrictions" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#rules-and-restrictions" name="rules-and-restrictions" shape="rect">3.2.1. Rules and Restrictions</a></h3>
<div class="body conbody"><a name="rules-and-restrictions__ol_x5g_nr1_n3" shape="rect">
<!-- --></a><ol class="ol" id="rules-and-restrictions__ol_x5g_nr1_n3">
<li class="li"><samp class="ph codeph">va_arg</samp> is not supported.
</li>
<li class="li">When an argument with width less than 32-bit is passed, the <samp class="ph codeph">zeroext/signext</samp>
parameter attribute should be set. <samp class="ph codeph">zeroext</samp> will be assumed if not set.
</li>
<li class="li">When a value with width less than 32-bit is returned, the
<samp class="ph codeph">zeroext/signext</samp> parameter attribute should be set.
<samp class="ph codeph">zeroext</samp> will be assumed if not set.
</li>
<li class="li">Arguments of aggregate or vector types that are passed by value can be passed by
pointer with the <samp class="ph codeph">byval</samp> attribute set (referred to as the
<samp class="ph codeph">by-pointer-byval</samp> case below). The align attribute must be
set if the type requires a non-natural alignment (natural alignment is the alignment
inferred for the aggregate type according to the <a class="xref" href="index.html#data-layout" shape="rect">Data Layout</a>
section).
</li>
<li class="li">If a function has an argument of aggregate or vector type that is passed by value
directly and the type has a non-natural alignment requirement, the alignment must be
annotated by the global property annotation <<samp class="ph codeph">align</samp>,
alignment>, where alignment is a 32-bit integer whose upper 16 bits
represent the argument position (starting from 1) and the lower 16 bits represent
the alignment.
</li>
<li class="li">If the return type of a function is an aggregate or a vector that has a non-natural
alignment, then the alignment requirement must be annotated by the global property
annotation <<samp class="ph codeph">align</samp>, alignment>, where the upper 16 bits is
0, and the lower 16 bits represent the alignment.
</li>
<li class="li">It is not required to annotate a function with <<samp class="ph codeph">align</samp>,
alignment> otherwise. If annotated, the alignment must match the
natural alignment or the align attribute in the <samp class="ph codeph">by-pointer-byval</samp>
case.
</li>
<li class="li" id="rules-and-restrictions__item-8-calling-conventions"><a name="rules-and-restrictions__item-8-calling-conventions" shape="rect">
<!-- --></a>For an indirect call instruction of a function that
has a non-natural alignment for its return value or one of its arguments that is not
expressed in alignment in the <samp class="ph codeph">by-pointer-byval</samp> case, the call
instruction must have an attached metadata of kind <samp class="ph codeph">callalign</samp>. The
metadata contains a sequence of <samp class="ph codeph">i32</samp> fields each of which represents
a non-natural alignment requirement. The upper 16 bits of an <samp class="ph codeph">i32</samp>
field represent the argument position (0 for return value, 1 for the first argument,
and so on) and the lower 16 bits represent the alignment. The <samp class="ph codeph">i32</samp>
fields must be sorted in the increasing order.
<p class="p">For
example,
</p><pre xml:space="preserve">
%call = call %struct.S %fp1(%struct.S* byval align 8 %arg1p, %struct.S %arg2),!callalign !10
!10 = metadata !{i32 0x0008, i32 0x208};
</pre><p class="p"></p>
</li>
<li class="li">It is not required to have an <samp class="ph codeph">i32</samp> metadata field for the other
arguments or the return value otherwise. If presented, the alignment must match the
natural alignment or the align attribute in the <samp class="ph codeph">by-pointer-byval
case</samp>.
</li>
<li class="li">It is not required to have a <samp class="ph codeph">callalign</samp> metadata attached to a
direct call instruction. If attached, the alignment must match the natural alignment
or the alignment in the <samp class="ph codeph">by-pointer-byval</samp> case.
</li>
<li class="li">The absence of the metadata in an indirect call instruction means using natural
alignment or the align attribute in the <samp class="ph codeph">by-pointer-byval</samp> case.
</li>
</ol>
</div>
</div>
</div>
<div class="topic concept nested1" id="visibility-styles"><a name="visibility-styles" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#visibility-styles" name="visibility-styles" shape="rect">3.3. Visibility Styles</a></h3>
<div class="body conbody">
<p class="p">All styles—default, hidden, and protected—are accepted and ignored.</p>
</div>
</div>
<div class="topic concept nested1" id="named-types"><a name="named-types" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#named-types" name="named-types" shape="rect">3.4. Named Types</a></h3>
<div class="body conbody">
<p class="p">Fully supported.</p>
</div>
</div>
<div class="topic concept nested1" id="global-variables"><a name="global-variables" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#global-variables" name="global-variables" shape="rect">3.5. Global Variables</a></h3>
<div class="body conbody">
<p class="p">A global variable, that is not an intrinsic global variable, may be optionally declared to reside
in one of the following address spaces:
</p><a name="global-variables__ul_p4q_4cb_n3" shape="rect">
<!-- --></a><ul class="ul" id="global-variables__ul_p4q_4cb_n3">
<li class="li"><samp class="ph codeph">global</samp></li>
<li class="li"><samp class="ph codeph">shared</samp></li>
<li class="li"><samp class="ph codeph">constant</samp></li>
</ul>
<p class="p">If no address space is explicitly specified, the global variable is assumed to reside in the
<samp class="ph codeph">global</samp> address space with a generic address value.
See <a class="xref" href="index.html#address-space" shape="rect">Address Space</a> for details.
</p>
<p class="p"><samp class="ph codeph">thread_local</samp> variables are not supported.
</p>
<p class="p">No explicit section (except for the metadata section) is allowed. </p>
</div>
</div>
<div class="topic concept nested1" id="functions"><a name="functions" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#functions" name="functions" shape="rect">3.6. Functions</a></h3>
<div class="body conbody">
<p class="p">The following are not supported on functions.</p><a name="functions__ul_gtp_gdb_n3" shape="rect">
<!-- --></a><ul class="ul" id="functions__ul_gtp_gdb_n3">
<li class="li"><samp class="ph codeph">explicit section</samp></li>
<li class="li"><samp class="ph codeph">alignment</samp></li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="aliases"><a name="aliases" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#aliases" name="aliases" shape="rect">3.7. Aliases</a></h3>
<div class="body conbody">
<p class="p">Fully supported.</p>
</div>
</div>
<div class="topic concept nested1" id="named-metadata"><a name="named-metadata" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#named-metadata" name="named-metadata" shape="rect">3.8. Named Metadata</a></h3>
<div class="body conbody">
<p class="p">Accepted and ignored, except for the following:</p><a name="named-metadata__ul_wzy_5db_n3" shape="rect">
<!-- --></a><ul class="ul" id="named-metadata__ul_wzy_5db_n3">
<li class="li"><samp class="ph codeph">nvvm.annotations</samp>: see <a class="xref" href="index.html#global-property-annotation-chapter-11" shape="rect">Global Property Annotation</a></li>
<li class="li"><samp class="ph codeph">nvvmir.version</samp></li>
<li class="li">debugging information</li>
</ul>
<p class="p">The NVVM IR version is specified using a named metadata called <samp class="ph codeph">!nvvmir.version</samp>. The
metadata node consists of two i32 values—the first denotes the major version number and
the second denotes the minor version number. If absent, the version number is assumed to
be 1.0, which can be specified as:
</p><pre xml:space="preserve">!nvvmir.version = !{!0}
!0 = metadata !{ i32 1, i32 0}</pre></div>
</div>
<div class="topic concept nested1" id="parameter-attributes"><a name="parameter-attributes" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#parameter-attributes" name="parameter-attributes" shape="rect">3.9. Parameter Attributes</a></h3>
<div class="body conbody">
<p class="p"> Fully supported, except that <samp class="ph codeph">inreg</samp> is accepted and ignored.
</p>
<p class="p">See <a class="xref" href="index.html#calling-conventions" shape="rect">Calling Conventions</a> for the use of the attributes.
</p>
</div>
</div>
<div class="topic concept nested1" id="garbage-collector-names"><a name="garbage-collector-names" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#garbage-collector-names" name="garbage-collector-names" shape="rect">3.10. Garbage Collector Names</a></h3>
<div class="body conbody">
<p class="p"> Not supported.</p>
</div>
</div>
<div class="topic concept nested1" id="function-attributes"><a name="function-attributes" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#function-attributes" name="function-attributes" shape="rect">3.11. Function Attributes</a></h3>
<div class="body conbody">
<p class="p">Supported:</p><a name="function-attributes__ul_ed5_hgb_n3" shape="rect">
<!-- --></a><ul class="ul" id="function-attributes__ul_ed5_hgb_n3">
<li class="li"><samp class="ph codeph">alwaysinline</samp></li>
<li class="li"><samp class="ph codeph">inlinehint</samp></li>
<li class="li"><samp class="ph codeph">noinline</samp></li>
<li class="li"><samp class="ph codeph">noreturn</samp></li>
<li class="li"><samp class="ph codeph">nounwind</samp></li>
<li class="li"><samp class="ph codeph">optsize</samp></li>
<li class="li"><samp class="ph codeph">readnone</samp></li>
<li class="li"><samp class="ph codeph">readonly</samp></li>
</ul>
<p class="p">Not Supported:</p><a name="function-attributes__ul_fvq_sgb_n3" shape="rect">
<!-- --></a><ul class="ul" id="function-attributes__ul_fvq_sgb_n3">
<li class="li"><samp class="ph codeph">address_safety</samp></li>
<li class="li"><samp class="ph codeph">alignstack</samp></li>
<li class="li"><samp class="ph codeph">nonlazybind</samp></li>
<li class="li"><samp class="ph codeph">naked</samp></li>
<li class="li"><samp class="ph codeph">noimplicitfloat</samp></li>
<li class="li"><samp class="ph codeph">noredzone</samp></li>
<li class="li"><samp class="ph codeph">returns_twice</samp></li>
<li class="li"><samp class="ph codeph">ssp</samp></li>
<li class="li"><samp class="ph codeph">sspreq</samp></li>
<li class="li"><samp class="ph codeph">uwtable</samp></li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="module-level-inline-assembly"><a name="module-level-inline-assembly" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#module-level-inline-assembly" name="module-level-inline-assembly" shape="rect">3.12. Module-Level Inline Assembly</a></h3>
<div class="body conbody">
<p class="p">Supported.</p>
</div>
</div>
<div class="topic concept nested1" id="data-layout"><a name="data-layout" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#data-layout" name="data-layout" shape="rect">3.13. Data Layout</a></h3>
<div class="body conbody">
<p class="p">Only the following data layouts are supported,</p><a name="data-layout__ul_p3j_ps1_n3" shape="rect">
<!-- --></a><ul class="ul" id="data-layout__ul_p3j_ps1_n3">
<li class="li">32-bit
<p class="p"><samp class="ph codeph">e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64
</samp></p>
<p class="p"></p>
<p class="p"></p>
</li>
<li class="li">64-bit
<p class="p"><samp class="ph codeph">e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64
</samp></p>
<p class="p"></p>
</li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="pointer-aliasing-rules"><a name="pointer-aliasing-rules" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#pointer-aliasing-rules" name="pointer-aliasing-rules" shape="rect">3.14. Pointer Aliasing Rules</a></h3>
<div class="body conbody">
<p class="p">Fully supported.</p>
</div>
</div>
<div class="topic concept nested1" id="volatile-memory-access"><a name="volatile-memory-access" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#volatile-memory-access" name="volatile-memory-access" shape="rect">3.15. Volatile Memory Access</a></h3>
<div class="body conbody">
<p class="p">Fully supported. Note that for code generation: <samp class="ph codeph">ld.volatile</samp> and
<samp class="ph codeph">st.volatile</samp> will be generated.
</p>
</div>
</div>
<div class="topic concept nested1" id="atomic-memory-ordering-constraints"><a name="atomic-memory-ordering-constraints" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#atomic-memory-ordering-constraints" name="atomic-memory-ordering-constraints" shape="rect">3.16. Atomic Memory Ordering Constraints</a></h3>
<div class="body conbody">
<p class="p">Atomic loads and stores are not supported. Other atomic operations on other than 32-bit
or 64-bit operands are not supported.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" id="type-system"><a name="type-system" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#type-system" name="type-system" shape="rect">4. Type System</a></h2>
<div class="body conbody">
<p class="p">Fully supported, except for the following:</p><a name="type-system__ul_all_y3b_n3" shape="rect">
<!-- --></a><ul class="ul" id="type-system__ul_all_y3b_n3">
<li class="li">Floating point types <samp class="ph codeph">half</samp>, <samp class="ph codeph">fp128</samp>, <samp class="ph codeph">x86_fp80</samp>, and
<samp class="ph codeph">ppc_fp128</samp> are not supported.
</li>
<li class="li">The <samp class="ph codeph">x86mmx</samp> type is not supported.
</li>
</ul>
</div>
</div>
<div class="topic concept nested0" id="constants"><a name="constants" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#constants" name="constants" shape="rect">5. Constants</a></h2>
<div class="body conbody">
<p class="p">Fully supported, except for the following:</p>
<ul class="ul">
<li class="li"><samp class="ph codeph">blockaddress(@function, %block)</samp> is not supported.
</li>
<li class="li">For a constant expression that is used as the initializer of a global variable
<samp class="ph codeph">@g1</samp>, if the constant expression contains a global identifier
<samp class="ph codeph">@g2</samp>, then the constant expression is supported if
it can be reduced to the form of <samp class="ph codeph">bitcast+offset</samp>, where offset
is an integer number (including <samp class="ph codeph">0</samp>)
</li>
</ul>
</div>
</div>
<div class="topic concept nested0" id="other-values"><a name="other-values" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#other-values" name="other-values" shape="rect">6. Other Values</a></h2>
<div class="body conbody"></div>
<div class="topic concept nested1" id="inline-assembler-expressions"><a name="inline-assembler-expressions" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#inline-assembler-expressions" name="inline-assembler-expressions" shape="rect">6.1. Inline Assembler Expressions</a></h3>
<div class="body conbody">
<div class="p">Inline assembler of PTX instructions is supported, with the following supported
constraints:
<div class="tablenoborder"><a name="inline-assembler-expressions__table_kcr_tpb_n3" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="inline-assembler-expressions__table_kcr_tpb_n3" class="table" frame="border" border="1" rules="all">
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="50%" id="d54e761" rowspan="1" colspan="1">Constraint</th>
<th class="entry" valign="top" width="50%" id="d54e764" rowspan="1" colspan="1">Type</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="50%" headers="d54e761" rowspan="1" colspan="1">c</td>
<td class="entry" valign="top" width="50%" headers="d54e764" rowspan="1" colspan="1">i8</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="50%" headers="d54e761" rowspan="1" colspan="1">h</td>
<td class="entry" valign="top" width="50%" headers="d54e764" rowspan="1" colspan="1">i16</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="50%" headers="d54e761" rowspan="1" colspan="1">r</td>
<td class="entry" valign="top" width="50%" headers="d54e764" rowspan="1" colspan="1">i32</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="50%" headers="d54e761" rowspan="1" colspan="1">l</td>
<td class="entry" valign="top" width="50%" headers="d54e764" rowspan="1" colspan="1">i64</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="50%" headers="d54e761" rowspan="1" colspan="1">f</td>
<td class="entry" valign="top" width="50%" headers="d54e764" rowspan="1" colspan="1">f32</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="50%" headers="d54e761" rowspan="1" colspan="1">d</td>
<td class="entry" valign="top" width="50%" headers="d54e764" rowspan="1" colspan="1">f64</td>
</tr>
</tbody>
</table>
</div>
</div>
<p class="p">The inline asm metadata <samp class="ph codeph">!srcloc</samp> is accepted and ignored.
</p>
<p class="p">The inline asm dialect <samp class="ph codeph">inteldialect</samp> is not supported.
</p>
</div>
</div>
<div class="topic concept nested1" id="metadata-nodes-and-metadata-strings"><a name="metadata-nodes-and-metadata-strings" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#metadata-nodes-and-metadata-strings" name="metadata-nodes-and-metadata-strings" shape="rect">6.2. Metadata Nodes and Metadata Strings</a></h3>
<div class="body conbody">
<p class="p">Fully supported. </p>
<p class="p">The following metadata are understood by the NVVM compiler:</p>
<ul class="ul">
<li class="li">Debug information.</li>
<li class="li"><samp class="ph codeph">pragma unroll</samp><p class="p">Attached to the branch instruction corresponding to
the backedge of a loop.
</p>
<p class="p">The kind of the MDNode is <samp class="ph codeph">pragma</samp>.
The first operand is a metadata string <samp class="ph codeph">!"unroll"</samp> and the second
operand is an <samp class="ph codeph">i32</samp> value which specifies the unroll factor. For
example,
</p><pre xml:space="preserve">br i1 %cond, label %BR1, label %BR2, !pragma !42
!42 = !{ !"unroll", i32 4}</pre></li>
</ul>
<ul class="ul">
<li class="li"><samp class="ph codeph">callalign</samp></li>
</ul>
<p class="p">See <a class="xref" href="index.html#rules-and-restrictions" shape="rect">Rules and Restrictions</a> for <dfn class="term">Calling Conventions</dfn>.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" id="intrinsic-global-variables"><a name="intrinsic-global-variables" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#intrinsic-global-variables" name="intrinsic-global-variables" shape="rect">7. Intrinsic Global Variables</a></h2>
<div class="body conbody">
<ul class="ul">
<li class="li">The <samp class="ph codeph">llvm.used</samp> global variable is supported.
</li>
<li class="li">The <samp class="ph codeph">llvm.compiler.used</samp> global variable is supported
</li>
<li class="li">The <samp class="ph codeph">llvm.global_ctors</samp> global variable is not supported
</li>
<li class="li">The <samp class="ph codeph">llvm.global_dtors</samp> global variable is not supported
</li>
</ul>
</div>
</div>
<div class="topic concept nested0" id="instructions"><a name="instructions" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#instructions" name="instructions" shape="rect">8. Instructions</a></h2>
<div class="body conbody">
<p class="p"></p>
</div>
<div class="topic concept nested1" id="terminator-instructions"><a name="terminator-instructions" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#terminator-instructions" name="terminator-instructions" shape="rect">8.1. Terminator Instructions</a></h3>
<div class="body conbody">
<p class="p">Supported: </p>
<ul class="ul">
<li class="li"><samp class="ph codeph">ret</samp></li>
<li class="li"><samp class="ph codeph">br</samp></li>
<li class="li"><samp class="ph codeph">switch</samp></li>
<li class="li"><samp class="ph codeph">unreachable</samp></li>
</ul>
<p class="p">Unsupported: </p>
<ul class="ul">
<li class="li"><samp class="ph codeph">indirectbr</samp></li>
<li class="li"><samp class="ph codeph">invoke</samp></li>
<li class="li"><samp class="ph codeph">unwind</samp></li>
<li class="li"><samp class="ph codeph">resume</samp></li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="binary-operations"><a name="binary-operations" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#binary-operations" name="binary-operations" shape="rect">8.2. Binary Operations</a></h3>
<div class="body conbody">
<p class="p">Supported: </p>
<ul class="ul">
<li class="li"><samp class="ph codeph">add</samp></li>
<li class="li"><samp class="ph codeph">fadd</samp></li>
<li class="li"><samp class="ph codeph">sub</samp></li>
<li class="li"><samp class="ph codeph">fsub</samp></li>
<li class="li"><samp class="ph codeph">mul</samp></li>
<li class="li"><samp class="ph codeph">fmul</samp></li>
<li class="li"><samp class="ph codeph">udiv</samp></li>
<li class="li"><samp class="ph codeph">sdiv</samp></li>
<li class="li"><samp class="ph codeph">fdiv</samp></li>
<li class="li"><samp class="ph codeph">urem</samp></li>
<li class="li"><samp class="ph codeph">srem</samp></li>
<li class="li"><samp class="ph codeph">frem</samp></li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="bitwise-binary-operations"><a name="bitwise-binary-operations" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#bitwise-binary-operations" name="bitwise-binary-operations" shape="rect">8.3. Bitwise Binary Operations</a></h3>
<div class="body conbody">
<p class="p">Supported: </p>
<ul class="ul">
<li class="li"><samp class="ph codeph">shl</samp></li>
<li class="li"><samp class="ph codeph">lshr</samp></li>
<li class="li"><samp class="ph codeph">ashr</samp></li>
<li class="li"><samp class="ph codeph">and</samp></li>
<li class="li"><samp class="ph codeph">or</samp></li>
<li class="li"><samp class="ph codeph">xor</samp></li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="vector-operations"><a name="vector-operations" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#vector-operations" name="vector-operations" shape="rect">8.4. Vector Operations</a></h3>
<div class="body conbody">
<p class="p">Supported: </p>
<ul class="ul">
<li class="li"><samp class="ph codeph">extractelement</samp></li>
<li class="li"><samp class="ph codeph">insertelement</samp></li>
<li class="li"><samp class="ph codeph">shufflevector</samp></li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="aggregate-operations"><a name="aggregate-operations" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#aggregate-operations" name="aggregate-operations" shape="rect">8.5. Aggregate Operations</a></h3>
<div class="body conbody">
<p class="p">Supported: </p>
<ul class="ul">
<li class="li"><samp class="ph codeph">extractvalue</samp></li>
<li class="li"><samp class="ph codeph">insertvalue</samp></li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="memory-access-and-addressing-operations"><a name="memory-access-and-addressing-operations" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memory-access-and-addressing-operations" name="memory-access-and-addressing-operations" shape="rect">8.6. Memory Access and Addressing Operations</a></h3>
<div class="body conbody">
<p class="p"></p>
</div>
<div class="topic concept nested2" id="alloca-instruction"><a name="alloca-instruction" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#alloca-instruction" name="alloca-instruction" shape="rect">8.6.1. alloca Instruction</a></h3>
<div class="body conbody">
<p class="p">The <samp class="ph codeph">alloca</samp> instruction returns a generic pointer to the local address
space. The number of elements, if specified, must be a compile-time constant, otherwise
it is not supported.
</p>
</div>
</div>
<div class="topic concept nested2" id="load-instruction"><a name="load-instruction" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#load-instruction" name="load-instruction" shape="rect">8.6.2. load Instruction</a></h3>
<div class="body conbody">
<p class="p"><samp class="ph codeph">load atomic</samp> is not supported.
</p>
</div>
</div>
<div class="topic concept nested2" id="store-instruction"><a name="store-instruction" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#store-instruction" name="store-instruction" shape="rect">8.6.3. store Instruction</a></h3>
<div class="body conbody">
<p class="p"><samp class="ph codeph">store atomic</samp> is not supported.
</p>
</div>
</div>
<div class="topic concept nested2" id="fence-instruction"><a name="fence-instruction" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#fence-instruction" name="fence-instruction" shape="rect">8.6.4. fence Instruction</a></h3>
<div class="body conbody">
<p class="p">Not supported. Use NVVM intrinsic functions instead. </p>
</div>
</div>
<div class="topic concept nested2" id="cmpxchg-instruction"><a name="cmpxchg-instruction" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#cmpxchg-instruction" name="cmpxchg-instruction" shape="rect">8.6.5. cmpxchg Instruction</a></h3>
<div class="body conbody">
<p class="p">Supported for <samp class="ph codeph">i32</samp> and <samp class="ph codeph">i64</samp> types, with the following
restrictions:
</p>
<ul class="ul">
<li class="li">
<p class="p">The pointer must be either a global pointer, a shared pointer, or a generic
pointer that points to either the global address space or the shared address
space.
</p>
</li>
<li class="li">
<p class="p">Requires CUDA architecture <samp class="ph codeph">sm_11</samp> or higher.
</p>
</li>
<li class="li">
<p class="p">Use of a shared pointer requires <samp class="ph codeph">sm_12</samp> or higher.
</p>
</li>
<li class="li">
<p class="p">Use of a generic pointer requires <samp class="ph codeph">sm_20</samp> or higher.
</p>
</li>
<li class="li">
<p class="p"><samp class="ph codeph">i64</samp> type with a global pointer requires <samp class="ph codeph">sm_12</samp>
or higher.
</p>
</li>
<li class="li">
<p class="p"><samp class="ph codeph">i64</samp> type with a shared pointer requires <samp class="ph codeph">sm_20</samp>
or higher.
</p>
</li>
</ul>
</div>
</div>
<div class="topic concept nested2" id="atomicrmw-instruction"><a name="atomicrmw-instruction" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#atomicrmw-instruction" name="atomicrmw-instruction" shape="rect">8.6.6. atomicrmw Instruction</a></h3>
<div class="body conbody">
<p class="p"><samp class="ph codeph">nand</samp> is not supported. The other keywords are supported for
<samp class="ph codeph">i32</samp> and <samp class="ph codeph">i64</samp> types, with the following
restrictions.
</p>
<ul class="ul">
<li class="li">The pointer must be either a global pointer, a shared pointer, or a generic pointer
that points to either the <samp class="ph codeph">global</samp> address space or the
<samp class="ph codeph">shared</samp> address space.
</li>
<li class="li">Requires CUDA architecture <samp class="ph codeph">sm_11</samp> or higher.
</li>
<li class="li">Use of a shared pointer requires <samp class="ph codeph">sm_12</samp> or higher.
</li>
<li class="li">Use of a generic pointer requires <samp class="ph codeph">sm_20</samp> or higher.
</li>
<li class="li"><samp class="ph codeph">i64</samp> type <samp class="ph codeph">xchg</samp>, add and sub with a global pointer
require <samp class="ph codeph">sm_12</samp> or higher.
</li>
<li class="li"><samp class="ph codeph">i64</samp> type <samp class="ph codeph">xchg</samp> , add and sub with a shared pointer
require <samp class="ph codeph">sm_20</samp> or higher.
</li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested1" id="conversion-operations"><a name="conversion-operations" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#conversion-operations" name="conversion-operations" shape="rect">8.7. Conversion Operations</a></h3>
<div class="body conbody">
<p class="p">Supported:</p>
<ul class="ul">
<li class="li"><samp class="ph codeph">trunc .. to</samp></li>
<li class="li"><samp class="ph codeph">zext .. to</samp></li>
<li class="li"><samp class="ph codeph">sext .. to</samp></li>
<li class="li"><samp class="ph codeph">fptrunc .. to</samp></li>
<li class="li"><samp class="ph codeph">fpext .. to</samp></li>
<li class="li"><samp class="ph codeph">fptoui .. to</samp></li>
<li class="li"><samp class="ph codeph">fptosi .. to</samp></li>
<li class="li"><samp class="ph codeph">uitofp .. to</samp></li>
<li class="li"><samp class="ph codeph">sitofp .. to</samp></li>
<li class="li"><samp class="ph codeph">ptrtoint .. to</samp></li>
<li class="li"><samp class="ph codeph">inttoptr .. to</samp></li>
<li class="li"><samp class="ph codeph">bitcast .. to</samp><p class="p">See <a class="xref" href="index.html#conversion" shape="rect">Conversion</a> for a special use case of
<samp class="ph codeph">bitcast</samp>.
</p>
</li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="other-operations"><a name="other-operations" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#other-operations" name="other-operations" shape="rect">8.8. Other Operations</a></h3>
<div class="body conbody">
<p class="p">Supported:</p>
<ul class="ul">
<li class="li"><samp class="ph codeph">icmp</samp></li>
<li class="li"><samp class="ph codeph">fcmp</samp></li>
<li class="li"><samp class="ph codeph">phi</samp></li>
<li class="li"><samp class="ph codeph">select</samp></li>
<li class="li"><samp class="ph codeph">call</samp> (See <a class="xref" href="index.html#calling-conventions" shape="rect">Calling Conventions</a> for rules and
restrictions.)
</li>
</ul>
<p class="p">Unsupported:</p>
<ul class="ul">
<li class="li"><samp class="ph codeph">va_arg</samp></li>
<li class="li"><samp class="ph codeph">landingpad</samp></li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested0" id="intrinsic-functions"><a name="intrinsic-functions" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#intrinsic-functions" name="intrinsic-functions" shape="rect">9. Intrinsic Functions</a></h2>
<div class="body conbody">
<p class="p"></p>
</div>
<div class="topic concept nested1" id="variable-argument-handling-intrinsics"><a name="variable-argument-handling-intrinsics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#variable-argument-handling-intrinsics" name="variable-argument-handling-intrinsics" shape="rect">9.1. Variable Argument Handling Intrinsics</a></h3>
<div class="body conbody">
<p class="p">Not supported.</p>
</div>
</div>
<div class="topic concept nested1" id="accurate-garbage-collection-intrinics"><a name="accurate-garbage-collection-intrinics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#accurate-garbage-collection-intrinics" name="accurate-garbage-collection-intrinics" shape="rect">9.2. Accurate Garbage Collection Intrinsics</a></h3>
<div class="body conbody">
<p class="p">Not supported.</p>
</div>
</div>
<div class="topic concept nested1" id="code-generator-intrinics"><a name="code-generator-intrinics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#code-generator-intrinics" name="code-generator-intrinics" shape="rect">9.3. Code Generator Intrinsics</a></h3>
<div class="body conbody">
<p class="p">Not supported.</p>
</div>
</div>
<div class="topic concept nested1" id="standard-c-library-intrinics"><a name="standard-c-library-intrinics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#standard-c-library-intrinics" name="standard-c-library-intrinics" shape="rect">9.4. Standard C Library Intrinsics</a></h3>
<div class="body conbody">
<ul class="ul">
<li class="li"><samp class="ph codeph">llvm.memcpy</samp><p class="p">Supported. Note that the constant address space
cannot be used as the destination since it is read-only.
</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.memmove </samp><p class="p">Supported. Note that the constant address space
cannot be used since it is read-only.
</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.memset</samp><p class="p">Supported. Note that the constant address space
cannot be used since it is read-only.
</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.sqrt</samp><p class="p">Supported for float/double and vector of float/double.
Mapped to PTX <samp class="ph codeph">sqrt.rn.f32</samp> and
<samp class="ph codeph">sqrt.rn.f64</samp>.
</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.powi</samp><p class="p">Not supported.</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.sin</samp><p class="p">Not supported.</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.cos</samp><p class="p">Not supported.</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.pow</samp><p class="p">Not supported.</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.exp</samp><p class="p">Not supported.</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.log</samp><p class="p">Not supported.</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.fma</samp><p class="p">Supported for float and vector of float for
<samp class="ph codeph">sm_20</samp> or higher, mapped to PTX <samp class="ph codeph">fma.rn.f32</samp>.
</p>
<p class="p">Supported for double and vector of double for <samp class="ph codeph">sm_13</samp> or
higher, mapped to PTX <samp class="ph codeph">fma.rn.f64</samp>.
</p>
<p class="p"></p>
</li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="bit-manipulations-intrinics"><a name="bit-manipulations-intrinics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#bit-manipulations-intrinics" name="bit-manipulations-intrinics" shape="rect">Bit Manipulations Intrinsics</a></h3>
<div class="body conbody"><a name="bit-manipulations-intrinics__ul_f2n_1th_n3" shape="rect">
<!-- --></a><ul class="ul" id="bit-manipulations-intrinics__ul_f2n_1th_n3">
<li class="li"><samp class="ph codeph">llvm.bswap</samp><p class="p">Supported for <samp class="ph codeph">i16</samp>,
<samp class="ph codeph">i32</samp>, and <samp class="ph codeph">i64</samp>.
</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.ctpop</samp><p class="p">Supported for <samp class="ph codeph">i8</samp>,
<samp class="ph codeph">i16</samp>, <samp class="ph codeph">i32</samp>, <samp class="ph codeph">i64</samp>, and
vectors of these types, for <samp class="ph codeph">sm_20</samp> or higher.
</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.ctlz</samp><p class="p">Supported for <samp class="ph codeph">i8</samp>,
<samp class="ph codeph">i16</samp>, <samp class="ph codeph">i32</samp>, <samp class="ph codeph">i64</samp>, and vectors of
these types, for <samp class="ph codeph">sm_20</samp> or higher.
</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.cttz</samp><p class="p">Supported for <samp class="ph codeph">i8</samp>,
<samp class="ph codeph">i16</samp>, <samp class="ph codeph">i32</samp>, <samp class="ph codeph">i64</samp>, and vectors of
these types, for <samp class="ph codeph">sm_20</samp> or higher.
</p>
<p class="p"></p>
</li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="arithmetic-with-overflow-intrinsics"><a name="arithmetic-with-overflow-intrinsics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#arithmetic-with-overflow-intrinsics" name="arithmetic-with-overflow-intrinsics" shape="rect">9.6. Arithmetic with Overflow Intrinsics</a></h3>
<div class="body conbody">
<p class="p">Not supported.</p>
</div>
</div>
<div class="topic concept nested1" id="half-precision-floating-point-intrinsics"><a name="half-precision-floating-point-intrinsics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#half-precision-floating-point-intrinsics" name="half-precision-floating-point-intrinsics" shape="rect">9.7. Half Precision Floating Point Intrinsics</a></h3>
<div class="body conbody">
<p class="p">Supported: <samp class="ph codeph">llvm.convert.to.fp16</samp> and
<samp class="ph codeph">llvm.convert.from.fp16</samp>.
</p>
</div>
</div>
<div class="topic concept nested1" id="debugger-intrinsics"><a name="debugger-intrinsics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#debugger-intrinsics" name="debugger-intrinsics" shape="rect">9.8. Debugger Intrinsics</a></h3>
<div class="body conbody">
<p class="p">Supported: <samp class="ph codeph">llvm.dbg.declare</samp> and <samp class="ph codeph">llvm.dbg.value</samp>.
</p>
</div>
</div>
<div class="topic concept nested1" id="exception-handling-intrinsics"><a name="exception-handling-intrinsics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#exception-handling-intrinsics" name="exception-handling-intrinsics" shape="rect">9.9. Exception Handling Intrinsics</a></h3>
<div class="body conbody">
<p class="p">Not supported.</p>
</div>
</div>
<div class="topic concept nested1" id="trampoline-intrinsics"><a name="trampoline-intrinsics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#trampoline-intrinsics" name="trampoline-intrinsics" shape="rect">9.10. Trampoline Intrinsics</a></h3>
<div class="body conbody">
<p class="p">Not supported.</p>
</div>
</div>
<div class="topic concept nested1" id="memory-use-markers"><a name="memory-use-markers" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memory-use-markers" name="memory-use-markers" shape="rect">9.11. Memory Use Markers</a></h3>
<div class="body conbody">
<p class="p">Supported: <samp class="ph codeph">llvm.lifetime.start</samp>, <samp class="ph codeph">llvm.lifetime.end</samp>,
<samp class="ph codeph">llvm.invariant.start</samp>, and <samp class="ph codeph">llvm.invariant.end</samp>.
</p>
</div>
</div>
<div class="topic concept nested1" id="general-intrinsics"><a name="general-intrinsics" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#general-intrinsics" name="general-intrinsics" shape="rect">9.12. General Intrinsics</a></h3>
<div class="body conbody">
<ul class="ul">
<li class="li">
<p class="p"><samp class="ph codeph">llvm.var.annotation</samp></p>
<p class="p">Accepted and ignored.</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.annotation</samp><p class="p">Accepted and ignored.</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.trap</samp><p class="p">Not supported.</p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.stackprotector</samp><p class="p">Not supported. </p>
<p class="p"></p>
</li>
<li class="li"><samp class="ph codeph">llvm.objectsize</samp><p class="p">Not supported.</p>
<p class="p"></p>
</li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested0" id="address-space"><a name="address-space" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#address-space" name="address-space" shape="rect">10. Address Space</a></h2>
<div class="body conbody"></div>
<div class="topic concept nested1" id="address-spaces"><a name="address-spaces" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#address-spaces" name="address-spaces" shape="rect">10.1. Address Spaces</a></h3>
<div class="body conbody">
<p class="p"> NVVM IR has a set of predefined memory address spaces, whose semantics are similar to
those defined in CUDA C/C++, OpenCL C and PTX. Any address space not listed below is not
supported .
</p>
<div class="tablenoborder"><a name="address-spaces__table_ecr_txh_n3" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="address-spaces__table_ecr_txh_n3" class="table" frame="border" border="1" rules="all">
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e1951" rowspan="1" colspan="1">Name</th>
<th class="entry" valign="top" width="33.33333333333333%" id="d54e1954" rowspan="1" colspan="1">Address Space Number</th>
<th class="entry" valign="top" width="33.33333333333333%" id="d54e1957" rowspan="1" colspan="1">Semantics/Example</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1951" rowspan="1" colspan="1">code</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1954" rowspan="1" colspan="1">0</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1957" rowspan="1" colspan="1">functions, code <a name="address-spaces__ul_ttv_dyh_n3" shape="rect">
<!-- --></a><ul class="ul" id="address-spaces__ul_ttv_dyh_n3">
<li class="li">CUDA C/C++ function</li>
<li class="li">OpenCL C function</li>
</ul>
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1951" rowspan="1" colspan="1"> generic </td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1954" rowspan="1" colspan="1">0</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1957" rowspan="1" colspan="1"> Can only be used to qualify the pointee of a pointer <a name="address-spaces__ul_srj_nyh_n3" shape="rect">
<!-- --></a><ul class="ul" id="address-spaces__ul_srj_nyh_n3">
<li class="li"> Pointers in CUDA C/C++</li>
</ul>
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1951" rowspan="1" colspan="1">global</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1954" rowspan="1" colspan="1">1</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1957" rowspan="1" colspan="1"><a name="address-spaces__ul_gpb_xyh_n3" shape="rect">
<!-- --></a><ul class="ul" id="address-spaces__ul_gpb_xyh_n3">
<li class="li">CUDA C/C++ __device__</li>
<li class="li">OpenCL C global</li>
</ul>
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1951" rowspan="1" colspan="1">shared</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1954" rowspan="1" colspan="1">3</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1957" rowspan="1" colspan="1"><a name="address-spaces__ul_ow2_zyh_n3" shape="rect">
<!-- --></a><ul class="ul" id="address-spaces__ul_ow2_zyh_n3">
<li class="li">CUDA C/C++ __shared__</li>
<li class="li">OpenCL C local</li>
</ul>
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1951" rowspan="1" colspan="1">constant</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1954" rowspan="1" colspan="1">4</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1957" rowspan="1" colspan="1"><a name="address-spaces__ul_w3z_1zh_n3" shape="rect">
<!-- --></a><ul class="ul" id="address-spaces__ul_w3z_1zh_n3">
<li class="li">CUDA C/C++ __constant__</li>
<li class="li">OpenCL C constant</li>
</ul>
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1951" rowspan="1" colspan="1">local</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1954" rowspan="1" colspan="1">5</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1957" rowspan="1" colspan="1"><a name="address-spaces__ul_tql_czh_n3" shape="rect">
<!-- --></a><ul class="ul" id="address-spaces__ul_tql_czh_n3">
<li class="li">CUDA C/C++ local</li>
<li class="li">OpenCL C private</li>
</ul>
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1951" rowspan="1" colspan="1"><reserved></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1954" rowspan="1" colspan="1">2, 101 and above</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1957" rowspan="1" colspan="1"> </td>
</tr>
</tbody>
</table>
</div>
<p class="p">Each global variable, that is not an intrinsic global variable, can be declared to
reside in a specific non-zero address space, which can only be one of the following:
<samp class="ph codeph">global</samp>, <samp class="ph codeph">shared</samp> or <samp class="ph codeph">constant</samp>.
</p>
<p class="p">If a non-intrinsic global variable is declared without any address space number or
with the address space number 0, then this global variable resides in address space
<samp class="ph codeph">global</samp> and the pointer of this global variable holds a
generic pointer value.
</p>
<p class="p">The predefined NVVM memory spaces are needed for the language front-ends to model the
memory spaces in the source languages. For example,
</p><pre xml:space="preserve">// CUDA C/C++
__constant__ int c;
__device__ int g;
; NVVM IR
@c = addrspace(4) global i32 0, align 4
@g = addrspace(1) global [2 x i32] zeroinitializer, align 4</pre><p class="p">Address space numbers 2 and 101 or higher are reserved for NVVM compiler internal use
only. No language front-end should generate code that uses these address spaces directly.
</p>
</div>
</div>
<div class="topic concept nested1" id="generic-pointers-and-non-generic-pointers"><a name="generic-pointers-and-non-generic-pointers" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#generic-pointers-and-non-generic-pointers" name="generic-pointers-and-non-generic-pointers" shape="rect">10.2. Generic Pointers and Non-Generic Pointers</a></h3>
<div class="body conbody">
<p class="p"></p>
</div>
<div class="topic concept nested2" id="generic-pointers-vs-non-generic-pointers"><a name="generic-pointers-vs-non-generic-pointers" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#generic-pointers-vs-non-generic-pointers" name="generic-pointers-vs-non-generic-pointers" shape="rect">10.2.1. Generic Pointers vs. Non-generic Pointers</a></h3>
<div class="body conbody">
<p class="p">There are generic pointers and non-generic pointers in NVVM IR. A generic pointer is a
pointer that may point to memory in any address space. A non-generic pointer points to
memory in a specific address space.
</p>
<p class="p">In NVVM IR, a generic pointer has a pointer type with the address space
<samp class="ph codeph">generic</samp>, while a non-generic pointer has a pointer type with a
non-generic address space.
</p>
<p class="p">Note that the address space number for the generic address space is 0—the default in both
NVVM IR and LLVM IR. The address space number for the code address space is also 0.
Function pointers are qualified by address space <samp class="ph codeph">code</samp>
(<samp class="ph codeph">addrspace(0)</samp>).
</p>
<p class="p">Loads/stores via generic pointers are supported, as well as loads/stores via non-generic
pointers. Loads/stores via function pointers are not supported
</p><pre xml:space="preserve">
@a = addrspace(1) global i32 0, align 4 ; 'global' addrspace, @a holds a specific value
@b = global i32 0, align 4 ; 'global' addrspace, @b holds a generic value
@c = addrspace(4) global i32 0, align 4 ; 'constant' addrspace, @c holds a specific value
... = load i32 addrspace(1)* @a, align 4 ; Correct
... = load i32* @a, align 4 ; Wrong
... = load i32* @b, align 4 ; Correct
... = load i32 addrspace(1)* @b, align 4 ; Wrong
... = load i32 addrspace(4)* @c, align4 ; Correct
... = load i32* @c, align 4 ; Wrong</pre></div>
</div>
<div class="topic concept nested2" id="conversion"><a name="conversion" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#conversion" name="conversion" shape="rect">10.2.2. Conversion</a></h3>
<div class="body conbody">
<p class="p">The bit value of a generic pointer that points to a specific object may be different from
the bit value of a specific pointer that points to the same object.
</p>
<div class="p">A generic pointer can be converted into a non-generic pointer using one of the following
NVVM intrinsic functions. <a name="conversion__ul_s4d_zb3_n3" shape="rect">
<!-- --></a><ul class="ul" id="conversion__ul_s4d_zb3_n3">
<li class="li"><samp class="ph codeph">llvm.nvvm.ptr.gen.to.global</samp></li>
<li class="li"><samp class="ph codeph">llvm.nvvm.ptr.gen.to.shared</samp></li>
<li class="li"><samp class="ph codeph">llvm.nvvm.ptr.gen.to.local</samp></li>
<li class="li"><samp class="ph codeph">llvm.nvvm.ptr.gen.to.constant</samp></li>
</ul>
</div>
<div class="p">A non-generic pointer can be converted into a generic pointer using one of the following
NVVM intrinsic functions.
<ul class="ul">
<li class="li"><samp class="ph codeph">llvm.nvvm.ptr.global.to.gen</samp></li>
<li class="li"><samp class="ph codeph">llvm.nvvm.ptr.shared.to.gen</samp></li>
<li class="li"><samp class="ph codeph">llvm.nvvm.ptr.local.to.gen</samp></li>
<li class="li"><samp class="ph codeph">llvm.nvvm.ptr.constant.to.gen</samp></li>
</ul>
</div>
<p class="p">The above conversion intrinsic functions are value changing functions, that is, the bit
patterns of the output value may be different from the input value.
</p>
<p class="p">One can convert a non-generic pointer to a generic pointer and then convert back to the
original non-generic space. But it is illegal to convert to a different non-generic
space.
</p>
<p class="p"><samp class="ph codeph">inttoptr</samp> and <samp class="ph codeph">ptrtoint</samp> are supported.
<samp class="ph codeph">inttoptr</samp> and <samp class="ph codeph">ptrtoin</samp> are value preserving
instructions when the two operands are of the same size.
</p>
<p class="p"><samp class="ph codeph">bitcast</samp> on pointers is supported. <samp class="ph codeph">bitcast</samp> is a value preserving
instruction. Although it is legal to bitcast a non-generic pointer to a non-generic
pointer that points to a different address space, or bitcast between a generic pointer
and a non-generic pointer, the producer of the NVVM IR code should make sure they are
used correctly. For example, accessing a memory location via a generic pointer bitcasted
from a non-generic pointer will result in an undefined result.
</p>
</div>
</div>
<div class="topic concept nested2" id="no-aliasing-between-two-different-specific-address-spaces"><a name="no-aliasing-between-two-different-specific-address-spaces" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#no-aliasing-between-two-different-specific-address-spaces" name="no-aliasing-between-two-different-specific-address-spaces" shape="rect">10.2.3. No Aliasing between Two Different Specific Address Spaces</a></h3>
<div class="body conbody">
<p class="p">Two different specific address spaces do not overlap. NVVM compiler assumes two memory accesses via non-generic pointers that
point to different address spaces are not aliased.
</p>
</div>
</div>
</div>
<div class="topic concept nested1" id="the-alloca-instruction"><a name="the-alloca-instruction" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#the-alloca-instruction" name="the-alloca-instruction" shape="rect">10.3. The alloca Instruction</a></h3>
<div class="body conbody">
<p class="p">The <samp class="ph codeph">alloca</samp> instruction returns a generic pointer that only points to
address space <samp class="ph codeph">local</samp>.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" id="global-property-annotation-chapter-11"><a name="global-property-annotation-chapter-11" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#global-property-annotation-chapter-11" name="global-property-annotation-chapter-11" shape="rect">11. Global Property Annotation</a></h2>
<div class="body conbody">
<p class="p"></p>
</div>
<div class="topic concept nested1" id="overview"><a name="overview" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#overview" name="overview" shape="rect">11.1. Overview</a></h3>
<div class="body conbody">
<p class="p">NVVM uses Named Metadata to annotate IR objects with properties that are otherwise not representable in the IR.
The NVVM IR producers can use the Named Metadata to annotate the IR with properties, which the NVVM compiler can process.
</p>
</div>
</div>
<div class="topic concept nested1" id="representation-of-properties"><a name="representation-of-properties" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#representation-of-properties" name="representation-of-properties" shape="rect">11.2. Representation of Properties</a></h3>
<div class="body conbody">
<p class="p">For each translation unit (that is, per bitcode file), there is a named metadata called
<samp class="ph codeph">nvvm.annotations</samp>.
</p>
<p class="p">This named metadata contains a list of MDNodes.</p>
<p class="p">The first operand of each MDNode is an entity that the node is annotating using the
remaining operands.
</p>
<p class="p">Multiple MDNodes may provide annotations for the same entity, in which case their first
operands will be same.
</p>
<div class="p">The remaining operands of the MDNode are organized in order as <property-name, value>.
<ul class="ul">
<li class="li">The property-name operand is MDString, while the value is <samp class="ph codeph">i32</samp>.
</li>
<li class="li">Starting with the operand after the annotated entity, every alternate operand
specifies a property.
</li>
<li class="li">The operand after a property is its value.
<p class="p">The following is an
example.
</p><pre xml:space="preserve">!nvvm.annotations = !{!12, !13}
!12 = metadata !{void (i32, i32)* @_Z6kernelii, metadata !"kernel", i32 1}
!13 = metadata !{void ()* @_Z7kernel2v, metadata !"kernel", i32 1, metadata !"maxntidx", i32 16}</pre></li>
</ul>
</div>
<p class="p">If two bitcode files are being linked and both have a named metadata
<samp class="ph codeph">nvvm.annotations</samp>, the linked file will have a single merged named
metadata. If both files define properties for the same entity foo , the linked file will
have two MDNodes defining properties for foo . It is illegal for the files to have
conflicting properties for the same entity.
</p>
</div>
</div>
<div class="topic concept nested1" id="supported-properties"><a name="supported-properties" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#supported-properties" name="supported-properties" shape="rect">11.3. Supported Properties</a></h3>
<div class="body conbody">
<div class="tablenoborder"><a name="supported-properties__table_usq_cp3_n3" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="supported-properties__table_usq_cp3_n3" class="table" frame="border" border="1" rules="all">
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e2383" rowspan="1" colspan="1">Property Name</th>
<th class="entry" valign="top" width="33.33333333333333%" id="d54e2386" rowspan="1" colspan="1">Annotated On</th>
<th class="entry" valign="top" width="33.33333333333333%" id="d54e2389" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2383" rowspan="1" colspan="1"><samp class="ph codeph">maxntid{x, y, z}</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2386" rowspan="1" colspan="1">kernel function</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2389" rowspan="1" colspan="1">Maximum expected CTA size from any launch.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2383" rowspan="1" colspan="1"><samp class="ph codeph">reqntid{x, y, z}</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2386" rowspan="1" colspan="1">kernel function</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2389" rowspan="1" colspan="1">Minimum expected CTA size from any launch.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2383" rowspan="1" colspan="1"><samp class="ph codeph">minctasm</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2386" rowspan="1" colspan="1">kernel function</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2389" rowspan="1" colspan="1">Hint/directive to the compiler/driver, asking it to put at least
these many CTAs on an SM.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2383" rowspan="1" colspan="1"><samp class="ph codeph">kernel</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2386" rowspan="1" colspan="1">function</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2389" rowspan="1" colspan="1">Signifies that this function is a kernel function.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2383" rowspan="1" colspan="1"><samp class="ph codeph">align</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2386" rowspan="1" colspan="1">function</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2389" rowspan="1" colspan="1">Signifies that the value in low 16-bits of the 32-bit value contains
alignment of n th parameter type if its alignment is not the natural
alignment. n is specified by high 16-bits of the value. For return type,
n is 0.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2383" rowspan="1" colspan="1"><samp class="ph codeph">texture</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2386" rowspan="1" colspan="1">global variable</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2389" rowspan="1" colspan="1">Signifies that variable is a texture.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2383" rowspan="1" colspan="1"><samp class="ph codeph">surface</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2386" rowspan="1" colspan="1">global variable</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2389" rowspan="1" colspan="1">Signifies that variable is a surface.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2383" rowspan="1" colspan="1"><samp class="ph codeph">managed</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2386" rowspan="1" colspan="1">global variable</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2389" rowspan="1" colspan="1">Signifies that variable is a UVM managed variable.</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="texture-and-surface"><a name="texture-and-surface" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#texture-and-surface" name="texture-and-surface" shape="rect">12. Texture and Surface</a></h2>
<div class="body conbody"></div>
<div class="topic concept nested1" id="texture-variable-and-surface-variable"><a name="texture-variable-and-surface-variable" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#texture-variable-and-surface-variable" name="texture-variable-and-surface-variable" shape="rect">12.1. Texture Variable and Surface Variable</a></h3>
<div class="body conbody">
<p class="p">A texture or a surface variable can be declared/defined as a global variable of
<samp class="ph codeph">i64</samp> type with annotation <samp class="ph codeph">texture</samp> or
<samp class="ph codeph">surface</samp> in the <samp class="ph codeph">global</samp> address space.
</p>
<p class="p">A texture or surface variable must have a name, which must follow identifier naming
conventions.
</p>
<div class="p">It is illegal to store to or load from the address of a texture or surface variable. A
texture or a surface variable may only have the following uses:<a name="texture-variable-and-surface-variable__ul_or2_bq3_n3" shape="rect">
<!-- --></a><ul class="ul" id="texture-variable-and-surface-variable__ul_or2_bq3_n3">
<li class="li">In a metadata node</li>
<li class="li">As an intrinsic function argument as shown below</li>
<li class="li">In <samp class="ph codeph">llvm.used</samp> Global Variable
</li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested1" id="accessing-texture-memory-or-surface-memory"><a name="accessing-texture-memory-or-surface-memory" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#accessing-texture-memory-or-surface-memory" name="accessing-texture-memory-or-surface-memory" shape="rect">12.2. Accessing Texture Memory or Surface Memory</a></h3>
<div class="body conbody">
<p class="p">Texture memory and surface memory can be accessed using texture or surface handles. NVVM
provides the following intrinsic function to get a texture or surface handle from a
texture or surface variable.
</p><pre xml:space="preserve">
delcare i64 %llvm.nvvm.texsurf.handle.p1i64(metadata, i64 addrspace(1)*)
</pre><p class="p">The first argument to the intrinsic is a metadata holding the texture or surface
variable. Such a metadata may hold only one texture or one surface variable. The second
argument to the intrinsic is the texture or surface variable itself. The intrinsic
returns a handle of <samp class="ph codeph">i64</samp> type.
</p>
<p class="p">The returned handle value from the intrinsic call can be used as an operand (with a
constraint of l) in a PTX inline asm to access the texture or surface memory.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" id="nvvm-specific-intrinsic-functions"><a name="nvvm-specific-intrinsic-functions" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#nvvm-specific-intrinsic-functions" name="nvvm-specific-intrinsic-functions" shape="rect">13. NVVM Specific Intrinsic Functions</a></h2>
<div class="body conbody"></div>
<div class="topic concept nested1" id="nvvm-intrin-atomic"><a name="nvvm-intrin-atomic" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#nvvm-intrin-atomic" name="nvvm-intrin-atomic" shape="rect">13.1. Atomic</a></h3>
<div class="body conbody">
<p class="p">Besides the atomic instructions, the following extra atomic intrinsic functions are supported.</p>
<div class="p"><pre xml:space="preserve">
declare float @llvm.nvvm.atomic.load.add.f32.p0f32(float* address, float val)
declare float @llvm.nvvm.atomic.load.add.f32.p1f32(float addrspace(1)* address, float val)
declare float @llvm.nvvm.atomic.load.add.f32.p3f32(float addrspace(3)* address, float val)
</pre></div>
<p class="p">reads the single precision floating point value <samp class="ph codeph">old</samp> located at the address
<samp class="ph codeph">address</samp>, computes <samp class="ph codeph">old+val</samp>, and stores the result back
to memory at the same address. These three operations are performed in one atomic transaction.
The function returns <samp class="ph codeph">old</samp>.
</p>
<div class="p"><pre xml:space="preserve">
declare i32 @llvm.nvvm.atomic.load.inc.32.p0i32(i32* address, i32 val)
declare i32 @llvm.nvvm.atomic.load.inc.32.p1i32(i32 addrspace(1)* address, i32 val)
declare i32 @llvm.nvvm.atomic.load.inc.32.p3i32(i32 addrspace(3)* address, i32 val)
</pre></div>
<p class="p">reads the 32-bit word <samp class="ph codeph">old</samp> located at the address <samp class="ph codeph">address</samp>,
computes <samp class="ph codeph">((old >= val) ? 0 : (old+1))</samp>, and stores the result back to memory
at the same address. These three operations are performed in one atomic transaction.
The function returns <samp class="ph codeph">old</samp>.
</p>
<div class="p"><pre xml:space="preserve">
declare i32 @llvm.nvvm.atomic.load.dec.32.p0i32(i32* address, i32 val)
declare i32 @llvm.nvvm.atomic.load.dec.32.p1i32(i32 addrspace(1)* address, i32 val)
declare i32 @llvm.nvvm.atomic.load.dec.32.p3i32(i32 addrspace(3)* address, i32 val)
</pre></div>
<p class="p">reads the 32-bit word <samp class="ph codeph">old</samp> located at the address <samp class="ph codeph">address</samp>,
computes <samp class="ph codeph">(((old == 0) | (old > val)) ? val : (old-1) )</samp>, and stores the
result back to memory at the same address. These three operations are performed in
one atomic transaction. The function returns <samp class="ph codeph">old</samp>.
</p>
</div>
</div>
<div class="topic concept nested1" id="nvvm-intrin-barrier"><a name="nvvm-intrin-barrier" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#nvvm-intrin-barrier" name="nvvm-intrin-barrier" shape="rect">13.2. Barrier and Memory Fence</a></h3>
<div class="body conbody">
<div class="p"><pre xml:space="preserve">
declare void @llvm.nvvm.barrier0()
</pre></div>
<p class="p">waits until all threads in the thread block have reached this point and all global and shared
memory accesses made by these threads prior to <samp class="ph codeph">llvm.nvvm.barrier0()</samp> are
visible to all threads in the block.
</p>
<div class="p"><pre xml:space="preserve">
declare i32 @llvm.nvvm.barrier0.popc(i32)
</pre></div>
<p class="p">is identical to <samp class="ph codeph">llvm.nvvm.barrier0()</samp> with the additional feature that it
evaluates predicate for all threads of the block and returns the number of threads
for which predicate evaluates to non-zero.
</p>
<div class="p"><pre xml:space="preserve">
declare i32 @llvm.nvvm.barrier0.and(i32)
</pre></div>
<p class="p">is identical to <samp class="ph codeph">llvm.nvvm.barrier0()</samp> with the additional feature that it
evaluates predicate for all threads of the block and returns non-zero if and only if
predicate evaluates to non-zero for all of them.
</p>
<div class="p"><pre xml:space="preserve">
declare i32 @llvm.nvvm.barrier0.or(i32)
</pre></div>
<p class="p">is identical to <samp class="ph codeph">llvm.nvvm.barrier0()</samp> with the additional feature that it
evaluates predicate for all threads of the block and returns non-zero if and only if
predicate evaluates to non-zero for any of them.
</p>
<div class="p"><pre xml:space="preserve">
declare void @llvm.nvvm.membar.cta()
</pre></div>
<p class="p">is a memory fence at the thread block level.</p>
<div class="p"><pre xml:space="preserve">
declare void @llvm.nvvm.membar.gl()
</pre></div>
<p class="p">is a memory fence at the device level.</p>
<div class="p"><pre xml:space="preserve">
declare void @llvm.nvvm.membar.sys()
</pre></div>
<p class="p">is a memory fence at the system level.</p>
</div>
</div>
<div class="topic concept nested1" id="nvvm-intrin-addrsp"><a name="nvvm-intrin-addrsp" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#nvvm-intrin-addrsp" name="nvvm-intrin-addrsp" shape="rect">13.3. Address space conversion</a></h3>
<div class="body conbody">
<p class="p">The following intrinsic functions are provided to support converting pointers from specific address
spaces to the generic address space.
</p>
<div class="p"><pre xml:space="preserve">
declare i8* @llvm.nvvm.ptr.global.to.gen.p0i8.p1i8(i8 addrspace(1))
declare i8* @llvm.nvvm.ptr.shared.to.gen.p0i8.p3i8(i8 addrspace(3)*)
declare i8* @llvm.nvvm.ptr.constant.to.gen.p0i8.p4i8(i8 addrspace(4)*)
declare i8* @llvm.nvvm.ptr.local.to.gen.p0i8.p5i8(i8 addrspace(5)*)
declare i16* @llvm.nvvm.ptr.global.to.gen.p0i16.p1i16(i16 addrspace(1))
declare i16* @llvm.nvvm.ptr.shared.to.gen.p0i16.p3i16(i16 addrspace(3)*)
declare i16* @llvm.nvvm.ptr.constant.to.gen.p0i16.p4i16(i16 addrspace(4)*)
declare i16* @llvm.nvvm.ptr.local.to.gen.p0i16.p5i16(i16 addrspace(5)*)
declare i32* @llvm.nvvm.ptr.global.to.gen.p0i32.p1i32(i32 addrspace(1))
declare i32* @llvm.nvvm.ptr.shared.to.gen.p0i32.p3i32(i32 addrspace(3)*)
declare i32* @llvm.nvvm.ptr.constant.to.gen.p0i32.p4i32(i32 addrspace(4)*)
declare i32* @llvm.nvvm.ptr.local.to.gen.p0i32.p5i32(i32 addrspace(5)*)
declare i64* @llvm.nvvm.ptr.global.to.gen.p0i64.p1i64(i64 addrspace(1))
declare i64* @llvm.nvvm.ptr.shared.to.gen.p0i64.p3i64(i64 addrspace(3)*)
declare i64* @llvm.nvvm.ptr.constant.to.gen.p0i64.p4i64(i64 addrspace(4)*)
declare i64* @llvm.nvvm.ptr.local.to.gen.p0i64.p5i64(i64 addrspace(5)*)
declare f32* @llvm.nvvm.ptr.global.to.gen.p0f32.p1f32(f32 addrspace(1))
declare f32* @llvm.nvvm.ptr.shared.to.gen.p0f32.p3f32(f32 addrspace(3)*)
declare f32* @llvm.nvvm.ptr.constant.to.gen.p0f32.p4f32(f32 addrspace(4)*)
declare f32* @llvm.nvvm.ptr.local.to.gen.p0f32.p5f32(f32 addrspace(5)*)
declare f64* @llvm.nvvm.ptr.global.to.gen.p0f64.p1f64(f64 addrspace(1))
declare f64* @llvm.nvvm.ptr.shared.to.gen.p0f64.p3f64(f64 addrspace(3)*)
declare f64* @llvm.nvvm.ptr.constant.to.gen.p0f64.p4f64(f64 addrspace(4)*)
declare f64* @llvm.nvvm.ptr.local.to.gen.p0f64.p5f64(f64 addrspace(5)*)
</pre></div>
<p class="p">The following intrinsic functions are provided to support converting pointers from the generic
address spaces to specific address spaces.
</p>
<div class="p"><pre xml:space="preserve">
declare i8 addrspace(1)* @llvm.nvvm.ptr.gen.to.global.p1i8.p0i8(i8*)
declare i8 addrspace(3)* @llvm.nvvm.ptr.gen.to.shared.p3i8.p0i8(i8*)
declare i8 addrspace(4)* @llvm.nvvm.ptr.gen.to.constant.p4i8.p0i8(i8*)
declare i8 addrspace(5)* @llvm.nvvm.ptr.gen.to.local.p5i8.p0i8(i8*)
declare i16 addrspace(1)* @llvm.nvvm.ptr.gen.to.global.p1i16.p0i16(i16*)
declare i16 addrspace(3)* @llvm.nvvm.ptr.gen.to.shared.p3i16.p0i16(i16*)
declare i16 addrspace(4)* @llvm.nvvm.ptr.gen.to.constant.p4i16.p0i16(i16*)
declare i16 addrspace(5)* @llvm.nvvm.ptr.gen.to.local.p5i16.p0i16(i16*)
declare i32 addrspace(1)* @llvm.nvvm.ptr.gen.to.global.p1i32.p0i32(i32*)
declare i32 addrspace(3)* @llvm.nvvm.ptr.gen.to.shared.p3i32.p0i32(i32*)
declare i32 addrspace(4)* @llvm.nvvm.ptr.gen.to.constant.p4i32.p0i32(i32*)
declare i32 addrspace(5)* @llvm.nvvm.ptr.gen.to.local.p5i32.p0i32(i32*)
declare i64 addrspace(1)* @llvm.nvvm.ptr.gen.to.global.p1i64.p0i64(i64*)
declare i64 addrspace(3)* @llvm.nvvm.ptr.gen.to.shared.p3i64.p0i64(i64*)
declare i64 addrspace(4)* @llvm.nvvm.ptr.gen.to.constant.p4i64.p0i64(i64*)
declare i64 addrspace(5)* @llvm.nvvm.ptr.gen.to.local.p5i64.p0i64(i64*)
declare f32 addrspace(1)* @llvm.nvvm.ptr.gen.to.global.p1f32.p0f32(f32*)
declare f32 addrspace(3)* @llvm.nvvm.ptr.gen.to.shared.p3f32.p0f32(f32*)
declare f32 addrspace(4)* @llvm.nvvm.ptr.gen.to.constant.p4f32.p0f32(f32*)
declare f32 addrspace(5)* @llvm.nvvm.ptr.gen.to.local.p5f32.p0f32(f32*)
declare f64 addrspace(1)* @llvm.nvvm.ptr.gen.to.global.p1f64.p0f64(f64*)
declare f64 addrspace(3)* @llvm.nvvm.ptr.gen.to.shared.p3f64.p0f64(f64*)
declare f64 addrspace(4)* @llvm.nvvm.ptr.gen.to.constant.p4f64.p0f64(f64*)
declare f64 addrspace(5)* @llvm.nvvm.ptr.gen.to.local.p5f64.p0f64(f64*)
</pre></div>
</div>
</div>
<div class="topic concept nested1" id="nvvm-intrin-spreg"><a name="nvvm-intrin-spreg" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#nvvm-intrin-spreg" name="nvvm-intrin-spreg" shape="rect">13.4. Special Registers</a></h3>
<div class="body conbody">
<p class="p">The following intrinsic functions are provided to support reading special PTX registers.</p>
<div class="p"><pre xml:space="preserve">
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
declare i32 @llvm.nvvm.read.ptx.sreg.warpsize()
</pre></div>
</div>
</div>
<div class="topic concept nested1" id="nvvm-intrin-texture"><a name="nvvm-intrin-texture" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#nvvm-intrin-texture" name="nvvm-intrin-texture" shape="rect">13.5. Texture/surface Access</a></h3>
<div class="body conbody">
<p class="p">The following intrinsic function is provided to support accessing the texture memory and the
surface memory.
</p>
<div class="p"><pre xml:space="preserve">
declare i64 %llvm.nvvm.texsurf.handle.p1i64(metadata, i64 addrspace(1)*)
</pre></div>
<p class="p">See <a class="xref" href="index.html#accessing-texture-memory-or-surface-memory" shape="rect">Accessing Texture Memory or Surface Memory</a> for details.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" id="nvvm-abi-for-ptx"><a name="nvvm-abi-for-ptx" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#nvvm-abi-for-ptx" name="nvvm-abi-for-ptx" shape="rect">14. NVVM ABI for PTX</a></h2>
<div class="topic concept nested1" id="linkage-types-nvvm"><a name="linkage-types-nvvm" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#linkage-types-nvvm" name="linkage-types-nvvm" shape="rect">14.1. Linkage Types</a></h3>
<div class="body conbody">
<p class="p">The following table provides the mapping of NVVM IR linkage types associated with
functions and global variables to PTX linker directives .
</p>
<div class="tablenoborder"><a name="linkage-types-nvvm__table_ech_f53_n3" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="linkage-types-nvvm__table_ech_f53_n3" class="table" frame="border" border="1" rules="all">
<thead class="thead" align="left">
<tr class="row">
<th class="entry" colspan="2" valign="top" id="d54e2825" rowspan="1">LLVM Linkage Type</th>
<th class="entry" valign="top" width="33.33333333333333%" id="d54e2828" rowspan="1" colspan="1">PTX Linker Directive</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2825" rowspan="1"><samp class="ph codeph">private</samp>,
<samp class="ph codeph">internal</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2828" rowspan="1" colspan="1">This is the default linkage type and does not require a linker
directive.
</td>
</tr>
<tr class="row">
<td class="entry" rowspan="4" valign="top" width="33.33333333333333%" headers="d54e2825" colspan="1"><samp class="ph codeph">external</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2825" rowspan="1" colspan="1">function with definition</td>
<td class="entry" rowspan="2" valign="top" width="33.33333333333333%" headers="d54e2828" colspan="1"><samp class="ph codeph">.visible</samp></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2825" rowspan="1" colspan="1">global variable with initialization</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2825" rowspan="1" colspan="1">function without definition</td>
<td class="entry" rowspan="2" valign="top" width="33.33333333333333%" headers="d54e2828" colspan="1"><samp class="ph codeph">.extern</samp></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2825" rowspan="1" colspan="1">global variable without initialization</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2825" rowspan="1"><samp class="ph codeph">available_externally</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2828" rowspan="1" colspan="1"><samp class="ph codeph">.extern</samp></td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2825" rowspan="1"><samp class="ph codeph">linkonce</samp>,
<samp class="ph codeph">linkonce_odr</samp>,
<samp class="ph codeph">weak</samp>,
<samp class="ph codeph">common</samp></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2828" rowspan="1" colspan="1"><samp class="ph codeph">.weak</samp></td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2825" rowspan="1">all other linkage types</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2828" rowspan="1" colspan="1">Not supported.</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="argument-passing-and-return"><a name="argument-passing-and-return" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#argument-passing-and-return" name="argument-passing-and-return" shape="rect">14.2. Argument for Passing and Return</a></h3>
<div class="body conbody">
<p class="p">The following table shows the mapping of function argument and return types in NVVM IR to
PTX types.
</p>
<div class="tablenoborder"><a name="argument-passing-and-return__table_ech_f53_n3" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="argument-passing-and-return__table_ech_f53_n3" class="table" frame="border" border="1" rules="all">
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e2956" rowspan="1" colspan="1">Source Type</th>
<th class="entry" valign="top" width="33.33333333333333%" id="d54e2959" rowspan="1" colspan="1">Size in Bits</th>
<th class="entry" valign="top" width="33.33333333333333%" id="d54e2962" rowspan="1" colspan="1">PTX Type </th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" rowspan="2" valign="top" width="33.33333333333333%" headers="d54e2956" colspan="1">Integer types</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2959" rowspan="1" colspan="1"><= 32</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2962" rowspan="1" colspan="1">
<p class="p"><samp class="ph codeph">.u32</samp> or <samp class="ph codeph">.b32</samp> (zero-extended if unsigned)
</p>
<p class="p"><samp class="ph codeph">.s32</samp> or <samp class="ph codeph">.b32</samp> (sign-extended if signed)
</p>
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2959" rowspan="1" colspan="1">64</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2962" rowspan="1" colspan="1">
<p class="p"><samp class="ph codeph">.u64</samp> or <samp class="ph codeph">.b64</samp> (if unsigned)
</p>
<p class="p"><samp class="ph codeph">.s64</samp> or <samp class="ph codeph">.b64</samp> (if signed)
</p>
</td>
</tr>
<tr class="row">
<td class="entry" rowspan="2" valign="top" width="33.33333333333333%" headers="d54e2956" colspan="1">Pointer types (without <samp class="ph codeph">byval</samp>
attribute)
</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2959" rowspan="1" colspan="1">32</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2962" rowspan="1" colspan="1"><samp class="ph codeph">.u32</samp> or <samp class="ph codeph">.b32</samp></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2959" rowspan="1" colspan="1">64</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2962" rowspan="1" colspan="1"><samp class="ph codeph">.u64</samp> or <samp class="ph codeph">.b64</samp></td>
</tr>
<tr class="row">
<td class="entry" rowspan="2" valign="top" width="33.33333333333333%" headers="d54e2956" colspan="1">Floating-point types</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2959" rowspan="1" colspan="1">32</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2962" rowspan="1" colspan="1"><samp class="ph codeph">.f32</samp> or <samp class="ph codeph">.b32</samp></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2959" rowspan="1" colspan="1">64</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2962" rowspan="1" colspan="1"><samp class="ph codeph">.f64</samp> or <samp class="ph codeph">.b64</samp></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2956" rowspan="1" colspan="1">Aggregate types</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2959" rowspan="1" colspan="1">Any size</td>
<td class="entry" rowspan="3" valign="top" width="33.33333333333333%" headers="d54e2962" colspan="1">
<p class="p"><samp class="ph codeph">.align </samp><em class="ph i">align </em><samp class="ph codeph">.b8 </samp><em class="ph i">name</em>[<em class="ph i">size</em>]
</p>
<p class="p">Where <em class="ph i">align</em> is overall aggregate or vector alignment in bytes,
<em class="ph i">name</em> is variable name associated with aggregate or
vector, and <em class="ph i">size</em> is the aggregate or vector size in
bytes.
</p>
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2956" rowspan="1" colspan="1">Pointer types to aggregate with <samp class="ph codeph">byval</samp>
attribute
</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2959" rowspan="1" colspan="1">32 or 64</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2956" rowspan="1" colspan="1">Vector type</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2959" rowspan="1" colspan="1">Any size</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic reference nested0" id="revision-history"><a name="revision-history" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#revision-history" name="revision-history" shape="rect">A. Revision History</a></h2>
<div class="body refbody">
<div class="section">
<h2 class="title sectiontitle">Version 1.0</h2>
<ul class="ul">
<li class="li">Initial Release</li>
</ul>
</div>
<div class="section">
<h2 class="title sectiontitle">Version 1.1</h2>
<ul class="ul">
<li class="li">Added support for UVM managed variables in global property annotation. See <a class="xref" href="index.html#supported-properties" shape="rect">Supported Properties</a>.
</li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested0" id="notices-header"><a name="notices-header" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#notices-header" name="notices-header" shape="rect">Notices</a></h2>
<div class="topic reference nested1" id="notice"><a name="notice" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#notice" name="notice" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Notice</h3>
<p class="p">ALL NVIDIA DESIGN SPECIFICATIONS, REFERENCE BOARDS, FILES, DRAWINGS, DIAGNOSTICS, LISTS, AND OTHER DOCUMENTS (TOGETHER AND
SEPARATELY, "MATERIALS") ARE BEING PROVIDED "AS IS." NVIDIA MAKES NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE
WITH RESPECT TO THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT, MERCHANTABILITY, AND FITNESS
FOR A PARTICULAR PURPOSE.
</p>
<p class="p">Information furnished is believed to be accurate and reliable. However, NVIDIA Corporation assumes no responsibility for the
consequences of use of such information or for any infringement of patents or other rights of third parties that may result
from its use. No license is granted by implication of otherwise under any patent rights of NVIDIA Corporation. Specifications
mentioned in this publication are subject to change without notice. This publication supersedes and replaces all other information
previously supplied. NVIDIA Corporation products are not authorized as critical components in life support devices or systems
without express written approval of NVIDIA Corporation.
</p>
</div>
</div>
</div>
<div class="topic reference nested1" id="trademarks"><a name="trademarks" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#trademarks" name="trademarks" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Trademarks</h3>
<p class="p">NVIDIA and the NVIDIA logo are trademarks or registered trademarks of NVIDIA Corporation
in the U.S. and other countries. Other company and product names may be trademarks of
the respective companies with which they are associated.
</p>
</div>
</div>
</div>
<div class="topic reference nested1" id="copyright-past-to-present"><a name="copyright-past-to-present" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#copyright-past-to-present" name="copyright-past-to-present" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Copyright</h3>
<p class="p">© <span class="ph">2012</span>-<span class="ph">2014</span> NVIDIA
Corporation. All rights reserved.
</p>
<p class="p">This product includes software developed by the Syncro Soft SRL (http://www.sync.ro/).</p>
</div>
</div>
</div>
</div>
<hr id="contents-end"></hr>
</article>
</div>
</div>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/formatting/common.min.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/google-analytics/google-analytics-write.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/google-analytics/google-analytics-tracker.js"></script>
<script type="text/javascript">var switchTo5x=true;</script><script type="text/javascript" src="http://w.sharethis.com/button/buttons.js"></script><script type="text/javascript">stLight.options({publisher: "998dc202-a267-4d8e-bce9-14debadb8d92", doNotHash: false, doNotCopy: false, hashAddressBar: false});</script></body>
</html>
distrib
>
Mageia
>
5
>
x86_64
>
media
>
nonfree-updates
>
by-pkgid
>
fd8445e7e4d58b8cfe6e0150bd441ee1
>
files
>
1701
