<!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="CUDA Binary Utilities"></meta>
<meta name="abstract" content="The application notes for cuobjdump, nvdisasm, and nvprune."></meta>
<meta name="description" content="The application notes for cuobjdump, nvdisasm, and nvprune."></meta>
<meta name="DC.Coverage" content="Tools"></meta>
<meta name="DC.subject" content="CUDA binary, cuobjdump, nvdisasm, nvprune, Fermi instruction set, Kepler instruction set, Maxwell and Pascal instruction set, Volta instruction set, Turing instruction set"></meta>
<meta name="keywords" content="CUDA binary, cuobjdump, nvdisasm, nvprune, Fermi instruction set, Kepler instruction set, Maxwell and Pascal instruction set, Volta instruction set, Turing instruction set"></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>CUDA Binary Utilities :: 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="//assets.adobedtm.com/b92787824f2e0e9b68dc2e993f9bd995339fe417/satelliteLib-7ba51e58dc61bcb0e9311aadd02a0108ab24cc6c.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/cuda-binary-utilities/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
v10.1.168</a></div>
<div class="category"><a href="index.html" title="CUDA Binary Utilities">CUDA Binary Utilities</a></div>
<ul>
<li>
<div class="section-link"><a href="#overview">1. Overview</a></div>
<ul>
<li>
<div class="section-link"><a href="#cuda-binary">1.1. What is a CUDA Binary?</a></div>
</li>
<li>
<div class="section-link"><a href="#differences">1.2. Differences between cuobjdump and nvdisasm</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#cuobjdump">2. cuobjdump</a></div>
<ul>
<li>
<div class="section-link"><a href="#cuobjdump-usage">2.1. Usage</a></div>
</li>
<li>
<div class="section-link"><a href="#cuobjdump-options">2.2. Command-line Options</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#nvdisasm">3. nvdisasm</a></div>
<ul>
<li>
<div class="section-link"><a href="#nvdisasm-usage">3.1. Usage</a></div>
</li>
<li>
<div class="section-link"><a href="#nvdisasm-options">3.2. Command-line Options</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#instruction-set-ref">4. Instruction Set Reference</a></div>
<ul>
<li>
<div class="section-link"><a href="#fermi">4.1. Fermi Instruction Set</a></div>
</li>
<li>
<div class="section-link"><a href="#kepler">4.2. Kepler Instruction Set</a></div>
</li>
<li>
<div class="section-link"><a href="#maxwell-pascal">4.3. Maxwell and Pascal Instruction Set</a></div>
</li>
<li>
<div class="section-link"><a href="#volta">4.4. Volta Instruction Set</a></div>
</li>
<li>
<div class="section-link"><a href="#turing">4.5. Turing Instruction Set</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#nvprune">5. nvprune</a></div>
<ul>
<li>
<div class="section-link"><a href="#nvprune-usage">5.1. Usage</a></div>
</li>
<li>
<div class="section-link"><a href="#nvprune-options">5.2. Command-line Options</a></div>
</li>
</ul>
</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">CUDA Binary Utilities
(<a href="../../pdf/CUDA_Binary_Utilities.pdf">PDF</a>)
-
v10.1.168
(<a href="https://developer.nvidia.com/cuda-toolkit-archive">older</a>)
-
Last updated April 24, 2019
-
<a href="mailto:CUDAIssues@nvidia.com?subject=CUDA Toolkit Documentation Feedback: CUDA Binary Utilities">Send Feedback</a></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">CUDA Binary Utilities</a></h2>
<div class="body conbody">
<p class="shortdesc">The application notes for cuobjdump, nvdisasm, and nvprune.</p>
</div>
</div>
<div class="topic concept nested0" id="overview"><a name="overview" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#overview" name="overview" shape="rect">1. Overview</a></h2>
<div class="body conbody">
<p class="p">This document introduces <samp class="ph codeph">cuobjdump</samp>, <samp class="ph codeph">nvdisasm</samp>, and <samp class="ph codeph">nvprune</samp>, three CUDA binary tools for Linux(x86 and ARM), Windows, Mac OS and Android.
</p>
</div>
<div class="topic concept nested1" id="cuda-binary"><a name="cuda-binary" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#cuda-binary" name="cuda-binary" shape="rect">1.1. What is a CUDA Binary?</a></h3>
<div class="body conbody">
<p class="p">A CUDA binary (also referred to as cubin) file is an ELF-formatted file which consists of CUDA executable code sections as
well as other sections containing symbols, relocators, debug info, etc. By default, the CUDA compiler
driver <samp class="ph codeph">nvcc</samp> embeds cubin files into the host executable file. But they can also be generated separately by
using the "<samp class="ph codeph">-cubin</samp>" option of <samp class="ph codeph">nvcc</samp>. cubin files are loaded at run time by the CUDA driver API.
</p>
<div class="note note"><span class="notetitle">Note:</span> For more details on cubin files or the CUDA compilation trajectory, refer to <a class="xref" href="http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html" target="_blank" shape="rect">NVIDIA CUDA Compiler Driver NVCC</a>.
</div>
</div>
</div>
<div class="topic concept nested1" id="differences"><a name="differences" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#differences" name="differences" shape="rect">1.2. Differences between <samp class="ph codeph">cuobjdump</samp> and <samp class="ph codeph">nvdisasm</samp></a></h3>
<div class="body conbody">
<p class="p">CUDA provides two binary utilities
for examining and disassembling cubin files and host executables: <samp class="ph codeph">cuobjdump</samp> and <samp class="ph codeph">nvdisasm</samp>.
Basically, <samp class="ph codeph">cuobjdump</samp> accepts both cubin files and host binaries while <samp class="ph codeph">nvdisasm</samp> only
accepts cubin files; but <samp class="ph codeph">nvdisasm</samp> provides richer output options.
</p>
<p class="p">Here's a quick comparison of the two tools:</p>
<p class="p"></p>
<p class="p"></p>
<p class="p"></p>
<div class="tablenoborder"><a name="differences__cuobjdump-nvdisasm-comparison" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="differences__cuobjdump-nvdisasm-comparison" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 1. Comparison of <samp class="ph codeph">cuobjdump</samp> and <samp class="ph codeph">nvdisasm</samp></span></caption>
<thead class="thead" align="left">
<tr class="row" valign="top">
<th class="entry" valign="top" width="60%" id="d54e173" rowspan="1" colspan="1"> </th>
<th class="entry" valign="top" width="20%" id="d54e175" rowspan="1" colspan="1"><samp class="ph codeph">cuobjdump</samp></th>
<th class="entry" valign="top" width="20%" id="d54e179" rowspan="1" colspan="1"><samp class="ph codeph">nvdisasm</samp></th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="60%" headers="d54e173" rowspan="1" colspan="1">Disassemble cubin</td>
<td class="entry" valign="top" width="20%" headers="d54e175" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="20%" headers="d54e179" rowspan="1" colspan="1">Yes</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="60%" headers="d54e173" rowspan="1" colspan="1">Extract ptx and extract and disassemble cubin from the following input files:
<ul class="ul">
<li class="li">Host binaries
<ul class="ul">
<li class="li">Executables</li>
<li class="li">Object files</li>
<li class="li">Static libraries</li>
</ul>
</li>
<li class="li">External fatbinary files</li>
</ul>
</td>
<td class="entry" valign="top" width="20%" headers="d54e175" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="20%" headers="d54e179" rowspan="1" colspan="1">No</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="60%" headers="d54e173" rowspan="1" colspan="1">Control flow analysis and output</td>
<td class="entry" valign="top" width="20%" headers="d54e175" rowspan="1" colspan="1">No</td>
<td class="entry" valign="top" width="20%" headers="d54e179" rowspan="1" colspan="1">Yes</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="60%" headers="d54e173" rowspan="1" colspan="1">Advanced display options</td>
<td class="entry" valign="top" width="20%" headers="d54e175" rowspan="1" colspan="1">No</td>
<td class="entry" valign="top" width="20%" headers="d54e179" rowspan="1" colspan="1">Yes</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="cuobjdump"><a name="cuobjdump" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#cuobjdump" name="cuobjdump" shape="rect">2. cuobjdump</a></h2>
<div class="body conbody">
<p class="p"><samp class="ph codeph">cuobjdump</samp> extracts information from CUDA binary files (both standalone and those embedded in
host binaries) and presents them in human readable format. The output of <samp class="ph codeph">cuobjdump</samp> includes
CUDA assembly code for each kernel, CUDA ELF section headers, string tables, relocators and other CUDA specific sections.
It
also extracts embedded ptx text from host binaries.
</p>
<p class="p">For a list of CUDA assembly instruction set of each GPU architecture, see
<a class="xref" href="index.html#instruction-set-ref" shape="rect">Instruction Set Reference</a>.
</p>
</div>
<div class="topic concept nested1" id="cuobjdump-usage"><a name="cuobjdump-usage" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#cuobjdump-usage" name="cuobjdump-usage" shape="rect">2.1. Usage</a></h3>
<div class="body conbody">
<p class="p"><samp class="ph codeph">cuobjdump</samp> accepts a single input file each time it's run. The basic usage is as following:
</p><pre class="pre screen" xml:space="preserve">cuobjdump [options] <file></pre><p class="p">To disassemble a standalone cubin or cubins embedded in a host executable and show CUDA assembly
of the kernels, use the following command:
</p><pre class="pre screen" xml:space="preserve">cuobjdump -sass <input file></pre><p class="p">To dump cuda elf sections in human readable format from a cubin file, use the following command:</p><pre class="pre screen" xml:space="preserve">cuobjdump -elf <cubin file></pre><p class="p">To extract ptx text from a host binary, use the following command:</p><pre class="pre screen" xml:space="preserve">cuobjdump -ptx <host binary></pre><p class="p">Here's a sample output of <samp class="ph codeph">cuobjdump</samp>:
</p><pre class="pre screen" xml:space="preserve">
$ cuobjdump a.out -ptx -sass
Fatbin elf code:
================
arch = sm_20
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
identifier = add.cu
code for sm_20
Function : _Z3addPiS_S_
.headerflags @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ MOV R6, c[0x0][0x20]; /* 0x2800400080019de4 */
/*0010*/ MOV R7, c[0x0][0x24]; /* 0x280040009001dde4 */
/*0018*/ MOV R2, c[0x0][0x28]; /* 0x28004000a0009de4 */
/*0020*/ MOV R3, c[0x0][0x2c]; /* 0x28004000b000dde4 */
/*0028*/ LDU.E R0, [R6]; /* 0x8c00000000601c85 */
/*0030*/ MOV R4, c[0x0][0x30]; /* 0x28004000c0011de4 */
/*0038*/ LDU.E R2, [R2]; /* 0x8c00000000209c85 */
/*0040*/ MOV R5, c[0x0][0x34]; /* 0x28004000d0015de4 */
/*0048*/ IADD R0, R2, R0; /* 0x4800000000201c03 */
/*0050*/ ST.E [R4], R0; /* 0x9400000000401c85 */
/*0058*/ EXIT; /* 0x8000000000001de7 */
.............................
Fatbin ptx code:
================
arch = sm_20
code version = [4,0]
producer = cuda
host = linux
compile_size = 64bit
compressed
identifier = add.cu
.version 4.0
.target sm_20
.address_size 64
.visible .entry _Z3addPiS_S_(
.param .u64 _Z3addPiS_S__param_0,
.param .u64 _Z3addPiS_S__param_1,
.param .u64 _Z3addPiS_S__param_2
)
{
.reg .s32 %r<4>;
.reg .s64 %rd<7>;
ld.param.u64 %rd1, [_Z3addPiS_S__param_0];
ld.param.u64 %rd2, [_Z3addPiS_S__param_1];
ld.param.u64 %rd3, [_Z3addPiS_S__param_2];
cvta.to.global.u64 %rd4, %rd3;
cvta.to.global.u64 %rd5, %rd2;
cvta.to.global.u64 %rd6, %rd1;
ldu.global.u32 %r1, [%rd6];
ldu.global.u32 %r2, [%rd5];
add.s32 %r3, %r2, %r1;
st.global.u32 [%rd4], %r3;
ret;
}
</pre><p class="p">As shown in the output, the <samp class="ph codeph">a.out</samp> host binary contains cubin and ptx code for sm_20.
</p>
<p class="p">To list cubin files in the host binary use <samp class="ph codeph">-lelf</samp> option:
</p><pre class="pre screen" xml:space="preserve">
$ cuobjdump a.out -lelf
ELF file 1: add_new.sm_20.cubin
ELF file 2: add_new.sm_30.cubin
ELF file 3: add_old.sm_20.cubin
ELF file 4: add_old.sm_30.cubin
</pre><p class="p">To extract all the cubins as files from the host binary use <samp class="ph codeph">-xelf all</samp> option:
</p><pre class="pre screen" xml:space="preserve">
$ cuobjdump a.out -xelf all
Extracting ELF file 1: add_new.sm_20.cubin
Extracting ELF file 2: add_new.sm_30.cubin
Extracting ELF file 3: add_old.sm_20.cubin
Extracting ELF file 4: add_old.sm_30.cubin
</pre><p class="p">To extract the cubin named <samp class="ph codeph">add_new.sm_30.cubin</samp>:
</p><pre class="pre screen" xml:space="preserve">
$ cuobjdump a.out -xelf add_new.sm_30.cubin
Extracting ELF file 1: add_old.sm_20.cubin
</pre><p class="p">To extract only the cubins containing <samp class="ph codeph">_old</samp> in their names:
</p><pre class="pre screen" xml:space="preserve">
$ cuobjdump a.out -xelf _old
Extracting ELF file 1: add_old.sm_20.cubin
Extracting ELF file 2: add_old.sm_30.cubin
</pre><p class="p">You can pass any substring to <samp class="ph codeph">-xelf</samp> and <samp class="ph codeph">-xptx</samp> options. Only the files having the substring in the name will be extracted from the input binary.
</p>
<p class="p">To dump common and per function resource usage information:</p><pre class="pre screen" xml:space="preserve">
$ cuobjdump test.cubin -res-usage
Resource usage:
Common:
GLOBAL:110 CONSTANT[2]:296 CONSTANT[14]:16
Function calculate:
REG:7 STACK:400 SHARED:0 LOCAL:0 CONSTANT[0]:328 TEXTURE:0 SURFACE:0 SAMPLER:0
Function mysurf_func:
REG:18 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:444 TEXTURE:0 SURFACE:1 SAMPLER:0
Function mytexsampler_func:
REG:42 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:472 TEXTURE:4 SURFACE:0 SAMPLER:1
Function mysharedfunc:
REG:30 STACK:0 SHARED:20 LOCAL:0 CONSTANT[0]:192 CONSTANT[16]:440 TEXTURE:0 SURFACE:0 SAMPLER:0
</pre><p class="p">Note that value for REG, TEXTURE, SURFACE and SAMPLER denotes the count and for other resources it denotes no. of byte(s)
used.
</p>
</div>
</div>
<div class="topic concept nested1" id="cuobjdump-options"><a name="cuobjdump-options" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#cuobjdump-options" name="cuobjdump-options" shape="rect">2.2. Command-line Options</a></h3>
<div class="body conbody">
<p class="p"><a class="xref" href="index.html#cuobjdump-options__cuobjdump-options-table" shape="rect">Table 2</a> contains supported command-line options of <samp class="ph codeph">cuobjdump</samp>, along
with a description of what each option does. Each option has a long name and
a short name, which can be used interchangeably.
</p>
<div class="tablenoborder"><a name="cuobjdump-options__cuobjdump-options-table" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="cuobjdump-options__cuobjdump-options-table" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 2. <samp class="ph codeph">cuobjdump</samp> Command-line Options</span></caption>
<thead class="thead" align="left">
<tr class="row" valign="top">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e425" rowspan="1" colspan="1">Option (long)</th>
<th class="entry" valign="top" width="13.333333333333334%" id="d54e428" rowspan="1" colspan="1">Option (short)</th>
<th class="entry" valign="top" width="53.333333333333336%" id="d54e431" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--all-fatbin</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-all</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Dump all fatbin sections. By default will only dump contents
of executable fatbin (if exists), else relocatable fatbin if no
executable fatbin.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--dump-elf</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-elf</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Dump ELF Object sections.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--dump-elf-symbols</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-symbols</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Dump ELF symbol names.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--dump-ptx</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-ptx</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Dump PTX for all listed device functions.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--dump-sass</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-sass</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Dump CUDA assembly for a single cubin file or all cubin files embedded in the binary.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--dump-resource-usage</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-res-usage</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Dump resource usage for each ELF. Useful in getting all the resource usage
information at one place.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--extract-elf <partial file name>,...</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-xelf</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Extract ELF file(s) name containing <partial file name> and save as file(s).
Use 'all' to extract all files. To get the list of ELF files use -lelf
option. Works with host executable/object/library and external fatbin.
All 'dump' and 'list' options are ignored with this option.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--extract-ptx <partial file name>,...</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-xptx</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Extract PTX file(s) name containing <partial file name> and save as file(s).
Use 'all' to extract all files. To get the list of PTX files use -lptx
option. Works with host executable/object/library and external fatbin.
All 'dump' and 'list' options are ignored with this option.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--function <function name>,...</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-fun</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Specify names of device functions whose fat binary structures must be dumped.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--function-index <function index>,...</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-findex</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Specify symbol table index of the function whose fat binary structures must
be dumped.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--gpu-architecture <gpu architecture name></samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-arch</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Specify GPU Architecture for which information should be dumped.
Allowed values for this option: 'sm_20','sm_21','sm_30','sm_32',
'sm_35','sm_37','sm_50','sm_52','sm_53'
<span class="ph">,'sm_60'</span><span class="ph">,'sm_61'</span><span class="ph">,'sm_62'</span><span class="ph">,'sm_70'</span><span class="ph">,'sm_72'</span><span class="ph">,'sm_75'</span>.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--help</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-h</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Print this help information on this tool.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--list-elf</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-lelf</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">List all the ELF files available in the fatbin. Works with host
executable/object/library and external fatbin. All other options are
ignored with this flag. This can be used to select particular ELF with
-xelf option later.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--list-ptx</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-lptx</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">List all the PTX files available in the fatbin. Works with host
executable/object/library and external fatbin. All other options are
ignored with this flag. This can be used to select particular PTX with
-xptx option later.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--options-file <file>,...</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-optf</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Include command line options from specified file.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--sort-functions</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-sort</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Sort functions when dumping sass.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e425" rowspan="1" colspan="1"><samp class="ph codeph">--version</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e428" rowspan="1" colspan="1"><samp class="ph codeph">-V</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e431" rowspan="1" colspan="1">Print version information on this tool.</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="nvdisasm"><a name="nvdisasm" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#nvdisasm" name="nvdisasm" shape="rect">3. nvdisasm</a></h2>
<div class="body conbody">
<p class="p"><samp class="ph codeph">nvdisasm</samp> extracts information from standalone cubin files and presents them
in human readable format. The output of <samp class="ph codeph">nvdisasm</samp> includes CUDA assembly code for
each kernel, listing of ELF data sections and other CUDA specific sections. Output style
and options are controlled through <samp class="ph codeph">nvdisasm</samp> command-line options.
<samp class="ph codeph">nvdisasm</samp> also does control flow analysis to annotate jump/branch targets and makes
the output easier to read.
</p>
<div class="note note"><span class="notetitle">Note:</span><samp class="ph codeph">nvdisasm</samp> requires complete relocation information to do control flow analysis.
If this information is missing from the CUDA binary, either use the <samp class="ph codeph">nvdisasm</samp> option "<samp class="ph codeph">-ndf</samp>"
to turn off control flow analysis, or use the <samp class="ph codeph">ptxas</samp> and <samp class="ph codeph">nvlink</samp> option "<samp class="ph codeph">-preserve-relocs</samp>"
to re-generate the cubin file.
</div>
<p class="p">For a list of CUDA assembly instruction set of each GPU architecture, see
<a class="xref" href="index.html#instruction-set-ref" shape="rect">Instruction Set Reference</a>.
</p>
</div>
<div class="topic concept nested1" id="nvdisasm-usage"><a name="nvdisasm-usage" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#nvdisasm-usage" name="nvdisasm-usage" shape="rect">3.1. Usage</a></h3>
<div class="body conbody">
<p class="p"><samp class="ph codeph">nvdisasm</samp> accepts a single input file each time it's run. The basic usage is as
following:
</p><pre class="pre screen" xml:space="preserve">nvdisasm [options] <input cubin file></pre><p class="p">To get the control flow graph of a kernel, use the following:</p><pre class="pre screen" xml:space="preserve">nvdisasm -cfg <input cubin file></pre><p class="p">Here's a sample output of <samp class="ph codeph">nvdisasm</samp>:
</p><pre class="pre screen" xml:space="preserve">
.headerflags @"EF_CUDA_TEXMODE_UNIFIED EF_CUDA_64BIT_ADDRESS EF_CUDA_SM30 EF_CUDA_PTX_SM(EF_CUDA_SM30) "
//--------------------- .nv.info --------------------------
.section .nv.info,"",@"SHT_CUDA_INFO "
.align 4
......
//--------------------- .text._Z4addXPii --------------------------
.section .text._Z4addXPii,"ax",@progbits
.sectioninfo @"SHI_REGISTERS=11 "
.align 4
.global _Z4addXPii
.type _Z4addXPii,@function
.size _Z4addXPii,(.L_19 - _Z4addXPii)
.other _Z4addXPii,@"STO_CUDA_ENTRY STV_DEFAULT "
_Z4addXPii:
.text._Z4addXPii:
/*0008*/ MOV R1, c[0x0][0x44];
/*0010*/ ISUB R1, R1, 0x8;
/*0018*/ MOV R0, c[0x0][0x148];
/*0020*/ IADD R6.CC, R1, c[0x0][0x24];
/*0028*/ ISETP.LT.AND P0, PT, R0, 0x1, PT;
/*0030*/ MOV R8, c[0x0][0x140];
/*0038*/ MOV R9, c[0x0][0x144];
/*0048*/ IADD.X R7, RZ, RZ;
/*0050*/ ISUB R10, R6, c[0x0][0x24];
/*0058*/ @P0 BRA `(.L_2);
/*0060*/ LD.E R0, [R8];
/*0068*/ MOV R2, RZ;
/*0070*/ NOP;
/*0078*/ NOP;
.L_3:
/*0088*/ IADD R2, R2, 0x1;
/*0090*/ MOV R3, R0;
/*0098*/ IADD R0, R0, 0x1;
/*00a0*/ ISETP.LT.AND P0, PT, R2, c[0x0][0x148], PT;
/*00a8*/ @P0 BRA `(.L_3);
/*00b0*/ IADD R0, R3, 0x1;
/*00b8*/ ST.E [R8], R0;
.L_2:
/*00c8*/ S2R R0, SR_TID.X;
/*00d0*/ ISETP.NE.AND P0, PT, R0, RZ, PT;
/*00d8*/ @P0 EXIT ;
/*00e0*/ LD.E R0, [R8];
/*00e8*/ MOV R4, c[0x0][0xf0];
/*00f0*/ MOV R5, c[0x0][0xf4];
/*00f8*/ STL [R10], R0;
/*0108*/ JCAL `(vprintf);
/*0110*/ EXIT ;
.L_4:
/*0118*/ BRA `(.L_4);
.L_19:
//--------------------- SYMBOLS --------------------------
.type vprintf,@function
</pre><p class="p"><samp class="ph codeph">nvdisasm</samp> is capable of generating control flow of CUDA assembly in the format of DOT graph description language.
The output of the control flow from nvdisasm can be directly imported to a DOT graph visualization tool such as
<a class="xref" href="http://www.graphviz.org" target="_blank" shape="rect">Graphviz</a>.
</p>
<div class="note note"><span class="notetitle">Note:</span> This feature is only supported on cubins generated for Compute Capability 3.0 and later.
</div>
<p class="p">Here's how you can generate a PNG image (<samp class="ph codeph">cfg.png</samp>) of the control flow of the above cubin (<samp class="ph codeph">a.cubin</samp>)
with <samp class="ph codeph">nvdisasm</samp> and Graphviz:
</p><pre class="pre screen" xml:space="preserve">nvdisasm -cfg a.cubin | dot -ocfg.png -Tpng</pre><p class="p">Here's the generated graph:</p>
<div class="fig fignone" id="nvdisasm-usage__control-flow-graph"><a name="nvdisasm-usage__control-flow-graph" shape="rect">
<!-- --></a><span class="figcap">Figure 1. Control Flow Graph</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/cfg.png" alt="Control Flow Graph"></img></div><br clear="none"></br></div>
<p class="p"><samp class="ph codeph">nvdisasm</samp> is capable of showing the register (CC, general and predicate) liveness range information.
For each line of CUDA assembly, <samp class="ph codeph">nvdisasm</samp> displays whether a given device register was assigned, accessed, live or re-assigned.
It also shows the total number of registers used. This is useful if the user
is interested in the life range of any particular register, or register usage in general.
</p>
<div class="note note"><span class="notetitle">Note:</span> This feature is only supported on cubins generated for Compute Capability 3.0 and later.
</div>
<p class="p">Here's a sample output (left columns are omitted):</p><pre xml:space="preserve">
// +------+---------------+-----+
// | CC | GPR |PRED |
// | | 0000000000 | |
// | # 01 | # 0123456789 | # 0 |
// +------+---------------+-----+
_main10acosParams // | | | |
_main10acosParams,@function // | | | |
_main10acosParams,(.L_17 - _Z9acos_main10acosParams) // | | | |
_main10acosParams,@"STO_CUDA_ENTRY STV_DEFAULT" // | | | |
// | | | |
// | | | |
MOV R1, c[0x0][0x44]; // | | 1 ^ | |
S2R R0, SR_CTAID.X; // | | 2 ^: | |
S2R R3, SR_TID.X; // | | 3 :: ^ | |
IMAD R3, R0, c[0x0][0x28], R3; // | | 3 v: x | |
MOV R0, c[0x0][0x28]; // | | 3 ^: : | |
ISETP.GE.AND P0, PT, R3, c[0x0][0x150], PT; // | | 3 :: v | 1 ^ |
IMUL R0, R0, c[0x0][0x34]; // | | 3 x: : | 1 : |
@P0 EXIT; // | | 3 :: : | 1 v |
MOV32I R8, 0x4; // | | 4 :: : ^ | |
MOV32I R9, 0x3c94d2e9; // | | 5 :: : :^ | |
NOP; // | | 5 :: : :: | |
NOP; // | | 5 :: : :: | |
NOP; // | | 5 :: : :: | |
NOP; // | | 5 :: : :: | |
// | | 5 :: : :: | |
IMAD R6.CC, R3, R8, c[0x0][0x140]; // | 1 ^ | 6 :: v ^ v: | |
IMAD.HI.X R7, R3, R8, c[0x0][0x144]; // | 1 v | 7 :: v :^v: | |
LD.E R2, [R6]; // | | 8 ::^: vv:: | |
FADD.FTZ R4, -|R2|, 1; // | | 7 ::v:^ :: | |
FSETP.GT.FTZ.AND P0, PT, |R2|, c[0x2][0x0], PT; // | | 7 ::v:: :: | 1 ^ |
FMUL.FTZ R4, R4, 0.5; // | | 7 ::::x :: | 1 : |
F2F.FTZ.F32.F32 R5, |R2|; // | | 8 ::v::^ :: | 1 : |
MUFU.RSQ R4, R4; // | | 8 ::::x: :: | 1 : |
@P0 MUFU.RCP R5, R4; // | | 8 ::::v^ :: | 1 v |
FMUL.FTZ R4, R5, R5; // | | 8 ::::^v :: | 1 : |
IMAD R6.CC, R3, R8, c[0x0][0x148]; // | 1 ^ | 9 :::v::^ v: | 1 : |
FFMA.FTZ R7, R4, c[0x2][0x4], R9; // | 1 : | 10 ::::v::^:v | 1 : |
FFMA.FTZ R7, R7, R4, c[0x2][0x8]; // | 1 : | 10 ::::v::x:: | 1 : |
FFMA.FTZ R7, R7, R4, c[0x2][0xc]; // | 1 : | 10 ::::v::x:: | 1 : |
FFMA.FTZ R7, R7, R4, c[0x2][0x10]; // | 1 : | 10 ::::v::x:: | 1 : |
FMUL.FTZ R4, R7, R4; // | 1 : | 10 ::::x::v:: | 1 : |
IMAD.HI.X R7, R3, R8, c[0x0][0x14c]; // | 1 v | 10 :::v:::^v: | 1 : |
FFMA.FTZ R4, R4, R5, R5; // | | 10 ::::xv:::: | 1 : |
IADD R3, R3, R0; // | | 9 v::x: :::: | 1 : |
FADD32I.FTZ R5, -R4, 1.5707963705062866211; // | | 10 ::::v^:::: | 1 : |
@P0 FADD.FTZ R5, R4, R4; // | | 10 ::::v^:::: | 1 v |
ISETP.LT.AND P0, PT, R3, c[0x0][0x150], PT; // | | 9 :::v ::::: | 1 ^ |
FADD32I.FTZ R4, -R5, 3.1415927410125732422; // | | 10 ::::^v:::: | 1 : |
FCMP.LT.FTZ R2, R4, R5, R2; // | | 10 ::x:vv:::: | 1 : |
ST.E [R6], R2; // | | 8 ::v: vv:: | 1 : |
@P0 BRA `(.L_1); // | | 5 :: : :: | 1 v |
MOV RZ, RZ; // | | 1 : | |
EXIT; // | | 1 : | |
// +......+...............+.....+
BRA `(.L_2); // | | | |
// +------+---------------+-----+
// Legend:
// ^ : Register assignment
// v : Register usage
// x : Register usage and reassignment
// : : Register in use
// <space> : Register not in use
// # : Number of occupied registers
</pre></div>
</div>
<div class="topic concept nested1" id="nvdisasm-options"><a name="nvdisasm-options" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#nvdisasm-options" name="nvdisasm-options" shape="rect">3.2. Command-line Options</a></h3>
<div class="body conbody">
<p class="p"><a class="xref" href="index.html#nvdisasm-options__nvdisasm-options-table" shape="rect">Table 3</a> contains the supported command-line options of <samp class="ph codeph">nvdisasm</samp>, along
with a description of what each option does. Each option has a long name and
a short name, which can be used interchangeably.
</p>
<div class="tablenoborder"><a name="nvdisasm-options__nvdisasm-options-table" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="nvdisasm-options__nvdisasm-options-table" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 3. nvdisasm Command-line Options</span></caption>
<thead class="thead" align="left">
<tr class="row" valign="top">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e867" rowspan="1" colspan="1">Option (long)</th>
<th class="entry" valign="top" width="13.333333333333334%" id="d54e870" rowspan="1" colspan="1">Option (short)</th>
<th class="entry" valign="top" width="53.333333333333336%" id="d54e873" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--base-address <value></samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-base</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Specify the logical base address of the image to disassemble.
This option is only valid when disassembling a raw instruction
binary (see option '--binary'), and is ignored when disassembling
an Elf file.
Default value: 0.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--binary <SMxy></samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-b</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">When this option is specified, the input file is assumed to
contain a raw instruction binary, that is, a sequence of binary
instruction encodings as they occur in instruction memory. The
value of this option must be the asserted architecture of the
raw binary.
Allowed values for this option: 'SM20','SM21','SM30','SM32',
'SM35','SM37','SM50','SM52','SM53'<span class="ph">,'SM60'</span><span class="ph">,'SM61'</span><span class="ph">,'SM62'</span><span class="ph">,'SM70'</span><span class="ph">,'SM72'</span><span class="ph">,'SM75'</span>.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--cuda-function-index <symbol index>,... </samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-fun</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Restrict the output to the CUDA functions represented by symbols with the
given indices. The CUDA function for a given symbol is the enclosing section.
This only restricts executable sections; all other sections will still be
printed..
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--help</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-h</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Print this help information on this tool.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--life-range-mode</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-lrm</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">This option implies option '--print-life-ranges', and determines how register
live range info should be printed. 'count': Not at all, leaving only the # column
(number of live registers); 'wide': Columns spaced out for readability (default);
'narrow': A one-character column for each register, economizing on table width
Allowed values for this option: 'count','narrow','wide'.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--no-dataflow</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-ndf</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Disable dataflow analyzer after disassembly. Dataflow analysis
is normally enabled to perform branch stack analysis and annotate
all instructions that jump via the GPU branch stack with inferred
branch target labels. However, it may occasionally fail when certain
restrictions on the input nvelf/cubin are not met.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--no-vliw</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-novliw</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Conventional mode; disassemble paired instructions in normal syntax, instead
of VLIW syntax.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--options-file <file>,...</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-optf</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Include command line options from specified file.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--output-control-flow-graph</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-cfg</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">When specified, output the control flow graph in a format
consumable by graphviz tools (such as dot).
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--print-code</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-c</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Only print code sections.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--print-instruction-encoding</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-hex</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">When specified, print the encoding bytes after each
disassembled operation.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--print-life-ranges</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-plr</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Print register life range information in a trailing column in the produced
disassembly.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--print-line-info</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-g</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Annotate disassembly with source line information obtained from .debug_line
section, if present.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--print-raw</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-raw</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Print the disassembly without any attempt to beautify it.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--separate-functions</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-sf</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Separate the code corresponding with function symbols by
some new lines to let them stand out in the printed disassembly.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e867" rowspan="1" colspan="1"><samp class="ph codeph">--version</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e870" rowspan="1" colspan="1"><samp class="ph codeph">-V</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e873" rowspan="1" colspan="1">Print version information on this tool.</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="instruction-set-ref"><a name="instruction-set-ref" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#instruction-set-ref" name="instruction-set-ref" shape="rect">4. Instruction Set Reference</a></h2>
<div class="body conbody">
<p class="p">This is an instruction set reference for NVIDIA<sup>®</sup>
GPU architectures Fermi, Kepler, Maxwell, Pascal<span class="ph">, Volta and Turing</span>.
</p>
</div>
<div class="topic concept nested1" id="fermi"><a name="fermi" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#fermi" name="fermi" shape="rect">4.1. Fermi Instruction Set</a></h3>
<div class="body conbody">
<div class="p">The Fermi architecture (Compute Capability 2.x) has the following instruction
set format:
<pre class="pre screen" xml:space="preserve"><samp class="ph codeph">(instruction) (destination) (source1), (source2) ...</samp></pre></div>
<div class="p">Valid destination and source locations include:
<ul class="ul">
<li class="li">RX for registers</li>
<li class="li">SRX for special system-controlled register</li>
<li class="li">PX for condition register</li>
<li class="li">c[X][Y] for constant memory</li>
</ul>
</div>
<p class="p"><a class="xref" href="index.html#fermi__fermi-instruction-set" shape="rect">Table 4</a> lists valid instructions for the Fermi GPUs.
</p>
<div class="tablenoborder"><a name="fermi__fermi-instruction-set" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="fermi__fermi-instruction-set" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 4. Fermi Instruction Set</span></caption>
<thead class="thead" align="left">
<tr class="row" valign="top">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e1195" rowspan="1" colspan="1">Opcode</th>
<th class="entry" valign="top" width="66.66666666666666%" id="d54e1198" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e1195 d54e1198" rowspan="1"><strong class="ph b">Floating Point Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">FFMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP32 Fused Multiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">FADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP32 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">FCMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP32 Compare</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">FMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP32 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">FMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP32 Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">FSWZ</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP32 Swizzle</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">FSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP32 Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">FSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP32 Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">RRO</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP Range Reduction Operator</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">MUFU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP Multi-Function Operator</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">DFMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP64 Fused Multiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">DADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP64 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">DMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP64 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">DMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP64 Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">DSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP64 Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">DSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">FP64 Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e1195 d54e1198" rowspan="1"><strong class="ph b">Integer Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">IMAD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Multiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">IMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">IADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">ISCADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Scaled Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">ISAD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Sum Of Abs Diff</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">IMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">BFE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Bit Field Extract</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">BFI</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Bit Field Insert</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">SHR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Shift Right</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">SHL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Shift Left</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Logic Op</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">FLO</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Find Leading One</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">ISET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">ISETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">ICMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer Compare and Select</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">POPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Population Count</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e1195 d54e1198" rowspan="1"><strong class="ph b">Conversion Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">F2F</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Float to Float</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">F2I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Float to Integer</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">I2F</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer to Float</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">I2I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Integer to Integer</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e1195 d54e1198" rowspan="1"><strong class="ph b">Movement Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">MOV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">SEL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Conditional Select/Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">PRMT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Permute</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e1195 d54e1198" rowspan="1"><strong class="ph b">Predicate/CC Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">P2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Predicate to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">R2P</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Register to Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">CSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">CC Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">CSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">CC Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">PSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Predicate Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">PSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Predicate Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e1195 d54e1198" rowspan="1"><strong class="ph b">Texture Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">TEX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Texture Fetch</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">TLD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Texture Load</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">TLD4</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Texture Load 4 Texels</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">TXQ</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Texture Query</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e1195 d54e1198" rowspan="1"><strong class="ph b">Compute Load/Store Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LDC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Load from Constant</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Load from Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LDU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Load Uniform</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LDL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Load from Local Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LDS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Load from Shared Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LDLK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Load and Lock</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LDSLK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Load from Shared Memory and Lock</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LD_LDU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">LD_LDU is a combination of a generic load LD with a load uniform LDU</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LDS_LDU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">LDS_LDU is combination of a Shared window load LDS with a load uniform LDU.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">ST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Store to Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">STL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Store to Local Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">STUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Store and Unlock</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">STS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Store to Shared Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">STSUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Store to Shared Memory and Unlock</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">ATOM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Atomic Memory Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">RED</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Atomic Memory Reduction Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">CCTL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Cache Control</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">CCTLL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Cache Control (Local)</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">MEMBAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Memory Barrier</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e1195 d54e1198" rowspan="1"><strong class="ph b">Surface Memory Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">SULD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Surface Load</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">SULEA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Surface Load Effective Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">SUST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Surface Store</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">SURED</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Surface Reduction</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">SUQ</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Surface Query</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e1195 d54e1198" rowspan="1"><strong class="ph b">Control Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">BRA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Branch to Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">BRX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Branch to Relative Indexed Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">JMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Jump to Absolute Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">JMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Jump to Absolute Indexed Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">CAL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Call to Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">JCAL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Call to Absolute Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">RET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Return from Call</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">BRK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Break from Loop</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">CONT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Continue in Loop</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LONGJMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Long Jump</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">SSY</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Set Sync Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">PBK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Pre-Break Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">PCNT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Pre-Continue Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">PRET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Pre-Return Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">PLONGJMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Pre-Long-Jump Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">BPT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Breakpoint/Trap</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">EXIT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Exit Program</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e1195 d54e1198" rowspan="1"><strong class="ph b">Miscellaneous Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">NOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">No Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">S2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Special Register to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">B2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Barrier to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">LEPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Load Effective PC</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">BAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Barrier Synchronization</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e1195" rowspan="1" colspan="1">VOTE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e1198" rowspan="1" colspan="1">Query condition across threads</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="kepler"><a name="kepler" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#kepler" name="kepler" shape="rect">4.2. Kepler Instruction Set</a></h3>
<div class="body conbody">
<div class="p">The Kepler architecture (Compute Capability 3.x) has the following instruction
set format:
<pre class="pre screen" xml:space="preserve"><samp class="ph codeph">(instruction) (destination) (source1), (source2) ...</samp></pre></div>
<div class="p">Valid destination and source locations include:
<ul class="ul">
<li class="li">RX for registers</li>
<li class="li">SRX for special system-controlled registers</li>
<li class="li">PX for condition registers</li>
<li class="li">c[X][Y] for constant memory</li>
</ul>
</div>
<p class="p"><a class="xref" href="index.html#kepler__kepler-instruction-set" shape="rect">Table 5</a> lists valid instructions for the Kepler GPUs.
</p>
<div class="tablenoborder"><a name="kepler__kepler-instruction-set" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="kepler__kepler-instruction-set" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 5. Kepler Instruction Set</span></caption>
<thead class="thead" align="left">
<tr class="row" valign="top">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e2213" rowspan="1" colspan="1">Opcode</th>
<th class="entry" valign="top" width="66.66666666666666%" id="d54e2216" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2213 d54e2216" rowspan="1"><strong class="ph b">Floating Point Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">FFMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP32 Fused Multiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">FADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP32 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">FCMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP32 Compare</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">FMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP32 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">FMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP32 Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">FSWZ</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP32 Swizzle</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">FSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP32 Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">FSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP32 Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">FCHK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP32 Division Test</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">RRO</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP Range Reduction Operator</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">MUFU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP Multi-Function Operator</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">DFMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP64 Fused Multiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">DADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP64 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">DMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP64 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">DMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP64 Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">DSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP64 Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">DSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">FP64 Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2213 d54e2216" rowspan="1"><strong class="ph b">Integer Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">IMAD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Multiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">IMADSP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Extract Multiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">IMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">IADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">ISCADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Scaled Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">ISAD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Sum Of Abs Diff</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">IMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">BFE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Bit Field Extract</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">BFI</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Bit Field Insert</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SHR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Shift Right</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SHL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Shift Left</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SHF</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Funnel Shift</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">LOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Logic Op</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">FLO</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Find Leading One</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">ISET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">ISETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">ICMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer Compare and Select</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">POPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Population Count</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2213 d54e2216" rowspan="1"><strong class="ph b">Conversion Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">F2F</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Float to Float</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">F2I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Float to Integer</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">I2F</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer to Float</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">I2I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Integer to Integer</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2213 d54e2216" rowspan="1"><strong class="ph b">Movement Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">MOV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SEL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Conditional Select/Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">PRMT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Permute</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SHFL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Warp Shuffle</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2213 d54e2216" rowspan="1"><strong class="ph b">Predicate/CC Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">P2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Predicate to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">R2P</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Register to Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">CSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">CC Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">CSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">CC Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">PSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Predicate Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">PSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Predicate Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2213 d54e2216" rowspan="1"><strong class="ph b">Texture Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">TEX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Texture Fetch</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">TLD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Texture Load</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">TLD4</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Texture Load 4 Texels</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">TXQ</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Texture Query</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2213 d54e2216" rowspan="1"><strong class="ph b">Compute Load/Store Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">LDC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Load from Constant</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">LD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Load from Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">LDG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Non-coherent Global Memory Load</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">LDL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Load from Local Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">LDS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Load from Shared Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">LDSLK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Load from Shared Memory and Lock</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">ST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Store to Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">STL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Store to Local Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">STS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Store to Shared Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">STSCUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Store to Shared Memory Conditionally and Unlock</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">ATOM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Atomic Memory Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">RED</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Atomic Memory Reduction Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">CCTL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Cache Control</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">CCTLL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Cache Control (Local)</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">MEMBAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Memory Barrier</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2213 d54e2216" rowspan="1"><strong class="ph b">Surface Memory Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SUCLAMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Surface Clamp </td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SUBFM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Surface Bit Field Merge </td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SUEAU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Surface Effective Address </td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SULDGA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Surface Load Generic Address </td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SUSTGA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Surface Store Generic Address </td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2213 d54e2216" rowspan="1"><strong class="ph b">Control Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">BRA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Branch to Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">BRX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Branch to Relative Indexed Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">JMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Jump to Absolute Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">JMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Jump to Absolute Indexed Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">CAL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Call to Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">JCAL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Call to Absolute Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">RET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Return from Call</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">BRK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Break from Loop</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">CONT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Continue in Loop</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">SSY</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Set Sync Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">PBK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Pre-Break Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">PCNT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Pre-Continue Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">PRET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Pre-Return Relative Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">BPT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Breakpoint/Trap</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">EXIT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Exit Program</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e2213 d54e2216" rowspan="1"><strong class="ph b">Miscellaneous Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">NOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">No Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">S2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Special Register to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">B2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Barrier to Register </td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">BAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Barrier Synchronization</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2213" rowspan="1" colspan="1">VOTE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e2216" rowspan="1" colspan="1">Query condition across threads</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="maxwell-pascal"><a name="maxwell-pascal" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#maxwell-pascal" name="maxwell-pascal" shape="rect">4.3. Maxwell and Pascal Instruction Set</a></h3>
<div class="body conbody">
<div class="p">The Maxwell (Compute Capability 5.x) and the Pascal (Compute Capability 6.x) architectures have the following instruction
set format:
<pre class="pre screen" xml:space="preserve"><samp class="ph codeph">(instruction) (destination) (source1), (source2) ...</samp></pre></div>
<div class="p">Valid destination and source locations include:
<ul class="ul">
<li class="li">RX for registers</li>
<li class="li">SRX for special system-controlled registers</li>
<li class="li">PX for condition registers</li>
<li class="li">c[X][Y] for constant memory</li>
</ul>
</div>
<p class="p"><a class="xref" href="index.html#maxwell-pascal__maxwell-pascal-instruction-set" shape="rect">Table 6</a> lists valid instructions for the Maxwell and Pascal GPUs.
</p>
<div class="tablenoborder"><a name="maxwell-pascal__maxwell-pascal-instruction-set" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="maxwell-pascal__maxwell-pascal-instruction-set" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 6. Maxwell and Pascal Instruction Set</span></caption>
<thead class="thead" align="left">
<tr class="row" valign="top">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e3204" rowspan="1" colspan="1">Opcode</th>
<th class="entry" valign="top" width="66.66666666666666%" id="d54e3207" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e3204 d54e3207" rowspan="1"><strong class="ph b">Floating Point Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">FADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP32 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">FCHK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Single Precision FP Divide Range Check</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">FCMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP32 Compare to Zero and Select Source</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">FFMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP32 Fused Multiply and Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">FMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP32 Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">FMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP32 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">FSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP32 Compare And Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">FSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP32 Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">FSWZADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP32 Add used for FSWZ emulation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">MUFU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Multi Function Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">RRO</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Range Reduction Operator FP</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">DADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP64 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">DFMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP64 Fused Mutiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">DMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP64 Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">DMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP64 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">DSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP64 Compare And Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">DSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP64 Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">HADD2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP16 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">HFMA2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP16 Fused Mutiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">HMUL2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP16 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">HSET2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP16 Compare And Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">HSETP2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">FP16 Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e3204 d54e3207" rowspan="1"><strong class="ph b">Integer Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">BFE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Bit Field Extract</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">BFI</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Bit Field Insert</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">FLO</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Find Leading One</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">IADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">IADD3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">3-input Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">ICMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Integer Compare to Zero and Select Source</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">IMAD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Integer Multiply And Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">IMADSP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Extracted Integer Multiply And Add.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">IMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Integer Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">IMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Integer Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">ISCADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Scaled Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">ISET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Integer Compare And Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">ISETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Integer Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">LEA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Compute Effective Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">LOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">LOP3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">3-input Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">POPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Population count</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SHF</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Funnel Shift</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SHL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Shift Left</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SHR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Shift Right</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">XMAD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Integer Short Multiply Add</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e3204 d54e3207" rowspan="1"><strong class="ph b">Conversion Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">F2F</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Floating Point To Floating Point Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">F2I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Floating Point To Integer Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">I2F</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Integer To Floating Point Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">I2I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Integer To Integer Conversion</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e3204 d54e3207" rowspan="1"><strong class="ph b">Movement Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">MOV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">PRMT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Permute Register Pair</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SEL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Select Source with Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SHFL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Warp Wide Register Shuffle</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e3204 d54e3207" rowspan="1"><strong class="ph b">Predicate/CC Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">CSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Test Condition Code And Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">CSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Test Condition Code and Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">PSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Combine Predicates and Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">PSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Combine Predicates and Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">P2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Move Predicate Register To Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">R2P</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Move Register To Predicate/CC Register</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e3204 d54e3207" rowspan="1"><strong class="ph b">Texture Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">TEX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Texture Fetch</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">TLD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Texture Load</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">TLD4</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Texture Load 4</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">TXQ</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Texture Query</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">TEXS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Texture Fetch with scalar/non-vec4 source/destinations</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">TLD4S</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Texture Load 4 with scalar/non-vec4 source/destinations</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">TLDS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Texture Load with scalar/non-vec4 source/destinations</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e3204 d54e3207" rowspan="1"><strong class="ph b">Compute Load/Store Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">LD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Load from generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">LDC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Load Constant</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">LDG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Load from Global Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">LDL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Load within Local Memory Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">LDS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Local within Shared Memory Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">ST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Store to generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">STG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Store to global Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">STL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Store within Local or Shared Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">STS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Store within Local or Shared Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">ATOM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Atomic Operation on generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">ATOMS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Atomic Operation on Shared Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">RED</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Reduction Operation on generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">CCTL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Cache Control</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">CCTLL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Cache Control</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">MEMBAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Memory Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">CCTLT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Texture Cache Control</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e3204 d54e3207" rowspan="1"><strong class="ph b">Surface Memory Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SUATOM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Surface Reduction</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SULD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Surface Load</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SURED</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Atomic Reduction on surface memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SUST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Surface Store</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e3204 d54e3207" rowspan="1"><strong class="ph b">Control Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">BRA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Relative Branch</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">BRX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Relative Branch Indirect</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">JMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Absolute Jump</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">JMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Absolute Jump Indirect</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SSY</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Set Synchronization Point</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">SYNC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Converge threads after conditional branch</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">CAL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Relative Call</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">JCAL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Absolute Call</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">PRET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Pre-Return From Subroutine</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">RET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Return From Subroutine</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">BRK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Break</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">PBK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Pre-Break</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">CONT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Continue</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">PCNT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Pre-continue</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">EXIT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Exit Program</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">PEXIT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Pre-Exit</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">BPT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">BreakPoint/Trap</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e3204 d54e3207" rowspan="1"><strong class="ph b">Miscellaneous Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">NOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">No Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">CS2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Move Special Register to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">S2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Move Special Register to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">B2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Move Barrier To Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">BAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Barrier Synchronization</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">R2B</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Move Register to Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e3204" rowspan="1" colspan="1">VOTE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e3207" rowspan="1" colspan="1">Vote Across SIMD Thread Group</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="volta"><a name="volta" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#volta" name="volta" shape="rect">4.4. Volta Instruction Set</a></h3>
<div class="body conbody">
<div class="p">The Volta architecture (Compute Capability 7.x) has the following instruction
set format:
<pre class="pre screen" xml:space="preserve"><samp class="ph codeph">(instruction) (destination) (source1), (source2) ...</samp></pre></div>
<div class="p">Valid destination and source locations include:
<ul class="ul">
<li class="li">RX for registers</li>
<li class="li">SRX for special system-controlled registers</li>
<li class="li">PX for predicate registers</li>
<li class="li">c[X][Y] for constant memory</li>
</ul>
</div>
<p class="p"><a class="xref" href="index.html#volta__volta-instruction-set" shape="rect">Table 7</a> lists valid instructions for the Volta GPUs.
</p>
<div class="tablenoborder"><a name="volta__volta-instruction-set" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="volta__volta-instruction-set" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 7. Volta Instruction Set</span></caption>
<thead class="thead" align="left">
<tr class="row" valign="top">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e4333" rowspan="1" colspan="1">Opcode</th>
<th class="entry" valign="top" width="66.66666666666666%" id="d54e4336" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e4333 d54e4336" rowspan="1"><strong class="ph b">Floating Point Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FADD32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FCHK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Floating-point Range Check</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FFMA32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Fused Multiply and Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FFMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Fused Multiply and Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FMUL32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FSEL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Floating Point Select</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Compare And Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FSWZADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Swizzle Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">MUFU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP32 Multi Function Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">HADD2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP16 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">HADD2_32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP16 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">HFMA2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP16 Fused Mutiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">HFMA2_32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP16 Fused Mutiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">HMMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Matrix Multiply and Accumulate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">HMUL2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP16 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">HMUL2_32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP16 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">HSET2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP16 Compare And Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">HSETP2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP16 Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">DADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP64 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">DFMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP64 Fused Mutiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">DMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP64 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">DSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">FP64 Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e4333 d54e4336" rowspan="1"><strong class="ph b">Integer Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">BMSK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Bitfield Mask</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">BREV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Bit Reverse</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FLO</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Find Leading One</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IABS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Absolute Value</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IADD3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">3-input Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IADD32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IDP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Dot Product and Accumulate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IDP4A</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Dot Product and Accumulate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IMAD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Multiply And Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IMMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Matrix Multiply and Accumulate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">IMUL32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">ISCADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Scaled Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">ISCADD32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Scaled Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">ISETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">LEA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">LOAD Effective Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">LOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">LOP3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">LOP32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">POPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Population count</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SHF</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Funnel Shift</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SHL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Shift Left</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SHR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Shift Right</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">VABSDIFF</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Absolute Difference</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">VABSDIFF4</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Absolute Difference</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e4333 d54e4336" rowspan="1"><strong class="ph b">Conversion Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">F2F</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Floating Point To Floating Point Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">F2I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Floating Point To Integer Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">I2F</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer To Floating Point Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">I2I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer To Integer Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">I2IP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Integer To Integer Conversion and Packing</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">FRND</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Round To Integer</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e4333 d54e4336" rowspan="1"><strong class="ph b">Movement Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">MOV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">MOV32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">PRMT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Permute Register Pair</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SEL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Select Source with Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SGXT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Sign Extend</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SHFL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Warp Wide Register Shuffle</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e4333 d54e4336" rowspan="1"><strong class="ph b">Predicate Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">PLOP3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Predicate Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">PSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Combine Predicates and Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">P2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Move Predicate Register To Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">R2P</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Move Register To Predicate Register</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e4333 d54e4336" rowspan="1"><strong class="ph b">Load/Store Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">LD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Load from generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">LDC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Load Constant</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">LDG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Load from Global Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">LDL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Load within Local Memory Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">LDS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Load within Shared Memory Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">ST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Store to Generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">STG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Store to Global Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">STL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Store within Local or Shared Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">STS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Store within Local or Shared Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">MATCH</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Match Register Values Across Thread Group</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">QSPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Query Space</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">ATOM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Atomic Operation on Generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">ATOMS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Atomic Operation on Shared Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">ATOMG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Atomic Operation on Global Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">RED</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Reduction Operation on Generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">CCTL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Cache Control</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">CCTLL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Cache Control</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">ERRBAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Error Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">MEMBAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Memory Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">CCTLT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Texture Cache Control</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e4333 d54e4336" rowspan="1"><strong class="ph b">Texture Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">TEX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Texture Fetch</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">TLD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Texture Load</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">TLD4</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Texture Load 4</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">TMML</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Texture MipMap Level</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">TXD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Texture Fetch With Derivatives</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">TXQ</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Texture Query</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e4333 d54e4336" rowspan="1"><strong class="ph b">Surface Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SUATOM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Surface Reduction</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SULD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Surface Load</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SURED</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Atomic Reduction on surface memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SUST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Surface Store</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e4333 d54e4336" rowspan="1"><strong class="ph b">Control Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">BMOV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Move Convergence Barrier State</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">BPT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">BreakPoint/Trap</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">BRA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Relative Branch</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">BREAK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Break out of the Specified Convergence Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">BRX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Relative Branch Indirect</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">BSSY</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Barrier Set Convergence Synchronization Point</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">BSYNC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Synchronize Threads on a Convergence Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">CALL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Call Function</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">EXIT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Exit Program</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">JMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Absolute Jump</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">JMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Absolute Jump Indirect</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">KILL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Kill Thread</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">NANOSLEEP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Suspend Execution</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">RET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Return From Subroutine</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">RPCMOV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">PC Register Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">RTT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Return From Trap</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">WARPSYNC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Synchronize Threads in Warp</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">YIELD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Yield Control</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e4333 d54e4336" rowspan="1"><strong class="ph b">Miscellaneous Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">B2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Move Barrier To Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">BAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Barrier Synchronization</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">CS2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Move Special Register to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">CSMTEST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Clip State Machine Test and Update</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">DEPBAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Dependency Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">GETLMEMBASE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Get Local Memory Base Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">LEPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Load Effective PC</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">NOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">No Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">PMTRIG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Performance Monitor Trigger</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">R2B</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Move Register to Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">S2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Move Special Register to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SETCTAID</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Set CTA ID</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">SETLMEMBASE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Set Local Memory Base Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">VOTE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Vote Across SIMD Thread Group</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e4333" rowspan="1" colspan="1">VOTE_VTG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e4336" rowspan="1" colspan="1">Clip State Machine Test and Update</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="turing"><a name="turing" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#turing" name="turing" shape="rect">4.5. Turing Instruction Set</a></h3>
<div class="body conbody">
<div class="p">The Turing architecture (Compute Capability <span class="ph"> and 7.5</span>) has the following instruction
set format:
<pre class="pre screen" xml:space="preserve"><samp class="ph codeph">(instruction) (destination) (source1), (source2) ...</samp></pre></div>
<div class="p">Valid destination and source locations include:
<ul class="ul">
<li class="li">RX for registers</li>
<li class="li">URX for uniform registers</li>
<li class="li">SRX for special system-controlled registers</li>
<li class="li">PX for predicate registers</li>
<li class="li">c[X][Y] for constant memory</li>
</ul>
</div>
<p class="p"><a class="xref" href="index.html#turing__turing-instruction-set" shape="rect">Table 8</a> lists valid instructions for the Turing GPUs.
</p>
<div class="tablenoborder"><a name="turing__turing-instruction-set" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="turing__turing-instruction-set" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 8. Turing Instruction Set</span></caption>
<thead class="thead" align="left">
<tr class="row" valign="top">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e5687" rowspan="1" colspan="1">Opcode</th>
<th class="entry" valign="top" width="66.66666666666666%" id="d54e5690" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Floating Point Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FADD32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FCHK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Floating-point Range Check</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FFMA32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Fused Multiply and Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FFMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Fused Multiply and Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FMUL32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FSEL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Floating Point Select</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FSET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Compare And Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FSWZADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Swizzle Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">MUFU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP32 Multi Function Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">HADD2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP16 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">HADD2_32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP16 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">HFMA2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP16 Fused Mutiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">HFMA2_32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP16 Fused Mutiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">HMMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Matrix Multiply and Accumulate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">HMUL2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP16 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">HMUL2_32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP16 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">HSET2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP16 Compare And Set</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">HSETP2</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP16 Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">DADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP64 Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">DFMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP64 Fused Mutiply Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">DMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP64 Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">DSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">FP64 Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Integer Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BMMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Bit Matrix Multiply and Accumulate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BMSK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Bitfield Mask</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BREV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Bit Reverse</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FLO</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Find Leading One</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IABS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Absolute Value</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IADD3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">3-input Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IADD32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IDP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Dot Product and Accumulate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IDP4A</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Dot Product and Accumulate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IMAD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Multiply And Add</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IMMA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Matrix Multiply and Accumulate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IMNMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Minimum/Maximum</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IMUL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">IMUL32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Multiply</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ISCADD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Scaled Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ISCADD32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Scaled Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ISETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Compare And Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LEA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">LOAD Effective Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LOP3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LOP32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">POPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Population count</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SHF</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Funnel Shift</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SHL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Shift Left</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SHR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Shift Right</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">VABSDIFF</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Absolute Difference</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">VABSDIFF4</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Absolute Difference</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Conversion Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">F2F</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Floating Point To Floating Point Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">F2I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Floating Point To Integer Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">I2F</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer To Floating Point Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">I2I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer To Integer Conversion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">I2IP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer To Integer Conversion and Packing</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">FRND</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Round To Integer</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Movement Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">MOV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">MOV32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">MOVM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move Matrix with Transposition or Expansion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">PRMT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Permute Register Pair</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SEL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Select Source with Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SGXT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Sign Extend</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SHFL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Warp Wide Register Shuffle</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Predicate Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">PLOP3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Predicate Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">PSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Combine Predicates and Set Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">P2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move Predicate Register To Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">R2P</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move Register To Predicate Register</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Load/Store Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Load from generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LDC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Load Constant</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LDG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Load from Global Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LDL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Load within Local Memory Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LDS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Load within Shared Memory Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LDSM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Load Matrix from Shared Memory with Element Size Expansion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Store to Generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">STG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Store to Global Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">STL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Store within Local or Shared Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">STS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Store within Local or Shared Window</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">MATCH</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Match Register Values Across Thread Group</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">QSPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Query Space</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ATOM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Atomic Operation on Generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ATOMS</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Atomic Operation on Shared Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ATOMG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Atomic Operation on Global Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">RED</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Reduction Operation on Generic Memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">CCTL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Cache Control</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">CCTLL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Cache Control</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ERRBAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Error Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">MEMBAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Memory Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">CCTLT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Texture Cache Control</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Uniform Datapath Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">R2UR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move from Vector Register to a Uniform Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">S2UR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move Special Register to Uniform Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UBMSK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Bitfield Mask</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UBREV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Bit Reverse</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UCLEA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Load Effective Address for a Constant</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UFLO</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Find Leading One</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UIADD3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UIADD3.64</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Integer Addition</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UIMAD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Integer Multiplication</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UISETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Integer Compare and Set Uniform Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ULDC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Load from Constant Memory into a Uniform Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ULEA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Load Effective Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ULOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ULOP3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">ULOP32I</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UMOV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UP2UR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Predicate to Uniform Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UPLOP3</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Predicate Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UPOPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Population Count</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UPRMT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Byte Permute</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UPSETP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Predicate Logic Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">UR2UP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Register to Uniform Predicate</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">USEL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Select</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">USGXT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Sign Extend</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">USHF</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Funnel Shift</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">USHL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Left Shift</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">USHR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Uniform Right Shift</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">VOTEU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Voting across SIMD Thread Group with Results in Uniform Destination</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Texture Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">TEX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Texture Fetch</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">TLD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Texture Load</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">TLD4</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Texture Load 4</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">TMML</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Texture MipMap Level</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">TXD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Texture Fetch With Derivatives</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">TXQ</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Texture Query</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Surface Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SUATOM</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Surface Reduction</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SULD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Surface Load</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SURED</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Atomic Reduction on surface memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SUST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Surface Store</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Control Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BMOV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move Convergence Barrier State</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BPT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">BreakPoint/Trap</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BRA</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Relative Branch</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BREAK</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Break out of the Specified Convergence Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BRX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Relative Branch Indirect</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BRXU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Relative Branch with Uniform Register Based Offset</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BSSY</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Barrier Set Convergence Synchronization Point</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BSYNC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Synchronize Threads on a Convergence Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">CALL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Call Function</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">EXIT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Exit Program</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">JMP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Absolute Jump</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">JMX</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Absolute Jump Indirect</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">JMXU</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Absolute Jump with Uniform Register Based Offset</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">KILL</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Kill Thread</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">NANOSLEEP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Suspend Execution</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">RET</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Return From Subroutine</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">RPCMOV</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">PC Register Move</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">RTT</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Return From Trap</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">WARPSYNC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Synchronize Threads in Warp</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">YIELD</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Yield Control</td>
</tr>
<tr class="row">
<td class="entry" colspan="2" valign="top" headers="d54e5687 d54e5690" rowspan="1"><strong class="ph b">Miscellaneous Instructions</strong></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">B2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move Barrier To Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">BAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Barrier Synchronization</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">CS2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move Special Register to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">CSMTEST</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Clip State Machine Test and Update</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">DEPBAR</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Dependency Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">GETLMEMBASE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Get Local Memory Base Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">LEPC</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Load Effective PC</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">NOP</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">No Operation</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">PMTRIG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Performance Monitor Trigger</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">R2B</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move Register to Barrier</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">S2R</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Move Special Register to Register</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SETCTAID</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Set CTA ID</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">SETLMEMBASE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Set Local Memory Base Address</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">VOTE</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Vote Across SIMD Thread Group</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e5687" rowspan="1" colspan="1">VOTE_VTG</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e5690" rowspan="1" colspan="1">Clip State Machine Test and Update</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="nvprune"><a name="nvprune" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#nvprune" name="nvprune" shape="rect">5. nvprune</a></h2>
<div class="body conbody">
<p class="p"><samp class="ph codeph">nvprune</samp> prunes host object files and libraries to only contain device code for the specified targets.
</p>
</div>
<div class="topic concept nested1" id="nvprune-usage"><a name="nvprune-usage" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#nvprune-usage" name="nvprune-usage" shape="rect">5.1. Usage</a></h3>
<div class="body conbody">
<p class="p"><samp class="ph codeph">nvprune</samp> accepts a single input file each time it's run, emitting a new output file. The basic usage is as following:
</p><pre class="pre screen" xml:space="preserve">nvprune [options] -o <outfile> <infile></pre><p class="p">The input file must be either a relocatable host object or static library (not a host executable), and the output file will
be the same format.
</p>
<p class="p">Either the --arch or --generate-code option must be used to specify the target(s) to keep. All other device code is discarded
from the file. The targets can be either a sm_NN arch (cubin) or compute_NN arch (ptx).
</p>
<p class="p">For example, the following will prune libcublas_static.a to only contain sm_35 cubin rather than all the targets which normally
exist:
</p><pre class="pre screen" xml:space="preserve">nvprune -arch sm_35 libcublas_static.a -o libcublas_static35.a</pre><p class="p">Note that this means that libcublas_static35.a will not run on any other architecture, so should only be used when you are
building for a single architecture.
</p>
</div>
</div>
<div class="topic concept nested1" id="nvprune-options"><a name="nvprune-options" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#nvprune-options" name="nvprune-options" shape="rect">5.2. Command-line Options</a></h3>
<div class="body conbody">
<p class="p"><a class="xref" href="index.html#nvprune-options__nvprune-options-table" shape="rect">Table 9</a> contains supported command-line options of <samp class="ph codeph">nvprune</samp>, along
with a description of what each option does. Each option has a long name and
a short name, which can be used interchangeably.
</p>
<div class="tablenoborder"><a name="nvprune-options__nvprune-options-table" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="nvprune-options__nvprune-options-table" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 9. <samp class="ph codeph">nvprune</samp> Command-line Options</span></caption>
<thead class="thead" align="left">
<tr class="row" valign="top">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e7375" rowspan="1" colspan="1">Option (long)</th>
<th class="entry" valign="top" width="13.333333333333334%" id="d54e7378" rowspan="1" colspan="1">Option (short)</th>
<th class="entry" valign="top" width="53.333333333333336%" id="d54e7381" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e7375" rowspan="1" colspan="1"><samp class="ph codeph">--arch <gpu architecture name>,...</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e7378" rowspan="1" colspan="1"><samp class="ph codeph">-arch</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e7381" rowspan="1" colspan="1">Specify the name of the NVIDIA GPU architecture
which will remain in the object or library.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e7375" rowspan="1" colspan="1"><samp class="ph codeph">--generate-code</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e7378" rowspan="1" colspan="1"><samp class="ph codeph">-gencode</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e7381" rowspan="1" colspan="1">This option is same format as nvcc --generate-code
option, and provides a way to specify multiple
architectures which should remain in the object or
library. Only the 'code' values are used as targets
to match. Allowed keywords for this option:
'arch','code'.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e7375" rowspan="1" colspan="1"><samp class="ph codeph">--output-file</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e7378" rowspan="1" colspan="1"><samp class="ph codeph">-o</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e7381" rowspan="1" colspan="1">Specify name and location of the output file.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e7375" rowspan="1" colspan="1"><samp class="ph codeph">--help</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e7378" rowspan="1" colspan="1"><samp class="ph codeph">-h</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e7381" rowspan="1" colspan="1">Print this help information on this tool.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e7375" rowspan="1" colspan="1"><samp class="ph codeph">--options-file <file>,...</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e7378" rowspan="1" colspan="1"><samp class="ph codeph">-optf</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e7381" rowspan="1" colspan="1">Include command line options from specified file.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e7375" rowspan="1" colspan="1"><samp class="ph codeph">--version</samp></td>
<td class="entry" valign="top" width="13.333333333333334%" headers="d54e7378" rowspan="1" colspan="1"><samp class="ph codeph">-V</samp></td>
<td class="entry" valign="top" width="53.333333333333336%" headers="d54e7381" rowspan="1" colspan="1">Print version information on this tool.</td>
</tr>
</tbody>
</table>
</div>
</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-present"><a name="copyright-present" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#copyright-present" name="copyright-present" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Copyright</h3>
<p class="p">© <span class="ph">2019</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><script type="text/javascript">_satellite.pageBottom();</script></body>
</html>
distrib
>
Mageia
>
7
>
x86_64
>
media
>
nonfree-updates
>
by-pkgid
>
b86a85131cc739c1c53d0b55840a4328
>
files
>
1309
