118 lines
16 KiB
HTML
118 lines
16 KiB
HTML
<!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">
|
|
<head>
|
|
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
|
|
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
|
|
<meta name="generator" content="Doxygen 1.8.11"/>
|
|
<title>CUTLASS: Main Page</title>
|
|
<link href="tabs.css" rel="stylesheet" type="text/css"/>
|
|
<script type="text/javascript" src="jquery.js"></script>
|
|
<script type="text/javascript" src="dynsections.js"></script>
|
|
<link href="search/search.css" rel="stylesheet" type="text/css"/>
|
|
<script type="text/javascript" src="search/searchdata.js"></script>
|
|
<script type="text/javascript" src="search/search.js"></script>
|
|
<script type="text/javascript">
|
|
$(document).ready(function() { init_search(); });
|
|
</script>
|
|
<script type="text/x-mathjax-config">
|
|
MathJax.Hub.Config({
|
|
extensions: ["tex2jax.js"],
|
|
jax: ["input/TeX","output/HTML-CSS"],
|
|
});
|
|
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
|
|
<link href="doxygen.css" rel="stylesheet" type="text/css" />
|
|
</head>
|
|
<body>
|
|
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
|
|
<div id="titlearea">
|
|
<table cellspacing="0" cellpadding="0">
|
|
<tbody>
|
|
<tr style="height: 56px;">
|
|
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
|
|
<td id="projectalign" style="padding-left: 0.5em;">
|
|
<div id="projectname">CUTLASS
|
|
</div>
|
|
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
|
|
</td>
|
|
</tr>
|
|
</tbody>
|
|
</table>
|
|
</div>
|
|
<!-- end header part -->
|
|
<!-- Generated by Doxygen 1.8.11 -->
|
|
<script type="text/javascript">
|
|
var searchBox = new SearchBox("searchBox", "search",false,'Search');
|
|
</script>
|
|
<div id="navrow1" class="tabs">
|
|
<ul class="tablist">
|
|
<li class="current"><a href="index.html"><span>Main Page</span></a></li>
|
|
<li><a href="modules.html"><span>Modules</span></a></li>
|
|
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
|
|
<li><a href="annotated.html"><span>Classes</span></a></li>
|
|
<li><a href="files.html"><span>Files</span></a></li>
|
|
<li>
|
|
<div id="MSearchBox" class="MSearchBoxInactive">
|
|
<span class="left">
|
|
<img id="MSearchSelect" src="search/mag_sel.png"
|
|
onmouseover="return searchBox.OnSearchSelectShow()"
|
|
onmouseout="return searchBox.OnSearchSelectHide()"
|
|
alt=""/>
|
|
<input type="text" id="MSearchField" value="Search" accesskey="S"
|
|
onfocus="searchBox.OnSearchFieldFocus(true)"
|
|
onblur="searchBox.OnSearchFieldFocus(false)"
|
|
onkeyup="searchBox.OnSearchFieldChange(event)"/>
|
|
</span><span class="right">
|
|
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
|
|
</span>
|
|
</div>
|
|
</li>
|
|
</ul>
|
|
</div>
|
|
</div><!-- top -->
|
|
<!-- window showing the filter options -->
|
|
<div id="MSearchSelectWindow"
|
|
onmouseover="return searchBox.OnSearchSelectShow()"
|
|
onmouseout="return searchBox.OnSearchSelectHide()"
|
|
onkeydown="return searchBox.OnSearchSelectKey(event)">
|
|
</div>
|
|
|
|
<!-- iframe showing the search results (closed by default) -->
|
|
<div id="MSearchResultsWindow">
|
|
<iframe src="javascript:void(0)" frameborder="0"
|
|
name="MSearchResults" id="MSearchResults">
|
|
</iframe>
|
|
</div>
|
|
|
|
<div class="header">
|
|
<div class="headertitle">
|
|
<div class="title">CUTLASS Documentation</div> </div>
|
|
</div><!--header-->
|
|
<div class="contents">
|
|
<div class="textblock"><div class="image">
|
|
<img src="/media/images/gemm-hierarchy-with-epilogue-no-labels.png" alt="ALT" title="CUTLASS 2.0"/>
|
|
</div>
|
|
<h1>CUTLASS 2.0</h1>
|
|
<p><em>CUTLASS 2.0 - November 2019</em></p>
|
|
<p>CUTLASS is a collection of CUDA C++ template abstractions for implementing high-performance matrix-multiplication (GEMM) at all levels and scales within CUDA. It incorporates strategies for hierarchical decomposition and data movement similar to those used to implement cuBLAS. CUTLASS decomposes these "moving parts" into reusable, modular software components abstracted by C++ template classes. These thread-wide, warp-wide, block-wide, and device-wide primitives can be specialized and tuned via custom tiling sizes, data types, and other algorithmic policy. The resulting flexibility simplifies their use as building blocks within custom kernels and applications.</p>
|
|
<p>To support a wide variety of applications, CUTLASS provides extensive support for mixed-precision computations, providing specialized data-movement and multiply-accumulate abstractions for 8-bit integer, half-precision floating point (FP16), single-precision floating point (FP32), and double-precision floating point (FP64) types. Furthermore, CUTLASS demonstrates warp-synchronous matrix multiply operations for targeting the programmable, high-throughput <em>Tensor Cores</em> implemented by NVIDIA's Volta and Turing architectures.</p>
|
|
<h1>What's New in CUTLASS 2.0</h1>
|
|
<p>CUTLASS 2.0 is a substantial refactoring from the previous version, intended to offer:</p>
|
|
<ul>
|
|
<li>Better performance over 1.x, particularly for kernels targeting Turing Tensor Cores</li>
|
|
<li>Robust and durable templates that reliably span the design space</li>
|
|
<li>Encapsulated functionality that may be reusable in other contexts</li>
|
|
</ul>
|
|
<h1>Example CUTLASS GEMM</h1>
|
|
<p>The following illustrates an example function that defines a CUTLASS GEMM kernel with single-precision inputs and outputs. This is an exercpt from the CUTLASS SDK <a href="https://github.com/NVIDIA/cutlass/tree/master/examples/00_basic_gemm/basic_gemm.cu">basic_gemm example</a>.</p>
|
|
<div class="fragment"><div class="line"><span class="comment">//</span></div><div class="line"><span class="comment">// CUTLASS includes needed for single-precision GEMM kernel</span></div><div class="line"><span class="comment">//</span></div><div class="line"></div><div class="line"><span class="comment">// Defines cutlass::gemm::device::Gemm, the generic Gemm computation template class.</span></div><div class="line"></div><div class="line"><span class="preprocessor">#include <<a class="code" href="include_2cutlass_2gemm_2device_2gemm_8h.html">cutlass/gemm/device/gemm.h</a>></span></div><div class="line"></div><div class="line">cudaError_t cutlass_sgemm_nn(</div><div class="line"> <span class="keywordtype">int</span> M,</div><div class="line"> <span class="keywordtype">int</span> N,</div><div class="line"> <span class="keywordtype">int</span> K,</div><div class="line"> <span class="keywordtype">float</span> alpha,</div><div class="line"> <span class="keywordtype">float</span> <span class="keyword">const</span> *A,</div><div class="line"> <span class="keywordtype">int</span> lda,</div><div class="line"> <span class="keywordtype">float</span> <span class="keyword">const</span> *B,</div><div class="line"> <span class="keywordtype">int</span> ldb,</div><div class="line"> <span class="keywordtype">float</span> beta,</div><div class="line"> <span class="keywordtype">float</span> *C,</div><div class="line"> <span class="keywordtype">int</span> ldc) {</div><div class="line"></div><div class="line"> <span class="comment">// Define type definition for single-precision CUTLASS GEMM with column-major</span></div><div class="line"> <span class="comment">// input matrices and 128x128x8 threadblock tile size (chosen by default).</span></div><div class="line"> <span class="comment">//</span></div><div class="line"> <span class="comment">// To keep the interface manageable, several helpers are defined for plausible compositions</span></div><div class="line"> <span class="comment">// including the following example for single-precision GEMM. Typical values are used as</span></div><div class="line"> <span class="comment">// default template arguments. See `cutlass/gemm/device/default_gemm_configuration.h` for more details.</span></div><div class="line"> <span class="comment">//</span></div><div class="line"> <span class="comment">// To view the full gemm device API interface, see `cutlass/gemm/device/gemm.h`</span></div><div class="line"></div><div class="line"> <span class="keyword">using</span> ColumnMajor = <a class="code" href="classcutlass_1_1layout_1_1ColumnMajor.html">cutlass::layout::ColumnMajor</a>;</div><div class="line"></div><div class="line"> <span class="keyword">using</span> CutlassGemm = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html">cutlass::gemm::device::Gemm</a><float, <span class="comment">// Data-type of A matrix</span></div><div class="line"> ColumnMajor, <span class="comment">// Layout of A matrix</span></div><div class="line"> float, <span class="comment">// Data-type of B matrix</span></div><div class="line"> ColumnMajor, <span class="comment">// Layout of B matrix</span></div><div class="line"> float, <span class="comment">// Data-type of C matrix</span></div><div class="line"> ColumnMajor>; <span class="comment">// Layout of C matrix</span></div><div class="line"></div><div class="line"> <span class="comment">// Define a CUTLASS GEMM type</span></div><div class="line"></div><div class="line"> CutlassGemm gemm_operator;</div><div class="line"></div><div class="line"> <span class="comment">// Construct the CUTLASS GEMM arguments object.</span></div><div class="line"> <span class="comment">//</span></div><div class="line"> <span class="comment">// One of CUTLASS's design patterns is to define gemm argument objects that are constructible</span></div><div class="line"> <span class="comment">// in host code and passed to kernels by value. These may include pointers, strides, scalars,</span></div><div class="line"> <span class="comment">// and other arguments needed by Gemm and its components.</span></div><div class="line"> <span class="comment">//</span></div><div class="line"> <span class="comment">// The benefits of this pattern are (1.) a structured, composable strategy for passing host-constructible</span></div><div class="line"> <span class="comment">// arguments to kernels and (2.) minimized initialization overhead on kernel entry.</span></div><div class="line"> <span class="comment">//</span></div><div class="line"></div><div class="line"> CutlassGemm::Arguments args({M , N, K}, <span class="comment">// Gemm Problem dimensions</span></div><div class="line"> {A, lda}, <span class="comment">// Tensor-ref for source matrix A</span></div><div class="line"> {B, ldb}, <span class="comment">// Tensor-ref for source matrix B</span></div><div class="line"> {C, ldc}, <span class="comment">// Tensor-ref for source matrix C</span></div><div class="line"> {C, ldc}, <span class="comment">// Tensor-ref for destination matrix D (may be different memory than source C matrix)</span></div><div class="line"> {alpha, beta}); <span class="comment">// Scalars used in the Epilogue</span></div><div class="line"></div><div class="line"> <span class="comment">//</span></div><div class="line"> <span class="comment">// Launch the CUTLASS GEMM kernel.</span></div><div class="line"> <span class="comment">//</span></div><div class="line"></div><div class="line"> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">cutlass::Status</a> status = gemm_operator(args);</div><div class="line"></div><div class="line"> <span class="comment">//</span></div><div class="line"> <span class="comment">// Return a cudaError_t if the CUTLASS GEMM operator returned an error code.</span></div><div class="line"> <span class="comment">//</span></div><div class="line"></div><div class="line"> <span class="keywordflow">if</span> (status != <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">cutlass::Status::kSuccess</a>) {</div><div class="line"> <span class="keywordflow">return</span> cudaErrorUnknown;</div><div class="line"> }</div><div class="line"></div><div class="line"> <span class="comment">// Return success, if no errors were encountered.</span></div><div class="line"></div><div class="line"> <span class="keywordflow">return</span> cudaSuccess;</div><div class="line">}</div></div><!-- fragment --><h1>Copyright</h1>
|
|
<p>Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.</p>
|
|
<div class="fragment"><div class="line"><a name="l00001"></a><span class="lineno"> 1</span> Redistribution and use in source and binary forms, with or without modification, are permitted</div><div class="line"><a name="l00002"></a><span class="lineno"> 2</span> provided that the following conditions are met:</div><div class="line"><a name="l00003"></a><span class="lineno"> 3</span>  * Redistributions of source code must retain the above copyright notice, this list of</div><div class="line"><a name="l00004"></a><span class="lineno"> 4</span>  conditions and the following disclaimer.</div><div class="line"><a name="l00005"></a><span class="lineno"> 5</span>  * Redistributions in binary form must reproduce the above copyright notice, this list of</div><div class="line"><a name="l00006"></a><span class="lineno"> 6</span>  conditions and the following disclaimer in the documentation and/or other materials</div><div class="line"><a name="l00007"></a><span class="lineno"> 7</span>  provided with the distribution.</div><div class="line"><a name="l00008"></a><span class="lineno"> 8</span>  * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used</div><div class="line"><a name="l00009"></a><span class="lineno"> 9</span>  to endorse or promote products derived from this software without specific prior written</div><div class="line"><a name="l00010"></a><span class="lineno"> 10</span>  permission.</div><div class="line"><a name="l00011"></a><span class="lineno"> 11</span> </div><div class="line"><a name="l00012"></a><span class="lineno"> 12</span> THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR</div><div class="line"><a name="l00013"></a><span class="lineno"> 13</span> IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND</div><div class="line"><a name="l00014"></a><span class="lineno"> 14</span> FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE</div><div class="line"><a name="l00015"></a><span class="lineno"> 15</span> FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,</div><div class="line"><a name="l00016"></a><span class="lineno"> 16</span> BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;</div><div class="line"><a name="l00017"></a><span class="lineno"> 17</span> OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,</div><div class="line"><a name="l00018"></a><span class="lineno"> 18</span> STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE</div><div class="line"><a name="l00019"></a><span class="lineno"> 19</span> OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.</div></div><!-- fragment --> </div></div><!-- contents -->
|
|
<!-- start footer part -->
|
|
<hr class="footer"/><address class="footer"><small>
|
|
Generated by  <a href="http://www.doxygen.org/index.html">
|
|
<img class="footer" src="doxygen.png" alt="doxygen"/>
|
|
</a> 1.8.11
|
|
</small></address>
|
|
</body>
|
|
</html>
|