dnrops
/
cutlass
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity
cutlass/docs/structcutlass_1_1layout_1_1...

572 lines
35 KiB
HTML
Raw Permalink Normal View History

CUTLASS 2.0 (#62) CUTLASS 2.0 Substantially refactored for - Better performance, particularly for native Turing Tensor Cores - Robust and durable templates spanning the design space - Encapsulated functionality embodying modern C++11 programming techniques - Optimized containers and data types for efficient, generic, portable device code Updates to: - Quick start guide - Documentation - Utilities - CUTLASS Profiler Native Turing Tensor Cores - Efficient GEMM kernels targeting Turing Tensor Cores - Mixed-precision floating point, 8-bit integer, 4-bit integer, and binarized operands Coverage of existing CUTLASS functionality: - GEMM kernels targeting CUDA and Tensor Cores in NVIDIA GPUs - Volta Tensor Cores through native mma.sync and through WMMA API - Optimizations such as parallel reductions, threadblock rasterization, and intra-threadblock reductions - Batched GEMM operations - Complex-valued GEMMs Note: this commit and all that follow require a host compiler supporting C++11 or greater.
2019-11-20 08:55:34 +08:00
<!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: cutlass::layout::ColumnMajorBlockLinear&lt; BlockRows, BlockColumns &gt; Struct Template Reference</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><a href="index.html"><span>Main&#160;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 class="current"><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 id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="annotated.html"><span>Class&#160;List</span></a></li>
<li><a href="classes.html"><span>Class&#160;Index</span></a></li>
<li><a href="inherits.html"><span>Class&#160;Hierarchy</span></a></li>
<li><a href="functions.html"><span>Class&#160;Members</span></a></li>
</ul>
</div>
<!-- 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 id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="namespacecutlass.html">cutlass</a></li><li class="navelem"><a class="el" href="namespacecutlass_1_1layout.html">layout</a></li><li class="navelem"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">ColumnMajorBlockLinear</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#pub-types">Public Types</a> &#124;
<a href="#pub-methods">Public Member Functions</a> &#124;
<a href="#pub-static-methods">Static Public Member Functions</a> &#124;
<a href="#pub-static-attribs">Static Public Attributes</a> &#124;
<a href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear-members.html">List of all members</a> </div>
<div class="headertitle">
<div class="title">cutlass::layout::ColumnMajorBlockLinear&lt; BlockRows, BlockColumns &gt; Struct Template Reference</div> </div>
</div><!--header-->
<div class="contents">
<p><code>#include &lt;<a class="el" href="layout_2matrix_8h_source.html">matrix.h</a>&gt;</code></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="pub-types"></a>
Public Types</h2></td></tr>
<tr class="memitem:a0ed618d4bb2250c01db5bf9ca61e6839"><td class="memItemLeft" align="right" valign="top">using&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0ed618d4bb2250c01db5bf9ca61e6839">Index</a> = int32_t</td></tr>
<tr class="memdesc:a0ed618d4bb2250c01db5bf9ca61e6839"><td class="mdescLeft">&#160;</td><td class="mdescRight">Index type used for coordinates. <a href="#a0ed618d4bb2250c01db5bf9ca61e6839">More...</a><br /></td></tr>
<tr class="separator:a0ed618d4bb2250c01db5bf9ca61e6839"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a3560440cff27c7723c731efd75b3e838"><td class="memItemLeft" align="right" valign="top">using&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a3560440cff27c7723c731efd75b3e838">LongIndex</a> = int64_t</td></tr>
<tr class="memdesc:a3560440cff27c7723c731efd75b3e838"><td class="mdescLeft">&#160;</td><td class="mdescRight">Long index type used for offsets. <a href="#a3560440cff27c7723c731efd75b3e838">More...</a><br /></td></tr>
<tr class="separator:a3560440cff27c7723c731efd75b3e838"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ae097874037352b7bf9b7640a9830a3aa"><td class="memItemLeft" align="right" valign="top">using&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#ae097874037352b7bf9b7640a9830a3aa">TensorCoord</a> = <a class="el" href="structcutlass_1_1MatrixCoord.html">MatrixCoord</a></td></tr>
<tr class="memdesc:ae097874037352b7bf9b7640a9830a3aa"><td class="mdescLeft">&#160;</td><td class="mdescRight">Logical coordinate. <a href="#ae097874037352b7bf9b7640a9830a3aa">More...</a><br /></td></tr>
<tr class="separator:ae097874037352b7bf9b7640a9830a3aa"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a64a52f31aa332650601a37643adea428"><td class="memItemLeft" align="right" valign="top">using&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a64a52f31aa332650601a37643adea428">Stride</a> = <a class="el" href="structcutlass_1_1Coord.html">Coord</a>&lt; <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0732c2b94df9a0c10453cba940684cca">kStrideRank</a>, <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0ed618d4bb2250c01db5bf9ca61e6839">Index</a> &gt;</td></tr>
<tr class="memdesc:a64a52f31aa332650601a37643adea428"><td class="mdescLeft">&#160;</td><td class="mdescRight">Stride vector. <a href="#a64a52f31aa332650601a37643adea428">More...</a><br /></td></tr>
<tr class="separator:a64a52f31aa332650601a37643adea428"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="pub-methods"></a>
Public Member Functions</h2></td></tr>
<tr class="memitem:a740b33237fd49dee4692946811e203ba"><td class="memItemLeft" align="right" valign="top"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a>&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a740b33237fd49dee4692946811e203ba">ColumnMajorBlockLinear</a> (<a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0ed618d4bb2250c01db5bf9ca61e6839">Index</a> ldm=0)</td></tr>
<tr class="memdesc:a740b33237fd49dee4692946811e203ba"><td class="mdescLeft">&#160;</td><td class="mdescRight">Ctor. <a href="#a740b33237fd49dee4692946811e203ba">More...</a><br /></td></tr>
<tr class="separator:a740b33237fd49dee4692946811e203ba"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a849bbda685abc0b855352c429a8f7e08"><td class="memItemLeft" align="right" valign="top"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a3560440cff27c7723c731efd75b3e838">LongIndex</a>&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a849bbda685abc0b855352c429a8f7e08">operator()</a> (<a class="el" href="structcutlass_1_1MatrixCoord.html">MatrixCoord</a> const &amp;coord) const </td></tr>
<tr class="separator:a849bbda685abc0b855352c429a8f7e08"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:aaf2b7966340292ba65e3cb6163b4348a"><td class="memItemLeft" align="right" valign="top"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1MatrixCoord.html">MatrixCoord</a>&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#aaf2b7966340292ba65e3cb6163b4348a">inverse</a> (<a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a3560440cff27c7723c731efd75b3e838">LongIndex</a> offset) const </td></tr>
<tr class="memdesc:aaf2b7966340292ba65e3cb6163b4348a"><td class="mdescLeft">&#160;</td><td class="mdescRight">Inverse of layout function, mapping linear offset to logical coordinate. <a href="#aaf2b7966340292ba65e3cb6163b4348a">More...</a><br /></td></tr>
<tr class="separator:aaf2b7966340292ba65e3cb6163b4348a"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a7cd57a34850bcf5badbbd760c89aed18"><td class="memItemLeft" align="right" valign="top"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a64a52f31aa332650601a37643adea428">Stride</a>&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a7cd57a34850bcf5badbbd760c89aed18">stride</a> () const </td></tr>
<tr class="memdesc:a7cd57a34850bcf5badbbd760c89aed18"><td class="mdescLeft">&#160;</td><td class="mdescRight">Returns the stride of the layout. <a href="#a7cd57a34850bcf5badbbd760c89aed18">More...</a><br /></td></tr>
<tr class="separator:a7cd57a34850bcf5badbbd760c89aed18"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a8c099a9478d79dbefb04b6918b2e53a1"><td class="memItemLeft" align="right" valign="top"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a64a52f31aa332650601a37643adea428">Stride</a> &amp;&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a8c099a9478d79dbefb04b6918b2e53a1">stride</a> ()</td></tr>
<tr class="memdesc:a8c099a9478d79dbefb04b6918b2e53a1"><td class="mdescLeft">&#160;</td><td class="mdescRight">Returns the stride of the layout. <a href="#a8c099a9478d79dbefb04b6918b2e53a1">More...</a><br /></td></tr>
<tr class="separator:a8c099a9478d79dbefb04b6918b2e53a1"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ae02aa695ed7a272c74153681d8d5c8e9"><td class="memItemLeft" align="right" valign="top"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0ed618d4bb2250c01db5bf9ca61e6839">Index</a>&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#ae02aa695ed7a272c74153681d8d5c8e9">stride</a> (int idx) const </td></tr>
<tr class="memdesc:ae02aa695ed7a272c74153681d8d5c8e9"><td class="mdescLeft">&#160;</td><td class="mdescRight">Returns the stride of the layout. <a href="#ae02aa695ed7a272c74153681d8d5c8e9">More...</a><br /></td></tr>
<tr class="separator:ae02aa695ed7a272c74153681d8d5c8e9"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:aac5e52b103e49eb50b69e93ce058ed59"><td class="memItemLeft" align="right" valign="top"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0ed618d4bb2250c01db5bf9ca61e6839">Index</a> &amp;&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#aac5e52b103e49eb50b69e93ce058ed59">stride</a> (int idx)</td></tr>
<tr class="memdesc:aac5e52b103e49eb50b69e93ce058ed59"><td class="mdescLeft">&#160;</td><td class="mdescRight">Returns the stride of the layout. <a href="#aac5e52b103e49eb50b69e93ce058ed59">More...</a><br /></td></tr>
<tr class="separator:aac5e52b103e49eb50b69e93ce058ed59"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:afa68813c9946c24aa27eae56848a52e6"><td class="memItemLeft" align="right" valign="top"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a3560440cff27c7723c731efd75b3e838">LongIndex</a>&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#afa68813c9946c24aa27eae56848a52e6">capacity</a> (<a class="el" href="structcutlass_1_1MatrixCoord.html">MatrixCoord</a> const &amp;extent) const </td></tr>
<tr class="memdesc:afa68813c9946c24aa27eae56848a52e6"><td class="mdescLeft">&#160;</td><td class="mdescRight">Compute the number of contiguous elements needed to store a tensor with the given size. <a href="#afa68813c9946c24aa27eae56848a52e6">More...</a><br /></td></tr>
<tr class="separator:afa68813c9946c24aa27eae56848a52e6"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="pub-static-methods"></a>
Static Public Member Functions</h2></td></tr>
<tr class="memitem:a17c4a1e21d92742ca725379a963a3fc7"><td class="memItemLeft" align="right" valign="top">static <a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">ColumnMajorBlockLinear</a>&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a17c4a1e21d92742ca725379a963a3fc7">packed</a> (<a class="el" href="structcutlass_1_1MatrixCoord.html">MatrixCoord</a> const &amp;extent)</td></tr>
<tr class="memdesc:a17c4a1e21d92742ca725379a963a3fc7"><td class="mdescLeft">&#160;</td><td class="mdescRight">Helper returns a layout to a tightly packed tensor. <a href="#a17c4a1e21d92742ca725379a963a3fc7">More...</a><br /></td></tr>
<tr class="separator:a17c4a1e21d92742ca725379a963a3fc7"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="pub-static-attribs"></a>
Static Public Attributes</h2></td></tr>
<tr class="memitem:a721cb27c7abd4ccdd476cf5c64cfe82d"><td class="memItemLeft" align="right" valign="top">static int const&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a721cb27c7abd4ccdd476cf5c64cfe82d">kRank</a> = 2</td></tr>
<tr class="memdesc:a721cb27c7abd4ccdd476cf5c64cfe82d"><td class="mdescLeft">&#160;</td><td class="mdescRight">Logical rank of tensor. <a href="#a721cb27c7abd4ccdd476cf5c64cfe82d">More...</a><br /></td></tr>
<tr class="separator:a721cb27c7abd4ccdd476cf5c64cfe82d"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a0732c2b94df9a0c10453cba940684cca"><td class="memItemLeft" align="right" valign="top">static int const&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0732c2b94df9a0c10453cba940684cca">kStrideRank</a> = 1</td></tr>
<tr class="memdesc:a0732c2b94df9a0c10453cba940684cca"><td class="mdescLeft">&#160;</td><td class="mdescRight">Rank of stride vector. <a href="#a0732c2b94df9a0c10453cba940684cca">More...</a><br /></td></tr>
<tr class="separator:a0732c2b94df9a0c10453cba940684cca"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a91fc032744a467348f1b89f31df7bbd6"><td class="memItemLeft" align="right" valign="top">static int const&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a91fc032744a467348f1b89f31df7bbd6">kBlockRows</a> = BlockRows</td></tr>
<tr class="memdesc:a91fc032744a467348f1b89f31df7bbd6"><td class="mdescLeft">&#160;</td><td class="mdescRight">Size of a block in rows. <a href="#a91fc032744a467348f1b89f31df7bbd6">More...</a><br /></td></tr>
<tr class="separator:a91fc032744a467348f1b89f31df7bbd6"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a67c82c42797bb2b0f03c2bc4ee6b5e04"><td class="memItemLeft" align="right" valign="top">static int const&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a67c82c42797bb2b0f03c2bc4ee6b5e04">kBlockColumns</a> = BlockColumns</td></tr>
<tr class="memdesc:a67c82c42797bb2b0f03c2bc4ee6b5e04"><td class="mdescLeft">&#160;</td><td class="mdescRight">Size of a block in columns. <a href="#a67c82c42797bb2b0f03c2bc4ee6b5e04">More...</a><br /></td></tr>
<tr class="separator:a67c82c42797bb2b0f03c2bc4ee6b5e04"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
<a name="details" id="details"></a><h2 class="groupheader">Detailed Description</h2>
<div class="textblock"><h3>template&lt;int BlockRows, int BlockColumns&gt;<br />
struct cutlass::layout::ColumnMajorBlockLinear&lt; BlockRows, BlockColumns &gt;</h3>
<p>Mapping function for block-linear matrices. Matrix is structured as column-major arrangement of 2D tiles (that are column-major). </p>
</div><h2 class="groupheader">Member Typedef Documentation</h2>
<a class="anchor" id="a0ed618d4bb2250c01db5bf9ca61e6839"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="memname">
<tr>
<td class="memname">using <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::<a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0ed618d4bb2250c01db5bf9ca61e6839">Index</a> = int32_t</td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="a3560440cff27c7723c731efd75b3e838"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="memname">
<tr>
<td class="memname">using <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::<a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a3560440cff27c7723c731efd75b3e838">LongIndex</a> = int64_t</td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="a64a52f31aa332650601a37643adea428"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="memname">
<tr>
<td class="memname">using <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::<a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a64a52f31aa332650601a37643adea428">Stride</a> = <a class="el" href="structcutlass_1_1Coord.html">Coord</a>&lt;<a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0732c2b94df9a0c10453cba940684cca">kStrideRank</a>, <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0ed618d4bb2250c01db5bf9ca61e6839">Index</a>&gt;</td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="ae097874037352b7bf9b7640a9830a3aa"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="memname">
<tr>
<td class="memname">using <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::<a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#ae097874037352b7bf9b7640a9830a3aa">TensorCoord</a> = <a class="el" href="structcutlass_1_1MatrixCoord.html">MatrixCoord</a></td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<h2 class="groupheader">Constructor &amp; Destructor Documentation</h2>
<a class="anchor" id="a740b33237fd49dee4692946811e203ba"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::<a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">ColumnMajorBlockLinear</a> </td>
<td>(</td>
<td class="paramtype"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0ed618d4bb2250c01db5bf9ca61e6839">Index</a>&#160;</td>
<td class="paramname"><em>ldm</em> = <code>0</code></td><td>)</td>
<td></td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<h2 class="groupheader">Member Function Documentation</h2>
<a class="anchor" id="afa68813c9946c24aa27eae56848a52e6"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a3560440cff27c7723c731efd75b3e838">LongIndex</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::capacity </td>
<td>(</td>
<td class="paramtype"><a class="el" href="structcutlass_1_1MatrixCoord.html">MatrixCoord</a> const &amp;&#160;</td>
<td class="paramname"><em>extent</em></td><td>)</td>
<td> const</td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="aaf2b7966340292ba65e3cb6163b4348a"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1MatrixCoord.html">MatrixCoord</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::inverse </td>
<td>(</td>
<td class="paramtype"><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a3560440cff27c7723c731efd75b3e838">LongIndex</a>&#160;</td>
<td class="paramname"><em>offset</em></td><td>)</td>
<td> const</td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="a849bbda685abc0b855352c429a8f7e08"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a3560440cff27c7723c731efd75b3e838">LongIndex</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::operator() </td>
<td>(</td>
<td class="paramtype"><a class="el" href="structcutlass_1_1MatrixCoord.html">MatrixCoord</a> const &amp;&#160;</td>
<td class="paramname"><em>coord</em></td><td>)</td>
<td> const</td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
<p>Returns the offset of a coordinate in linear memory. Assumes coordinate has convention (row, column) </p>
</div>
</div>
<a class="anchor" id="a17c4a1e21d92742ca725379a963a3fc7"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname">static <a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">ColumnMajorBlockLinear</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::packed </td>
<td>(</td>
<td class="paramtype"><a class="el" href="structcutlass_1_1MatrixCoord.html">MatrixCoord</a> const &amp;&#160;</td>
<td class="paramname"><em>extent</em></td><td>)</td>
<td></td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span><span class="mlabel">static</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="a7cd57a34850bcf5badbbd760c89aed18"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a64a52f31aa332650601a37643adea428">Stride</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::stride </td>
<td>(</td>
<td class="paramname"></td><td>)</td>
<td> const</td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="a8c099a9478d79dbefb04b6918b2e53a1"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a64a52f31aa332650601a37643adea428">Stride</a>&amp; <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::stride </td>
<td>(</td>
<td class="paramname"></td><td>)</td>
<td></td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="ae02aa695ed7a272c74153681d8d5c8e9"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0ed618d4bb2250c01db5bf9ca61e6839">Index</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::stride </td>
<td>(</td>
<td class="paramtype">int&#160;</td>
<td class="paramname"><em>idx</em></td><td>)</td>
<td> const</td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="aac5e52b103e49eb50b69e93ce058ed59"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html#a0ed618d4bb2250c01db5bf9ca61e6839">Index</a>&amp; <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::stride </td>
<td>(</td>
<td class="paramtype">int&#160;</td>
<td class="paramname"><em>idx</em></td><td>)</td>
<td></td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<h2 class="groupheader">Member Data Documentation</h2>
<a class="anchor" id="a67c82c42797bb2b0f03c2bc4ee6b5e04"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname">int const <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::kBlockColumns = BlockColumns</td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">static</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="a91fc032744a467348f1b89f31df7bbd6"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname">int const <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::kBlockRows = BlockRows</td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">static</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="a721cb27c7abd4ccdd476cf5c64cfe82d"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname">int const <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::kRank = 2</td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">static</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<a class="anchor" id="a0732c2b94df9a0c10453cba940684cca"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;int BlockRows, int BlockColumns&gt; </div>
<table class="mlabels">
<tr>
<td class="mlabels-left">
<table class="memname">
<tr>
<td class="memname">int const <a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html">cutlass::layout::ColumnMajorBlockLinear</a>&lt; BlockRows, BlockColumns &gt;::kStrideRank = 1</td>
</tr>
</table>
</td>
<td class="mlabels-right">
<span class="mlabels"><span class="mlabel">static</span></span> </td>
</tr>
</table>
</div><div class="memdoc">
</div>
</div>
<hr/>The documentation for this struct was generated from the following file:<ul>
<li><a class="el" href="layout_2matrix_8h_source.html">layout/matrix.h</a></li>
</ul>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>