194 lines
157 KiB
HTML
194 lines
157 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: gemm.h Source File</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 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 class="current"><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="files.html"><span>File List</span></a></li>
|
|
<li><a href="globals.html"><span>File 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="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li><li class="navelem"><a class="el" href="dir_9aa36bd9cfad59a1f88859a38871c977.html">gemm</a></li><li class="navelem"><a class="el" href="dir_36528dc2736efa40b421028b7309c671.html">device</a></li> </ul>
|
|
</div>
|
|
</div><!-- top -->
|
|
<div class="header">
|
|
<div class="headertitle">
|
|
<div class="title">include/cutlass/gemm/device/gemm.h</div> </div>
|
|
</div><!--header-->
|
|
<div class="contents">
|
|
<a href="include_2cutlass_2gemm_2device_2gemm_8h.html">Go to the documentation of this file.</a><div class="fragment"><div class="line"><a name="l00001"></a><span class="lineno"> 1</span> <span class="comment">/***************************************************************************************************</span></div><div class="line"><a name="l00002"></a><span class="lineno"> 2</span> <span class="comment"> * Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.</span></div><div class="line"><a name="l00003"></a><span class="lineno"> 3</span> <span class="comment"> *</span></div><div class="line"><a name="l00004"></a><span class="lineno"> 4</span> <span class="comment"> * Redistribution and use in source and binary forms, with or without modification, are permitted</span></div><div class="line"><a name="l00005"></a><span class="lineno"> 5</span> <span class="comment"> * provided that the following conditions are met:</span></div><div class="line"><a name="l00006"></a><span class="lineno"> 6</span> <span class="comment"> * * Redistributions of source code must retain the above copyright notice, this list of</span></div><div class="line"><a name="l00007"></a><span class="lineno"> 7</span> <span class="comment"> * conditions and the following disclaimer.</span></div><div class="line"><a name="l00008"></a><span class="lineno"> 8</span> <span class="comment"> * * Redistributions in binary form must reproduce the above copyright notice, this list of</span></div><div class="line"><a name="l00009"></a><span class="lineno"> 9</span> <span class="comment"> * conditions and the following disclaimer in the documentation and/or other materials</span></div><div class="line"><a name="l00010"></a><span class="lineno"> 10</span> <span class="comment"> * provided with the distribution.</span></div><div class="line"><a name="l00011"></a><span class="lineno"> 11</span> <span class="comment"> * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used</span></div><div class="line"><a name="l00012"></a><span class="lineno"> 12</span> <span class="comment"> * to endorse or promote products derived from this software without specific prior written</span></div><div class="line"><a name="l00013"></a><span class="lineno"> 13</span> <span class="comment"> * permission.</span></div><div class="line"><a name="l00014"></a><span class="lineno"> 14</span> <span class="comment"> *</span></div><div class="line"><a name="l00015"></a><span class="lineno"> 15</span> <span class="comment"> * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR</span></div><div class="line"><a name="l00016"></a><span class="lineno"> 16</span> <span class="comment"> * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND</span></div><div class="line"><a name="l00017"></a><span class="lineno"> 17</span> <span class="comment"> * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE</span></div><div class="line"><a name="l00018"></a><span class="lineno"> 18</span> <span class="comment"> * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,</span></div><div class="line"><a name="l00019"></a><span class="lineno"> 19</span> <span class="comment"> * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;</span></div><div class="line"><a name="l00020"></a><span class="lineno"> 20</span> <span class="comment"> * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,</span></div><div class="line"><a name="l00021"></a><span class="lineno"> 21</span> <span class="comment"> * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE</span></div><div class="line"><a name="l00022"></a><span class="lineno"> 22</span> <span class="comment"> * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.</span></div><div class="line"><a name="l00023"></a><span class="lineno"> 23</span> <span class="comment"> *</span></div><div class="line"><a name="l00024"></a><span class="lineno"> 24</span> <span class="comment"> **************************************************************************************************/</span></div><div class="line"><a name="l00029"></a><span class="lineno"> 29</span> <span class="preprocessor">#pragma once</span></div><div class="line"><a name="l00030"></a><span class="lineno"> 30</span> </div><div class="line"><a name="l00031"></a><span class="lineno"> 31</span> <span class="preprocessor">#include "<a class="code" href="cutlass_8h.html">cutlass/cutlass.h</a>"</span></div><div class="line"><a name="l00032"></a><span class="lineno"> 32</span> <span class="preprocessor">#include "<a class="code" href="numeric__types_8h.html">cutlass/numeric_types.h</a>"</span></div><div class="line"><a name="l00033"></a><span class="lineno"> 33</span> <span class="preprocessor">#include "<a class="code" href="arch_8h.html">cutlass/arch/arch.h</a>"</span></div><div class="line"><a name="l00034"></a><span class="lineno"> 34</span> <span class="preprocessor">#include "<a class="code" href="device__kernel_8h.html">cutlass/device_kernel.h</a>"</span></div><div class="line"><a name="l00035"></a><span class="lineno"> 35</span> </div><div class="line"><a name="l00036"></a><span class="lineno"> 36</span> <span class="preprocessor">#include "<a class="code" href="gemm_2threadblock_2threadblock__swizzle_8h.html">cutlass/gemm/threadblock/threadblock_swizzle.h</a>"</span></div><div class="line"><a name="l00037"></a><span class="lineno"> 37</span> <span class="preprocessor">#include "<a class="code" href="include_2cutlass_2gemm_2kernel_2gemm_8h.html">cutlass/gemm/kernel/gemm.h</a>"</span></div><div class="line"><a name="l00038"></a><span class="lineno"> 38</span> </div><div class="line"><a name="l00039"></a><span class="lineno"> 39</span> <span class="preprocessor">#include "<a class="code" href="default__gemm_8h.html">cutlass/gemm/kernel/default_gemm.h</a>"</span></div><div class="line"><a name="l00040"></a><span class="lineno"> 40</span> <span class="preprocessor">#include "<a class="code" href="default__gemm__configuration_8h.html">cutlass/gemm/device/default_gemm_configuration.h</a>"</span></div><div class="line"><a name="l00041"></a><span class="lineno"> 41</span> </div><div class="line"><a name="l00043"></a><span class="lineno"> 43</span> </div><div class="line"><a name="l00044"></a><span class="lineno"> 44</span> <span class="keyword">namespace </span><a class="code" href="namespacecutlass.html">cutlass</a> {</div><div class="line"><a name="l00045"></a><span class="lineno"> 45</span> <span class="keyword">namespace </span>gemm {</div><div class="line"><a name="l00046"></a><span class="lineno"> 46</span> <span class="keyword">namespace </span>device {</div><div class="line"><a name="l00047"></a><span class="lineno"> 47</span> </div><div class="line"><a name="l00049"></a><span class="lineno"> 49</span> </div><div class="line"><a name="l00113"></a><span class="lineno"> 113</span> </div><div class="line"><a name="l00116"></a><span class="lineno"> 116</span> </div><div class="line"><a name="l00119"></a><span class="lineno"> 119</span> </div><div class="line"><a name="l00122"></a><span class="lineno"> 122</span> </div><div class="line"><a name="l00125"></a><span class="lineno"> 125</span> </div><div class="line"><a name="l00128"></a><span class="lineno"> 128</span> </div><div class="line"><a name="l00131"></a><span class="lineno"> 131</span> </div><div class="line"><a name="l00134"></a><span class="lineno"> 134</span> </div><div class="line"><a name="l00137"></a><span class="lineno"> 137</span> </div><div class="line"><a name="l00140"></a><span class="lineno"> 140</span> </div><div class="line"><a name="l00143"></a><span class="lineno"> 143</span> </div><div class="line"><a name="l00146"></a><span class="lineno"> 146</span> </div><div class="line"><a name="l00149"></a><span class="lineno"> 149</span> </div><div class="line"><a name="l00152"></a><span class="lineno"> 152</span> </div><div class="line"><a name="l00155"></a><span class="lineno"> 155</span> </div><div class="line"><a name="l00159"></a><span class="lineno"> 159</span> <span class="keyword">template</span> <</div><div class="line"><a name="l00161"></a><span class="lineno"> 161</span>  <span class="keyword">typename</span> ElementA_,</div><div class="line"><a name="l00163"></a><span class="lineno"> 163</span>  <span class="keyword">typename</span> LayoutA_,</div><div class="line"><a name="l00165"></a><span class="lineno"> 165</span>  <span class="keyword">typename</span> ElementB_,</div><div class="line"><a name="l00167"></a><span class="lineno"> 167</span>  <span class="keyword">typename</span> LayoutB_,</div><div class="line"><a name="l00169"></a><span class="lineno"> 169</span>  <span class="keyword">typename</span> ElementC_,</div><div class="line"><a name="l00171"></a><span class="lineno"> 171</span>  <span class="keyword">typename</span> LayoutC_,</div><div class="line"><a name="l00173"></a><span class="lineno"> 173</span>  <span class="keyword">typename</span> ElementAccumulator_ = ElementC_,</div><div class="line"><a name="l00175"></a><span class="lineno"> 175</span>  <span class="keyword">typename</span> OperatorClass_ = arch::OpClassSimt,</div><div class="line"><a name="l00177"></a><span class="lineno"> 177</span>  <span class="keyword">typename</span> ArchTag_ = arch::Sm70,</div><div class="line"><a name="l00179"></a><span class="lineno"> 179</span>  <span class="keyword">typename</span> ThreadblockShape_ = <span class="keyword">typename</span> DefaultGemmConfiguration<</div><div class="line"><a name="l00180"></a><span class="lineno"> 180</span>  OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_,</div><div class="line"><a name="l00181"></a><span class="lineno"> 181</span>  ElementAccumulator_>::ThreadblockShape,</div><div class="line"><a name="l00183"></a><span class="lineno"> 183</span>  <span class="keyword">typename</span> WarpShape_ = <span class="keyword">typename</span> DefaultGemmConfiguration<</div><div class="line"><a name="l00184"></a><span class="lineno"> 184</span>  OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_,</div><div class="line"><a name="l00185"></a><span class="lineno"> 185</span>  ElementAccumulator_>::WarpShape,</div><div class="line"><a name="l00187"></a><span class="lineno"> 187</span>  <span class="keyword">typename</span> InstructionShape_ = <span class="keyword">typename</span> DefaultGemmConfiguration<</div><div class="line"><a name="l00188"></a><span class="lineno"> 188</span>  OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_,</div><div class="line"><a name="l00189"></a><span class="lineno"> 189</span>  ElementAccumulator_>::InstructionShape,</div><div class="line"><a name="l00191"></a><span class="lineno"> 191</span>  <span class="keyword">typename</span> EpilogueOutputOp_ = <span class="keyword">typename</span> DefaultGemmConfiguration<</div><div class="line"><a name="l00192"></a><span class="lineno"> 192</span>  OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_,</div><div class="line"><a name="l00193"></a><span class="lineno"> 193</span>  ElementAccumulator_>::EpilogueOutputOp,</div><div class="line"><a name="l00195"></a><span class="lineno"> 195</span>  <span class="keyword">typename</span> ThreadblockSwizzle_ = threadblock::GemmIdentityThreadblockSwizzle,</div><div class="line"><a name="l00197"></a><span class="lineno"> 197</span>  <span class="keywordtype">int</span> Stages =</div><div class="line"><a name="l00198"></a><span class="lineno"> 198</span>  DefaultGemmConfiguration<OperatorClass_, ArchTag_, ElementA_, ElementB_,</div><div class="line"><a name="l00199"></a><span class="lineno"> 199</span>  ElementC_, ElementAccumulator_>::kStages,</div><div class="line"><a name="l00201"></a><span class="lineno"> 201</span>  <span class="keywordtype">int</span> AlignmentA =</div><div class="line"><a name="l00202"></a><span class="lineno"> 202</span>  DefaultGemmConfiguration<OperatorClass_, ArchTag_, ElementA_, ElementB_,</div><div class="line"><a name="l00203"></a><span class="lineno"> 203</span>  ElementC_, ElementAccumulator_>::kAlignmentA,</div><div class="line"><a name="l00205"></a><span class="lineno"> 205</span>  <span class="keywordtype">int</span> AlignmentB =</div><div class="line"><a name="l00206"></a><span class="lineno"> 206</span>  DefaultGemmConfiguration<OperatorClass_, ArchTag_, ElementA_, ElementB_,</div><div class="line"><a name="l00207"></a><span class="lineno"> 207</span>  ElementC_, ElementAccumulator_>::kAlignmentB,</div><div class="line"><a name="l00209"></a><span class="lineno"> 209</span>  <span class="keywordtype">bool</span> SplitKSerial = <span class="keyword">false</span>,</div><div class="line"><a name="l00211"></a><span class="lineno"> 211</span>  <span class="keyword">typename</span> Operator_ = <span class="keyword">typename</span> DefaultGemmConfiguration<</div><div class="line"><a name="l00212"></a><span class="lineno"> 212</span>  OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_,</div><div class="line"><a name="l00213"></a><span class="lineno"> 213</span>  ElementAccumulator_>::Operator,</div><div class="line"><a name="l00215"></a><span class="lineno"> 215</span>  <span class="keywordtype">bool</span> IsBetaZero = <span class="keyword">false</span>></div><div class="line"><a name="l00216"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html"> 216</a></span> <span class="keyword">class </span><a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html">Gemm</a> {</div><div class="line"><a name="l00217"></a><span class="lineno"> 217</span>  <span class="keyword">public</span>:</div><div class="line"><a name="l00218"></a><span class="lineno"> 218</span> </div><div class="line"><a name="l00219"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#adf1ec9ab942ed20321cb91f02154abf0"> 219</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#adf1ec9ab942ed20321cb91f02154abf0">ElementA</a> = ElementA_;</div><div class="line"><a name="l00220"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a6bbefc44e1ed6822dc8a4e2edc2b276b"> 220</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a6bbefc44e1ed6822dc8a4e2edc2b276b">LayoutA</a> = LayoutA_;</div><div class="line"><a name="l00221"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#afc766f8e09f9eb55d50ad23184022c2d"> 221</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1TensorRef.html">TensorRefA</a> = <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementA const, LayoutA></a>;</div><div class="line"><a name="l00222"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a7d04e3640e81787a28956c75acde997e"> 222</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a7d04e3640e81787a28956c75acde997e">ElementB</a> = ElementB_;</div><div class="line"><a name="l00223"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ab389fd311c5763587315c210f40a11f2"> 223</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ab389fd311c5763587315c210f40a11f2">LayoutB</a> = LayoutB_;</div><div class="line"><a name="l00224"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a137892b19acbb4aff34198ab8fdd7c0d"> 224</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1TensorRef.html">TensorRefB</a> = <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementB const, LayoutB></a>;</div><div class="line"><a name="l00225"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af4c31798e982518db6a8fe66e63962e6"> 225</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af4c31798e982518db6a8fe66e63962e6">ElementC</a> = ElementC_;</div><div class="line"><a name="l00226"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a93acaff17a5916bc11040a9ddf04d9f2"> 226</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1layout_1_1RowMajor.html">LayoutC</a> = LayoutC_;</div><div class="line"><a name="l00227"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a2d593dfa19efc9ec65373031ed9f9202"> 227</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1TensorRef.html">TensorRefC</a> = <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC const, LayoutC></a>;</div><div class="line"><a name="l00228"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af39d838d8ab3bba97f555afb714663f2"> 228</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1TensorRef.html">TensorRefD</a> = <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC, LayoutC></a>;</div><div class="line"><a name="l00229"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ad35f92eaf14c268f8bda2a0e48399a9a"> 229</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ad35f92eaf14c268f8bda2a0e48399a9a">ElementAccumulator</a> = ElementAccumulator_;</div><div class="line"><a name="l00230"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a359e8eaddc486fcd0e15b879b11247ad"> 230</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a359e8eaddc486fcd0e15b879b11247ad">OperatorClass</a> = OperatorClass_;</div><div class="line"><a name="l00231"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#afe7ab8c15e83c6cd59b6bcf3fe6e48c0"> 231</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#afe7ab8c15e83c6cd59b6bcf3fe6e48c0">ArchTag</a> = ArchTag_;</div><div class="line"><a name="l00232"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a01ef9ea6588f488e8d727e36bdec8ba8"> 232</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a01ef9ea6588f488e8d727e36bdec8ba8">ThreadblockShape</a> = ThreadblockShape_;</div><div class="line"><a name="l00233"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#adb5bd27efee1575710e6e4157e32449d"> 233</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#adb5bd27efee1575710e6e4157e32449d">WarpShape</a> = WarpShape_;</div><div class="line"><a name="l00234"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ac98211b8adf5c18b4b6c54c5d1cdbb1a"> 234</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ac98211b8adf5c18b4b6c54c5d1cdbb1a">InstructionShape</a> = InstructionShape_;</div><div class="line"><a name="l00235"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af34b612740d98d3066b2b14785f9e415"> 235</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af34b612740d98d3066b2b14785f9e415">EpilogueOutputOp</a> = EpilogueOutputOp_;</div><div class="line"><a name="l00236"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a1eda40e6a86fb3ebeabed2f717e47ced"> 236</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a1eda40e6a86fb3ebeabed2f717e47ced">ThreadblockSwizzle</a> = ThreadblockSwizzle_;</div><div class="line"><a name="l00237"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a6f8a0ff6be313d9f2aac2de4259f65b4"> 237</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a6f8a0ff6be313d9f2aac2de4259f65b4">Operator</a> = Operator_;</div><div class="line"><a name="l00238"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a5cd83a90660626f6c446d45881f0fc22"> 238</a></span>  <span class="keyword">static</span> <span class="keywordtype">int</span> <span class="keyword">const</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a5cd83a90660626f6c446d45881f0fc22">kStages</a> = Stages;</div><div class="line"><a name="l00239"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ad0f1a2669ecc62f24922613e9bad857f"> 239</a></span>  <span class="keyword">static</span> <span class="keywordtype">int</span> <span class="keyword">const</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ad0f1a2669ecc62f24922613e9bad857f">kAlignmentA</a> = AlignmentA;</div><div class="line"><a name="l00240"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a30809c4843226401b9d2da5ecb8d838b"> 240</a></span>  <span class="keyword">static</span> <span class="keywordtype">int</span> <span class="keyword">const</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a30809c4843226401b9d2da5ecb8d838b">kAlignmentB</a> = AlignmentB;</div><div class="line"><a name="l00241"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ac85244d4a91ea8e8c1c3187fad1004db"> 241</a></span>  <span class="keyword">static</span> <span class="keywordtype">int</span> <span class="keyword">const</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ac85244d4a91ea8e8c1c3187fad1004db">kAlignmentC</a> = EpilogueOutputOp::kCount;</div><div class="line"><a name="l00242"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af3c79a0271e684d93c3dca5ad230f45f"> 242</a></span>  <span class="keyword">static</span> <span class="keywordtype">bool</span> <span class="keyword">const</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af3c79a0271e684d93c3dca5ad230f45f">kSplitKSerial</a> = SplitKSerial;</div><div class="line"><a name="l00243"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a0303e85432228170eba7dc4b418c86b4"> 243</a></span>  <span class="keyword">static</span> <span class="keywordtype">bool</span> <span class="keyword">const</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a0303e85432228170eba7dc4b418c86b4">kIsBetaZero</a> = IsBetaZero;</div><div class="line"><a name="l00244"></a><span class="lineno"> 244</span> </div><div class="line"><a name="l00246"></a><span class="lineno"> 246</span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a431ea9cb851566df881457a0f496b976">GemmKernel</a> = <span class="keyword">typename</span> <a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemm.html">kernel::DefaultGemm</a><</div><div class="line"><a name="l00247"></a><span class="lineno"> 247</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#adf1ec9ab942ed20321cb91f02154abf0">ElementA</a>,</div><div class="line"><a name="l00248"></a><span class="lineno"> 248</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a6bbefc44e1ed6822dc8a4e2edc2b276b">LayoutA</a>,</div><div class="line"><a name="l00249"></a><span class="lineno"> 249</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ad0f1a2669ecc62f24922613e9bad857f">kAlignmentA</a>,</div><div class="line"><a name="l00250"></a><span class="lineno"> 250</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a7d04e3640e81787a28956c75acde997e">ElementB</a>,</div><div class="line"><a name="l00251"></a><span class="lineno"> 251</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ab389fd311c5763587315c210f40a11f2">LayoutB</a>,</div><div class="line"><a name="l00252"></a><span class="lineno"> 252</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a30809c4843226401b9d2da5ecb8d838b">kAlignmentB</a>,</div><div class="line"><a name="l00253"></a><span class="lineno"> 253</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af4c31798e982518db6a8fe66e63962e6">ElementC</a>,</div><div class="line"><a name="l00254"></a><span class="lineno"> 254</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a93acaff17a5916bc11040a9ddf04d9f2">LayoutC</a>,</div><div class="line"><a name="l00255"></a><span class="lineno"> 255</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ad35f92eaf14c268f8bda2a0e48399a9a">ElementAccumulator</a>,</div><div class="line"><a name="l00256"></a><span class="lineno"> 256</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a359e8eaddc486fcd0e15b879b11247ad">OperatorClass</a>,</div><div class="line"><a name="l00257"></a><span class="lineno"> 257</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#afe7ab8c15e83c6cd59b6bcf3fe6e48c0">ArchTag</a>,</div><div class="line"><a name="l00258"></a><span class="lineno"> 258</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a01ef9ea6588f488e8d727e36bdec8ba8">ThreadblockShape</a>,</div><div class="line"><a name="l00259"></a><span class="lineno"> 259</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#adb5bd27efee1575710e6e4157e32449d">WarpShape</a>,</div><div class="line"><a name="l00260"></a><span class="lineno"> 260</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ac98211b8adf5c18b4b6c54c5d1cdbb1a">InstructionShape</a>,</div><div class="line"><a name="l00261"></a><span class="lineno"> 261</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af34b612740d98d3066b2b14785f9e415">EpilogueOutputOp</a>,</div><div class="line"><a name="l00262"></a><span class="lineno"> 262</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a1eda40e6a86fb3ebeabed2f717e47ced">ThreadblockSwizzle</a>,</div><div class="line"><a name="l00263"></a><span class="lineno"> 263</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a5cd83a90660626f6c446d45881f0fc22">kStages</a>,</div><div class="line"><a name="l00264"></a><span class="lineno"> 264</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af3c79a0271e684d93c3dca5ad230f45f">kSplitKSerial</a>,</div><div class="line"><a name="l00265"></a><span class="lineno"> 265</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a6f8a0ff6be313d9f2aac2de4259f65b4">Operator</a>,</div><div class="line"><a name="l00266"></a><span class="lineno"> 266</span>  kIsBetaZero</div><div class="line"><a name="l00267"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a431ea9cb851566df881457a0f496b976"> 267</a></span>  ><a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a431ea9cb851566df881457a0f496b976">::GemmKernel</a>;</div><div class="line"><a name="l00268"></a><span class="lineno"> 268</span> </div><div class="line"><a name="l00270"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html"> 270</a></span>  <span class="keyword">struct </span><a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html">Arguments</a> {</div><div class="line"><a name="l00271"></a><span class="lineno"> 271</span> </div><div class="line"><a name="l00272"></a><span class="lineno"> 272</span>  <span class="comment">//</span></div><div class="line"><a name="l00273"></a><span class="lineno"> 273</span>  <span class="comment">// Data members</span></div><div class="line"><a name="l00274"></a><span class="lineno"> 274</span>  <span class="comment">//</span></div><div class="line"><a name="l00275"></a><span class="lineno"> 275</span> </div><div class="line"><a name="l00276"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#ad68cd06dadc163a13f5ed29e07d6535b"> 276</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html">GemmCoord</a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#ad68cd06dadc163a13f5ed29e07d6535b">problem_size</a>;</div><div class="line"><a name="l00277"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a390abae1ca0e01a4b6e58f3724b48eed"> 277</a></span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementA const, LayoutA></a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a390abae1ca0e01a4b6e58f3724b48eed">ref_A</a>;</div><div class="line"><a name="l00278"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#ae712c362f83fbd45679a6e989315d3dc"> 278</a></span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementB const, LayoutB></a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#ae712c362f83fbd45679a6e989315d3dc">ref_B</a>;</div><div class="line"><a name="l00279"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#abdba57a68d6982fffbb1cc3db34ef0f9"> 279</a></span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC const, LayoutC></a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#abdba57a68d6982fffbb1cc3db34ef0f9">ref_C</a>;</div><div class="line"><a name="l00280"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a66f9983db4a09ac0d90291c0f8723897"> 280</a></span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC, LayoutC></a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a66f9983db4a09ac0d90291c0f8723897">ref_D</a>;</div><div class="line"><a name="l00281"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a98c7946bf72f054b026bb6fd49175e90"> 281</a></span>  <span class="keyword">typename</span> EpilogueOutputOp::Params <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a98c7946bf72f054b026bb6fd49175e90">epilogue</a>;</div><div class="line"><a name="l00282"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#ac9c3c764f72e29c3aea99a8f3998e6cd"> 282</a></span>  <span class="keywordtype">int</span> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#ac9c3c764f72e29c3aea99a8f3998e6cd">split_k_slices</a>;</div><div class="line"><a name="l00283"></a><span class="lineno"> 283</span> </div><div class="line"><a name="l00284"></a><span class="lineno"> 284</span>  <span class="comment">//</span></div><div class="line"><a name="l00285"></a><span class="lineno"> 285</span>  <span class="comment">// Methods</span></div><div class="line"><a name="l00286"></a><span class="lineno"> 286</span>  <span class="comment">//</span></div><div class="line"><a name="l00287"></a><span class="lineno"> 287</span> </div><div class="line"><a name="l00289"></a><span class="lineno"> 289</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00290"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa"> 290</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a>(): problem_size(0, 0, 0), split_k_slices(1) {</div><div class="line"><a name="l00291"></a><span class="lineno"> 291</span> </div><div class="line"><a name="l00292"></a><span class="lineno"> 292</span>  }</div><div class="line"><a name="l00293"></a><span class="lineno"> 293</span> </div><div class="line"><a name="l00295"></a><span class="lineno"> 295</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00296"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a975d1d55ab7d42eb02eb226e2bbb5812"> 296</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a975d1d55ab7d42eb02eb226e2bbb5812">Arguments</a>(</div><div class="line"><a name="l00297"></a><span class="lineno"> 297</span>  <a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html">GemmCoord</a> problem_size_,</div><div class="line"><a name="l00298"></a><span class="lineno"> 298</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementA const, LayoutA></a> ref_A_,</div><div class="line"><a name="l00299"></a><span class="lineno"> 299</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementB const, LayoutB></a> ref_B_,</div><div class="line"><a name="l00300"></a><span class="lineno"> 300</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC const, LayoutC></a> ref_C_,</div><div class="line"><a name="l00301"></a><span class="lineno"> 301</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC, LayoutC></a> ref_D_,</div><div class="line"><a name="l00302"></a><span class="lineno"> 302</span>  <span class="keyword">typename</span> EpilogueOutputOp::Params epilogue_ = </div><div class="line"><a name="l00303"></a><span class="lineno"> 303</span>  <span class="keyword">typename</span> EpilogueOutputOp::Params(),</div><div class="line"><a name="l00304"></a><span class="lineno"> 304</span>  <span class="keywordtype">int</span> split_k_slices = 1</div><div class="line"><a name="l00305"></a><span class="lineno"> 305</span>  ):</div><div class="line"><a name="l00306"></a><span class="lineno"> 306</span>  problem_size(problem_size_),</div><div class="line"><a name="l00307"></a><span class="lineno"> 307</span>  ref_A(ref_A_),</div><div class="line"><a name="l00308"></a><span class="lineno"> 308</span>  ref_B(ref_B_),</div><div class="line"><a name="l00309"></a><span class="lineno"> 309</span>  ref_C(ref_C_),</div><div class="line"><a name="l00310"></a><span class="lineno"> 310</span>  ref_D(ref_D_),</div><div class="line"><a name="l00311"></a><span class="lineno"> 311</span>  epilogue(epilogue_),</div><div class="line"><a name="l00312"></a><span class="lineno"> 312</span>  split_k_slices(split_k_slices) {</div><div class="line"><a name="l00313"></a><span class="lineno"> 313</span> </div><div class="line"><a name="l00314"></a><span class="lineno"> 314</span>  }</div><div class="line"><a name="l00315"></a><span class="lineno"> 315</span>  };</div><div class="line"><a name="l00316"></a><span class="lineno"> 316</span> </div><div class="line"><a name="l00317"></a><span class="lineno"> 317</span> <span class="keyword">private</span>:</div><div class="line"><a name="l00318"></a><span class="lineno"> 318</span> </div><div class="line"><a name="l00320"></a><span class="lineno"> 320</span>  <span class="keyword">typename</span> GemmKernel::Params params_;</div><div class="line"><a name="l00321"></a><span class="lineno"> 321</span> </div><div class="line"><a name="l00322"></a><span class="lineno"> 322</span> <span class="keyword">public</span>:</div><div class="line"><a name="l00323"></a><span class="lineno"> 323</span> </div><div class="line"><a name="l00325"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aec04d65c6265eb5f63d703f2dd99cb3f"> 325</a></span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aec04d65c6265eb5f63d703f2dd99cb3f">Gemm</a>() { }</div><div class="line"><a name="l00326"></a><span class="lineno"> 326</span> </div><div class="line"><a name="l00328"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a40ad889da7ff420fd9f9000cd9f98e32"> 328</a></span>  <span class="keyword">static</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a40ad889da7ff420fd9f9000cd9f98e32">can_implement</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args) {</div><div class="line"><a name="l00329"></a><span class="lineno"> 329</span> </div><div class="line"><a name="l00330"></a><span class="lineno"> 330</span>  <span class="keywordflow">if</span> (!kSplitKSerial && args.split_k_slices > 1) {</div><div class="line"><a name="l00331"></a><span class="lineno"> 331</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da58c16acfe39125423ff22c907487ba2c">Status::kErrorInvalidProblem</a>;</div><div class="line"><a name="l00332"></a><span class="lineno"> 332</span>  }</div><div class="line"><a name="l00333"></a><span class="lineno"> 333</span> </div><div class="line"><a name="l00334"></a><span class="lineno"> 334</span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> status = GemmKernel::can_implement(</div><div class="line"><a name="l00335"></a><span class="lineno"> 335</span>  args.problem_size,</div><div class="line"><a name="l00336"></a><span class="lineno"> 336</span>  args.ref_A.non_const_ref(),</div><div class="line"><a name="l00337"></a><span class="lineno"> 337</span>  args.ref_B.non_const_ref(),</div><div class="line"><a name="l00338"></a><span class="lineno"> 338</span>  args.ref_C.non_const_ref(),</div><div class="line"><a name="l00339"></a><span class="lineno"> 339</span>  args.ref_D</div><div class="line"><a name="l00340"></a><span class="lineno"> 340</span>  );</div><div class="line"><a name="l00341"></a><span class="lineno"> 341</span> </div><div class="line"><a name="l00342"></a><span class="lineno"> 342</span>  <span class="keywordflow">if</span> (status != <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>) {</div><div class="line"><a name="l00343"></a><span class="lineno"> 343</span>  <span class="keywordflow">return</span> status;</div><div class="line"><a name="l00344"></a><span class="lineno"> 344</span>  }</div><div class="line"><a name="l00345"></a><span class="lineno"> 345</span> </div><div class="line"><a name="l00346"></a><span class="lineno"> 346</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>;</div><div class="line"><a name="l00347"></a><span class="lineno"> 347</span>  }</div><div class="line"><a name="l00348"></a><span class="lineno"> 348</span> </div><div class="line"><a name="l00350"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a0fec423a58e8de8ff7b015e5167ac614"> 350</a></span>  <span class="keyword">static</span> <span class="keywordtype">size_t</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a0fec423a58e8de8ff7b015e5167ac614">get_workspace_size</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args) {</div><div class="line"><a name="l00351"></a><span class="lineno"> 351</span> </div><div class="line"><a name="l00352"></a><span class="lineno"> 352</span>  <span class="keywordflow">if</span> (kSplitKSerial && args.split_k_slices > 1) {</div><div class="line"><a name="l00353"></a><span class="lineno"> 353</span> </div><div class="line"><a name="l00354"></a><span class="lineno"> 354</span>  <span class="comment">// Determine grid shape</span></div><div class="line"><a name="l00355"></a><span class="lineno"> 355</span>  ThreadblockSwizzle threadblock_swizzle;</div><div class="line"><a name="l00356"></a><span class="lineno"> 356</span> </div><div class="line"><a name="l00357"></a><span class="lineno"> 357</span>  <a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html">cutlass::gemm::GemmCoord</a> tiled_shape = threadblock_swizzle.get_tiled_shape(</div><div class="line"><a name="l00358"></a><span class="lineno"> 358</span>  args.problem_size, </div><div class="line"><a name="l00359"></a><span class="lineno"> 359</span>  {ThreadblockShape::kM, ThreadblockShape::kN, ThreadblockShape::kK},</div><div class="line"><a name="l00360"></a><span class="lineno"> 360</span>  args.split_k_slices);</div><div class="line"><a name="l00361"></a><span class="lineno"> 361</span> </div><div class="line"><a name="l00362"></a><span class="lineno"> 362</span>  <span class="keywordflow">return</span> <span class="keyword">sizeof</span>(int) * <span class="keywordtype">size_t</span>(tiled_shape.<a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html#a93515a41db6c4b7e9101067f60d41b8c">m</a>()) * <span class="keywordtype">size_t</span>(tiled_shape.<a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html#a1b29d2cb15360ad5499216859ad5436a">n</a>());</div><div class="line"><a name="l00363"></a><span class="lineno"> 363</span>  }</div><div class="line"><a name="l00364"></a><span class="lineno"> 364</span> </div><div class="line"><a name="l00365"></a><span class="lineno"> 365</span>  <span class="keywordflow">return</span> 0;</div><div class="line"><a name="l00366"></a><span class="lineno"> 366</span>  }</div><div class="line"><a name="l00367"></a><span class="lineno"> 367</span> </div><div class="line"><a name="l00369"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a53d79d1b434100da1e466e6378ec43ab"> 369</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a53d79d1b434100da1e466e6378ec43ab">initialize</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args, <span class="keywordtype">void</span> *workspace = <span class="keyword">nullptr</span>, cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00370"></a><span class="lineno"> 370</span> </div><div class="line"><a name="l00371"></a><span class="lineno"> 371</span>  <span class="comment">// Determine grid shape</span></div><div class="line"><a name="l00372"></a><span class="lineno"> 372</span>  ThreadblockSwizzle threadblock_swizzle;</div><div class="line"><a name="l00373"></a><span class="lineno"> 373</span> </div><div class="line"><a name="l00374"></a><span class="lineno"> 374</span>  <a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html">cutlass::gemm::GemmCoord</a> grid_shape = threadblock_swizzle.get_tiled_shape(</div><div class="line"><a name="l00375"></a><span class="lineno"> 375</span>  args.problem_size, </div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span>  {ThreadblockShape::kM, ThreadblockShape::kN, ThreadblockShape::kK},</div><div class="line"><a name="l00377"></a><span class="lineno"> 377</span>  args.split_k_slices);</div><div class="line"><a name="l00378"></a><span class="lineno"> 378</span> </div><div class="line"><a name="l00379"></a><span class="lineno"> 379</span>  <span class="keywordflow">if</span> (kSplitKSerial) {</div><div class="line"><a name="l00380"></a><span class="lineno"> 380</span>  <span class="keywordflow">if</span> (args.split_k_slices > 1) {</div><div class="line"><a name="l00381"></a><span class="lineno"> 381</span>  <span class="keywordflow">if</span> (!workspace) {</div><div class="line"><a name="l00382"></a><span class="lineno"> 382</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da6467125354474c5cff2f3c920abb7d3f">Status::kErrorWorkspaceNull</a>;</div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>  }</div><div class="line"><a name="l00384"></a><span class="lineno"> 384</span> </div><div class="line"><a name="l00385"></a><span class="lineno"> 385</span>  <span class="keywordtype">size_t</span> bytes = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a0fec423a58e8de8ff7b015e5167ac614">get_workspace_size</a>(args);</div><div class="line"><a name="l00386"></a><span class="lineno"> 386</span>  </div><div class="line"><a name="l00387"></a><span class="lineno"> 387</span>  cudaError_t result = cudaMemsetAsync(workspace, 0, bytes, stream);</div><div class="line"><a name="l00388"></a><span class="lineno"> 388</span> </div><div class="line"><a name="l00389"></a><span class="lineno"> 389</span>  <span class="keywordflow">if</span> (result != cudaSuccess) {</div><div class="line"><a name="l00390"></a><span class="lineno"> 390</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da4d3b5847a0b17037c6b69faf8b1d4d71">Status::kErrorInternal</a>;</div><div class="line"><a name="l00391"></a><span class="lineno"> 391</span>  }</div><div class="line"><a name="l00392"></a><span class="lineno"> 392</span>  }</div><div class="line"><a name="l00393"></a><span class="lineno"> 393</span>  }</div><div class="line"><a name="l00394"></a><span class="lineno"> 394</span>  <span class="keywordflow">else</span> {</div><div class="line"><a name="l00395"></a><span class="lineno"> 395</span> </div><div class="line"><a name="l00396"></a><span class="lineno"> 396</span>  <span class="keywordflow">if</span> (args.split_k_slices > 1) {</div><div class="line"><a name="l00397"></a><span class="lineno"> 397</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da58c16acfe39125423ff22c907487ba2c">Status::kErrorInvalidProblem</a>;</div><div class="line"><a name="l00398"></a><span class="lineno"> 398</span>  }</div><div class="line"><a name="l00399"></a><span class="lineno"> 399</span>  }</div><div class="line"><a name="l00400"></a><span class="lineno"> 400</span> </div><div class="line"><a name="l00401"></a><span class="lineno"> 401</span>  <span class="comment">// Initialize the Params structure</span></div><div class="line"><a name="l00402"></a><span class="lineno"> 402</span>  params_ = <span class="keyword">typename</span> GemmKernel::Params{</div><div class="line"><a name="l00403"></a><span class="lineno"> 403</span>  args.problem_size,</div><div class="line"><a name="l00404"></a><span class="lineno"> 404</span>  grid_shape,</div><div class="line"><a name="l00405"></a><span class="lineno"> 405</span>  args.ref_A.non_const_ref(),</div><div class="line"><a name="l00406"></a><span class="lineno"> 406</span>  args.ref_B.non_const_ref(),</div><div class="line"><a name="l00407"></a><span class="lineno"> 407</span>  args.ref_C.non_const_ref(),</div><div class="line"><a name="l00408"></a><span class="lineno"> 408</span>  args.ref_D,</div><div class="line"><a name="l00409"></a><span class="lineno"> 409</span>  args.epilogue,</div><div class="line"><a name="l00410"></a><span class="lineno"> 410</span>  <span class="keyword">static_cast<</span><span class="keywordtype">int</span> *<span class="keyword">></span>(workspace)</div><div class="line"><a name="l00411"></a><span class="lineno"> 411</span>  };</div><div class="line"><a name="l00412"></a><span class="lineno"> 412</span> </div><div class="line"><a name="l00413"></a><span class="lineno"> 413</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>;</div><div class="line"><a name="l00414"></a><span class="lineno"> 414</span>  }</div><div class="line"><a name="l00415"></a><span class="lineno"> 415</span> </div><div class="line"><a name="l00417"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aaaa871717d2fbe254a434160bc5acc65"> 417</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aaaa871717d2fbe254a434160bc5acc65">update</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args, <span class="keywordtype">void</span> *workspace = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00418"></a><span class="lineno"> 418</span>  </div><div class="line"><a name="l00419"></a><span class="lineno"> 419</span>  <span class="keywordflow">if</span> (kSplitKSerial && args.split_k_slices > 1) { </div><div class="line"><a name="l00420"></a><span class="lineno"> 420</span>  <span class="keywordflow">if</span> (!workspace) {</div><div class="line"><a name="l00421"></a><span class="lineno"> 421</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da6467125354474c5cff2f3c920abb7d3f">Status::kErrorWorkspaceNull</a>;</div><div class="line"><a name="l00422"></a><span class="lineno"> 422</span>  }</div><div class="line"><a name="l00423"></a><span class="lineno"> 423</span>  }</div><div class="line"><a name="l00424"></a><span class="lineno"> 424</span> </div><div class="line"><a name="l00425"></a><span class="lineno"> 425</span>  params_.ref_A.reset(args.ref_A.non_const_ref().data());</div><div class="line"><a name="l00426"></a><span class="lineno"> 426</span>  params_.ref_B.reset(args.ref_B.non_const_ref().data());</div><div class="line"><a name="l00427"></a><span class="lineno"> 427</span>  params_.ref_C.reset(args.ref_C.non_const_ref().data());</div><div class="line"><a name="l00428"></a><span class="lineno"> 428</span>  params_.ref_D.reset(args.ref_D.data());</div><div class="line"><a name="l00429"></a><span class="lineno"> 429</span>  params_.semaphore = <span class="keyword">static_cast<</span><span class="keywordtype">int</span> *<span class="keyword">></span>(workspace);</div><div class="line"><a name="l00430"></a><span class="lineno"> 430</span> </div><div class="line"><a name="l00431"></a><span class="lineno"> 431</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>;</div><div class="line"><a name="l00432"></a><span class="lineno"> 432</span>  }</div><div class="line"><a name="l00433"></a><span class="lineno"> 433</span> </div><div class="line"><a name="l00435"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aef8c133e539ef91efc9dba9012118147"> 435</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aef8c133e539ef91efc9dba9012118147">run</a>(cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00436"></a><span class="lineno"> 436</span> </div><div class="line"><a name="l00437"></a><span class="lineno"> 437</span>  ThreadblockSwizzle threadblock_swizzle;</div><div class="line"><a name="l00438"></a><span class="lineno"> 438</span> </div><div class="line"><a name="l00439"></a><span class="lineno"> 439</span>  dim3 grid = threadblock_swizzle.get_grid_shape(params_.grid_tiled_shape);</div><div class="line"><a name="l00440"></a><span class="lineno"> 440</span>  dim3 block(GemmKernel::kThreadCount, 1, 1);</div><div class="line"><a name="l00441"></a><span class="lineno"> 441</span> </div><div class="line"><a name="l00442"></a><span class="lineno"> 442</span>  cudaError_t result;</div><div class="line"><a name="l00443"></a><span class="lineno"> 443</span> </div><div class="line"><a name="l00444"></a><span class="lineno"> 444</span>  <span class="keywordtype">int</span> smem_size = int(<span class="keyword">sizeof</span>(<span class="keyword">typename</span> GemmKernel::SharedStorage));</div><div class="line"><a name="l00445"></a><span class="lineno"> 445</span>  <span class="keywordflow">if</span> (smem_size >= (48 << 10)) {</div><div class="line"><a name="l00446"></a><span class="lineno"> 446</span>  result = cudaFuncSetAttribute(Kernel<GemmKernel>,</div><div class="line"><a name="l00447"></a><span class="lineno"> 447</span>  cudaFuncAttributeMaxDynamicSharedMemorySize,</div><div class="line"><a name="l00448"></a><span class="lineno"> 448</span>  smem_size);</div><div class="line"><a name="l00449"></a><span class="lineno"> 449</span> </div><div class="line"><a name="l00450"></a><span class="lineno"> 450</span>  <span class="keywordflow">if</span> (result != cudaSuccess) {</div><div class="line"><a name="l00451"></a><span class="lineno"> 451</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da4d3b5847a0b17037c6b69faf8b1d4d71">Status::kErrorInternal</a>;</div><div class="line"><a name="l00452"></a><span class="lineno"> 452</span>  }</div><div class="line"><a name="l00453"></a><span class="lineno"> 453</span> </div><div class="line"><a name="l00454"></a><span class="lineno"> 454</span>  result = cudaFuncSetAttribute(</div><div class="line"><a name="l00455"></a><span class="lineno"> 455</span>  Kernel<GemmKernel>,</div><div class="line"><a name="l00456"></a><span class="lineno"> 456</span>  cudaFuncAttributePreferredSharedMemoryCarveout, 100);</div><div class="line"><a name="l00457"></a><span class="lineno"> 457</span> </div><div class="line"><a name="l00458"></a><span class="lineno"> 458</span>  <span class="keywordflow">if</span> (result != cudaSuccess) {</div><div class="line"><a name="l00459"></a><span class="lineno"> 459</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da4d3b5847a0b17037c6b69faf8b1d4d71">Status::kErrorInternal</a>;</div><div class="line"><a name="l00460"></a><span class="lineno"> 460</span>  }</div><div class="line"><a name="l00461"></a><span class="lineno"> 461</span>  }</div><div class="line"><a name="l00462"></a><span class="lineno"> 462</span> </div><div class="line"><a name="l00463"></a><span class="lineno"> 463</span>  cutlass::Kernel<GemmKernel><<<grid, block, smem_size, stream>>>(params_);</div><div class="line"><a name="l00464"></a><span class="lineno"> 464</span> </div><div class="line"><a name="l00465"></a><span class="lineno"> 465</span>  result = cudaGetLastError();</div><div class="line"><a name="l00466"></a><span class="lineno"> 466</span> </div><div class="line"><a name="l00467"></a><span class="lineno"> 467</span>  <span class="keywordflow">return</span> result == cudaSuccess ? <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a> : <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da4d3b5847a0b17037c6b69faf8b1d4d71">Status::kErrorInternal</a>;</div><div class="line"><a name="l00468"></a><span class="lineno"> 468</span>  }</div><div class="line"><a name="l00469"></a><span class="lineno"> 469</span> </div><div class="line"><a name="l00471"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a114b9c14e102d333ef3dcad7865a4efb"> 471</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a114b9c14e102d333ef3dcad7865a4efb">operator()</a>(cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00472"></a><span class="lineno"> 472</span>  <span class="keywordflow">return</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aef8c133e539ef91efc9dba9012118147">run</a>(stream);</div><div class="line"><a name="l00473"></a><span class="lineno"> 473</span>  }</div><div class="line"><a name="l00474"></a><span class="lineno"> 474</span> </div><div class="line"><a name="l00476"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a35d9f803fdfbbd4608243881c04316dc"> 476</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a35d9f803fdfbbd4608243881c04316dc">operator()</a>(</div><div class="line"><a name="l00477"></a><span class="lineno"> 477</span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args, </div><div class="line"><a name="l00478"></a><span class="lineno"> 478</span>  <span class="keywordtype">void</span> *workspace = <span class="keyword">nullptr</span>, </div><div class="line"><a name="l00479"></a><span class="lineno"> 479</span>  cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00480"></a><span class="lineno"> 480</span>  </div><div class="line"><a name="l00481"></a><span class="lineno"> 481</span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> status = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a53d79d1b434100da1e466e6378ec43ab">initialize</a>(args, workspace);</div><div class="line"><a name="l00482"></a><span class="lineno"> 482</span>  </div><div class="line"><a name="l00483"></a><span class="lineno"> 483</span>  <span class="keywordflow">if</span> (status == <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>) {</div><div class="line"><a name="l00484"></a><span class="lineno"> 484</span>  status = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aef8c133e539ef91efc9dba9012118147">run</a>(stream);</div><div class="line"><a name="l00485"></a><span class="lineno"> 485</span>  }</div><div class="line"><a name="l00486"></a><span class="lineno"> 486</span> </div><div class="line"><a name="l00487"></a><span class="lineno"> 487</span>  <span class="keywordflow">return</span> status;</div><div class="line"><a name="l00488"></a><span class="lineno"> 488</span>  }</div><div class="line"><a name="l00489"></a><span class="lineno"> 489</span> };</div><div class="line"><a name="l00490"></a><span class="lineno"> 490</span> </div><div class="line"><a name="l00492"></a><span class="lineno"> 492</span> </div><div class="line"><a name="l00494"></a><span class="lineno"> 494</span> <span class="keyword">template</span> <</div><div class="line"><a name="l00496"></a><span class="lineno"> 496</span>  <span class="keyword">typename</span> ElementA_,</div><div class="line"><a name="l00498"></a><span class="lineno"> 498</span>  <span class="keyword">typename</span> LayoutA_,</div><div class="line"><a name="l00500"></a><span class="lineno"> 500</span>  <span class="keyword">typename</span> ElementB_,</div><div class="line"><a name="l00502"></a><span class="lineno"> 502</span>  <span class="keyword">typename</span> LayoutB_,</div><div class="line"><a name="l00504"></a><span class="lineno"> 504</span>  <span class="keyword">typename</span> ElementC_,</div><div class="line"><a name="l00506"></a><span class="lineno"> 506</span>  <span class="keyword">typename</span> ElementAccumulator_,</div><div class="line"><a name="l00508"></a><span class="lineno"> 508</span>  <span class="keyword">typename</span> OperatorClass_,</div><div class="line"><a name="l00510"></a><span class="lineno"> 510</span>  <span class="keyword">typename</span> ArchTag_,</div><div class="line"><a name="l00512"></a><span class="lineno"> 512</span>  <span class="keyword">typename</span> ThreadblockShape_,</div><div class="line"><a name="l00514"></a><span class="lineno"> 514</span>  <span class="keyword">typename</span> WarpShape_,</div><div class="line"><a name="l00516"></a><span class="lineno"> 516</span>  <span class="keyword">typename</span> InstructionShape_,</div><div class="line"><a name="l00518"></a><span class="lineno"> 518</span>  <span class="keyword">typename</span> EpilogueOutputOp_,</div><div class="line"><a name="l00520"></a><span class="lineno"> 520</span>  <span class="keyword">typename</span> ThreadblockSwizzle_,</div><div class="line"><a name="l00522"></a><span class="lineno"> 522</span>  <span class="keywordtype">int</span> Stages,</div><div class="line"><a name="l00524"></a><span class="lineno"> 524</span>  <span class="keywordtype">int</span> AlignmentA,</div><div class="line"><a name="l00526"></a><span class="lineno"> 526</span>  <span class="keywordtype">int</span> AlignmentB,</div><div class="line"><a name="l00528"></a><span class="lineno"> 528</span>  <span class="keywordtype">bool</span> SplitKSerial,</div><div class="line"><a name="l00530"></a><span class="lineno"> 530</span>  <span class="keyword">typename</span> Operator_,</div><div class="line"><a name="l00532"></a><span class="lineno"> 532</span>  <span class="keywordtype">bool</span> IsBetaZero></div><div class="line"><a name="l00533"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html"> 533</a></span> <span class="keyword">class </span><a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html">Gemm</a><ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_,</div><div class="line"><a name="l00534"></a><span class="lineno"> 534</span>  layout::ColumnMajor, <span class="comment">// partially specialized on LayoutC</span></div><div class="line"><a name="l00535"></a><span class="lineno"> 535</span>  ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_,</div><div class="line"><a name="l00536"></a><span class="lineno"> 536</span>  WarpShape_, InstructionShape_, EpilogueOutputOp_,</div><div class="line"><a name="l00537"></a><span class="lineno"> 537</span>  ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial,</div><div class="line"><a name="l00538"></a><span class="lineno"> 538</span>  Operator_, IsBetaZero> {</div><div class="line"><a name="l00539"></a><span class="lineno"> 539</span>  <span class="keyword">public</span>:</div><div class="line"><a name="l00540"></a><span class="lineno"> 540</span> </div><div class="line"><a name="l00541"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a09db4f8f255d272e7350394d568f4a01"> 541</a></span>  <span class="keyword">using</span> ElementA = ElementA_;</div><div class="line"><a name="l00542"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a5212eb5b3af32e5bc43cc4179bb346ef"> 542</a></span>  <span class="keyword">using</span> LayoutA = LayoutA_;</div><div class="line"><a name="l00543"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a82da12cbdf6f75499d315ee530f5330e"> 543</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1TensorRef.html">TensorRefA</a> = <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementA const, LayoutA></a>;</div><div class="line"><a name="l00544"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a0f904b72a3ff91f7ae6ad1a91e915b6d"> 544</a></span>  <span class="keyword">using</span> ElementB = ElementB_;</div><div class="line"><a name="l00545"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a1841e0e97e59862c7a92fc8d2ab7c9bc"> 545</a></span>  <span class="keyword">using</span> LayoutB = LayoutB_;</div><div class="line"><a name="l00546"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#ab4bd8a2bb7be0fa2b583cf34b63b62eb"> 546</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1TensorRef.html">TensorRefB</a> = <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementB const, LayoutB></a>;</div><div class="line"><a name="l00547"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#ad58a37fecfeb982d20fc209a0df4c1fa"> 547</a></span>  <span class="keyword">using</span> ElementC = ElementC_;</div><div class="line"><a name="l00548"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#afe4685fea6a4603a7459bbe9923c9cb3"> 548</a></span>  <span class="keyword">using</span> LayoutC = <a class="code" href="classcutlass_1_1layout_1_1ColumnMajor.html">layout::ColumnMajor</a>;</div><div class="line"><a name="l00549"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#aff1ad6d93937a9e4b261eb69322449e7"> 549</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1TensorRef.html">TensorRefC</a> = <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC const, LayoutC></a>;</div><div class="line"><a name="l00550"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#abaa02d78437ae0f42260848d722c134f"> 550</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1TensorRef.html">TensorRefD</a> = <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC, LayoutC></a>;</div><div class="line"><a name="l00551"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#acbc61142b95f4d33bc0b8518857ab7be"> 551</a></span>  <span class="keyword">using</span> ElementAccumulator = ElementAccumulator_;</div><div class="line"><a name="l00552"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#aa816137b589b8bf2204ace73d49b7ded"> 552</a></span>  <span class="keyword">using</span> OperatorClass = OperatorClass_;</div><div class="line"><a name="l00553"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a1b502a4097e745c12d0d628d080ba447"> 553</a></span>  <span class="keyword">using</span> ArchTag = ArchTag_;</div><div class="line"><a name="l00554"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#ab798f409ba80eab4a0140fdf43e768ee"> 554</a></span>  <span class="keyword">using</span> ThreadblockShape = ThreadblockShape_;</div><div class="line"><a name="l00555"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#abdc15293a8b083372e5395049440d01c"> 555</a></span>  <span class="keyword">using</span> WarpShape = WarpShape_;</div><div class="line"><a name="l00556"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a71b84f983b94b50a48bd0890f1e0ed59"> 556</a></span>  <span class="keyword">using</span> InstructionShape = InstructionShape_;</div><div class="line"><a name="l00557"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#ad1d60f7381ae03803a078a26604bd8be"> 557</a></span>  <span class="keyword">using</span> EpilogueOutputOp = EpilogueOutputOp_;</div><div class="line"><a name="l00558"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#af55d56acaa01ce303c22d6e9e0b0f895"> 558</a></span>  <span class="keyword">using</span> ThreadblockSwizzle = ThreadblockSwizzle_;</div><div class="line"><a name="l00559"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a029f48f17ec3fb98067bfacd7e06f3d2"> 559</a></span>  <span class="keyword">using</span> Operator = Operator_;</div><div class="line"><a name="l00560"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a5de0cfa9c3831daebbdc8326c239dd33"> 560</a></span>  <span class="keyword">static</span> <span class="keywordtype">int</span> <span class="keyword">const</span> kStages = Stages;</div><div class="line"><a name="l00561"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a78660ed036162b8455546ff5718968d0"> 561</a></span>  <span class="keyword">static</span> <span class="keywordtype">int</span> <span class="keyword">const</span> kAlignmentA = AlignmentA;</div><div class="line"><a name="l00562"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#aa39221ab9fa4248c613b7222f764072e"> 562</a></span>  <span class="keyword">static</span> <span class="keywordtype">int</span> <span class="keyword">const</span> kAlignmentB = AlignmentB;</div><div class="line"><a name="l00563"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a0b609010f97cb53cf4d8f1ecb4bb0b79"> 563</a></span>  <span class="keyword">static</span> <span class="keywordtype">bool</span> <span class="keyword">const</span> kSplitKSerial = SplitKSerial;</div><div class="line"><a name="l00564"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a98d1d07f32f29b29e883775fcd276833"> 564</a></span>  <span class="keyword">static</span> <span class="keywordtype">bool</span> <span class="keyword">const</span> kIsBetaZero = IsBetaZero;</div><div class="line"><a name="l00565"></a><span class="lineno"> 565</span> </div><div class="line"><a name="l00566"></a><span class="lineno"> 566</span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html">UnderlyingOperator</a> = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html">Gemm</a>< </div><div class="line"><a name="l00567"></a><span class="lineno"> 567</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a7d04e3640e81787a28956c75acde997e">ElementB</a>,</div><div class="line"><a name="l00568"></a><span class="lineno"> 568</span>  <span class="keyword">typename</span> <a class="code" href="structcutlass_1_1layout_1_1LayoutTranspose.html">layout::LayoutTranspose<LayoutB>::type</a>,</div><div class="line"><a name="l00569"></a><span class="lineno"> 569</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#adf1ec9ab942ed20321cb91f02154abf0">ElementA</a>,</div><div class="line"><a name="l00570"></a><span class="lineno"> 570</span>  <span class="keyword">typename</span> <a class="code" href="structcutlass_1_1layout_1_1LayoutTranspose.html">layout::LayoutTranspose<LayoutA>::type</a>,</div><div class="line"><a name="l00571"></a><span class="lineno"> 571</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af4c31798e982518db6a8fe66e63962e6">ElementC</a>,</div><div class="line"><a name="l00572"></a><span class="lineno"> 572</span>  <a class="code" href="classcutlass_1_1layout_1_1RowMajor.html">layout::RowMajor</a>, </div><div class="line"><a name="l00573"></a><span class="lineno"> 573</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ad35f92eaf14c268f8bda2a0e48399a9a">ElementAccumulator</a>,</div><div class="line"><a name="l00574"></a><span class="lineno"> 574</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a359e8eaddc486fcd0e15b879b11247ad">OperatorClass</a>,</div><div class="line"><a name="l00575"></a><span class="lineno"> 575</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#afe7ab8c15e83c6cd59b6bcf3fe6e48c0">ArchTag</a>,</div><div class="line"><a name="l00576"></a><span class="lineno"> 576</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a01ef9ea6588f488e8d727e36bdec8ba8">ThreadblockShape</a>,</div><div class="line"><a name="l00577"></a><span class="lineno"> 577</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#adb5bd27efee1575710e6e4157e32449d">WarpShape</a>,</div><div class="line"><a name="l00578"></a><span class="lineno"> 578</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ac98211b8adf5c18b4b6c54c5d1cdbb1a">InstructionShape</a>,</div><div class="line"><a name="l00579"></a><span class="lineno"> 579</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af34b612740d98d3066b2b14785f9e415">EpilogueOutputOp</a>,</div><div class="line"><a name="l00580"></a><span class="lineno"> 580</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a1eda40e6a86fb3ebeabed2f717e47ced">ThreadblockSwizzle</a>,</div><div class="line"><a name="l00581"></a><span class="lineno"> 581</span>  Stages,</div><div class="line"><a name="l00582"></a><span class="lineno"> 582</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a30809c4843226401b9d2da5ecb8d838b">kAlignmentB</a>,</div><div class="line"><a name="l00583"></a><span class="lineno"> 583</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ad0f1a2669ecc62f24922613e9bad857f">kAlignmentA</a>,</div><div class="line"><a name="l00584"></a><span class="lineno"> 584</span>  SplitKSerial,</div><div class="line"><a name="l00585"></a><span class="lineno"> 585</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a6f8a0ff6be313d9f2aac2de4259f65b4">Operator</a>,</div><div class="line"><a name="l00586"></a><span class="lineno"> 586</span>  kIsBetaZero</div><div class="line"><a name="l00587"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#aa48a0b3645f2ef103cb1b2d41218d865"> 587</a></span>  >;</div><div class="line"><a name="l00588"></a><span class="lineno"> 588</span> </div><div class="line"><a name="l00589"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a7615ad046304360243729c29c65e878a"> 589</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a7615ad046304360243729c29c65e878a">UnderlyingArguments</a> = <span class="keyword">typename</span> UnderlyingOperator::Arguments;</div><div class="line"><a name="l00590"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a2bdbc5e737f9bfd1e09a7cfb30e60e29"> 590</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a2bdbc5e737f9bfd1e09a7cfb30e60e29">GemmKernel</a> = <span class="keyword">typename</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a431ea9cb851566df881457a0f496b976">UnderlyingOperator::GemmKernel</a>;</div><div class="line"><a name="l00591"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#ad74a049e26f4b9224362b4d1c93ca14b"> 591</a></span>  <span class="keyword">static</span> <span class="keywordtype">int</span> <span class="keyword">const</span> kAlignmentC = UnderlyingOperator::kAlignmentC;</div><div class="line"><a name="l00592"></a><span class="lineno"> 592</span> </div><div class="line"><a name="l00594"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html"> 594</a></span>  <span class="keyword">struct </span><a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> {</div><div class="line"><a name="l00595"></a><span class="lineno"> 595</span> </div><div class="line"><a name="l00596"></a><span class="lineno"> 596</span>  <span class="comment">//</span></div><div class="line"><a name="l00597"></a><span class="lineno"> 597</span>  <span class="comment">// Data members</span></div><div class="line"><a name="l00598"></a><span class="lineno"> 598</span>  <span class="comment">//</span></div><div class="line"><a name="l00599"></a><span class="lineno"> 599</span> </div><div class="line"><a name="l00600"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#acd02e86dfff866eade08415e0043ccc3"> 600</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html">GemmCoord</a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#acd02e86dfff866eade08415e0043ccc3">problem_size</a>;</div><div class="line"><a name="l00601"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a9bdaf3563983efcca649460be169b334"> 601</a></span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementA const, LayoutA></a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a9bdaf3563983efcca649460be169b334">ref_A</a>;</div><div class="line"><a name="l00602"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#ab77204c1010b17c6643d26a89f41c3d0"> 602</a></span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementB const, LayoutB></a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#ab77204c1010b17c6643d26a89f41c3d0">ref_B</a>;</div><div class="line"><a name="l00603"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a590b8da88ae9350042838451e3e37a22"> 603</a></span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC const, LayoutC></a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a590b8da88ae9350042838451e3e37a22">ref_C</a>;</div><div class="line"><a name="l00604"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#ab1d4d5865786a415f87db1def1b029e7"> 604</a></span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC, LayoutC></a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#ab1d4d5865786a415f87db1def1b029e7">ref_D</a>;</div><div class="line"><a name="l00605"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a426f402c08be99849a4477a07f010a5e"> 605</a></span>  <span class="keyword">typename</span> EpilogueOutputOp::Params <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a426f402c08be99849a4477a07f010a5e">epilogue</a>;</div><div class="line"><a name="l00606"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#aaef8450711318fa1a53fe3cb72b59263"> 606</a></span>  <span class="keywordtype">int</span> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#aaef8450711318fa1a53fe3cb72b59263">split_k_slices</a>;</div><div class="line"><a name="l00607"></a><span class="lineno"> 607</span> </div><div class="line"><a name="l00608"></a><span class="lineno"> 608</span>  <span class="comment">//</span></div><div class="line"><a name="l00609"></a><span class="lineno"> 609</span>  <span class="comment">// Methods</span></div><div class="line"><a name="l00610"></a><span class="lineno"> 610</span>  <span class="comment">//</span></div><div class="line"><a name="l00611"></a><span class="lineno"> 611</span> </div><div class="line"><a name="l00613"></a><span class="lineno"> 613</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00614"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#ac6c397a181a52c0dbb39bf3710ee4658"> 614</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#ac6c397a181a52c0dbb39bf3710ee4658">Arguments</a>() { }</div><div class="line"><a name="l00615"></a><span class="lineno"> 615</span> </div><div class="line"><a name="l00617"></a><span class="lineno"> 617</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00618"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a331de1adfdcbea6d0137afe64a4f6f4c"> 618</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a331de1adfdcbea6d0137afe64a4f6f4c">Arguments</a>(</div><div class="line"><a name="l00619"></a><span class="lineno"> 619</span>  <a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html">GemmCoord</a> problem_size_,</div><div class="line"><a name="l00620"></a><span class="lineno"> 620</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementA const, LayoutA></a> ref_A_,</div><div class="line"><a name="l00621"></a><span class="lineno"> 621</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementB const, LayoutB></a> ref_B_,</div><div class="line"><a name="l00622"></a><span class="lineno"> 622</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC const, LayoutC></a> ref_C_,</div><div class="line"><a name="l00623"></a><span class="lineno"> 623</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC, LayoutC></a> ref_D_,</div><div class="line"><a name="l00624"></a><span class="lineno"> 624</span>  <span class="keyword">typename</span> EpilogueOutputOp::Params epilogue_ = </div><div class="line"><a name="l00625"></a><span class="lineno"> 625</span>  <span class="keyword">typename</span> EpilogueOutputOp::Params(),</div><div class="line"><a name="l00626"></a><span class="lineno"> 626</span>  <span class="keywordtype">int</span> split_k_slices = 1</div><div class="line"><a name="l00627"></a><span class="lineno"> 627</span>  ):</div><div class="line"><a name="l00628"></a><span class="lineno"> 628</span>  problem_size(problem_size_),</div><div class="line"><a name="l00629"></a><span class="lineno"> 629</span>  ref_A(ref_A_),</div><div class="line"><a name="l00630"></a><span class="lineno"> 630</span>  ref_B(ref_B_),</div><div class="line"><a name="l00631"></a><span class="lineno"> 631</span>  ref_C(ref_C_),</div><div class="line"><a name="l00632"></a><span class="lineno"> 632</span>  ref_D(ref_D_),</div><div class="line"><a name="l00633"></a><span class="lineno"> 633</span>  epilogue(epilogue_),</div><div class="line"><a name="l00634"></a><span class="lineno"> 634</span>  split_k_slices(split_k_slices) { }</div><div class="line"><a name="l00635"></a><span class="lineno"> 635</span>  };</div><div class="line"><a name="l00636"></a><span class="lineno"> 636</span> </div><div class="line"><a name="l00637"></a><span class="lineno"> 637</span> <span class="keyword">private</span>:</div><div class="line"><a name="l00638"></a><span class="lineno"> 638</span> </div><div class="line"><a name="l00639"></a><span class="lineno"> 639</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html">UnderlyingOperator</a> underlying_operator_;</div><div class="line"><a name="l00640"></a><span class="lineno"> 640</span> </div><div class="line"><a name="l00641"></a><span class="lineno"> 641</span> <span class="keyword">public</span>:</div><div class="line"><a name="l00642"></a><span class="lineno"> 642</span> </div><div class="line"><a name="l00644"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#abcacf502806db50eb17a6d925aee16d5"> 644</a></span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#abcacf502806db50eb17a6d925aee16d5">Gemm</a>() { }</div><div class="line"><a name="l00645"></a><span class="lineno"> 645</span> </div><div class="line"><a name="l00647"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#aa9313915a6129f0c43b43ef3698b3ee4"> 647</a></span>  <span class="keyword">static</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a7615ad046304360243729c29c65e878a">UnderlyingArguments</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#aa9313915a6129f0c43b43ef3698b3ee4">to_underlying_arguments</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args) {</div><div class="line"><a name="l00648"></a><span class="lineno"> 648</span>  <span class="keywordflow">return</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a7615ad046304360243729c29c65e878a">UnderlyingArguments</a>(</div><div class="line"><a name="l00649"></a><span class="lineno"> 649</span>  {args.problem_size.n(), args.problem_size.m(), args.problem_size.k()},</div><div class="line"><a name="l00650"></a><span class="lineno"> 650</span>  {args.ref_B.data(), args.ref_B.stride(0)},</div><div class="line"><a name="l00651"></a><span class="lineno"> 651</span>  {args.ref_A.data(), args.ref_A.stride(0)},</div><div class="line"><a name="l00652"></a><span class="lineno"> 652</span>  {args.ref_C.data(), args.ref_C.stride(0)},</div><div class="line"><a name="l00653"></a><span class="lineno"> 653</span>  {args.ref_D.data(), args.ref_D.stride(0)},</div><div class="line"><a name="l00654"></a><span class="lineno"> 654</span>  args.epilogue,</div><div class="line"><a name="l00655"></a><span class="lineno"> 655</span>  args.split_k_slices</div><div class="line"><a name="l00656"></a><span class="lineno"> 656</span>  );</div><div class="line"><a name="l00657"></a><span class="lineno"> 657</span>  }</div><div class="line"><a name="l00658"></a><span class="lineno"> 658</span> </div><div class="line"><a name="l00660"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a662bcbcb6164c803ab490c86e69b9ee1"> 660</a></span>  <span class="keyword">static</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a662bcbcb6164c803ab490c86e69b9ee1">can_implement</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args) {</div><div class="line"><a name="l00661"></a><span class="lineno"> 661</span> </div><div class="line"><a name="l00662"></a><span class="lineno"> 662</span>  <span class="keywordflow">return</span> UnderlyingOperator::can_implement(to_underlying_arguments(args));</div><div class="line"><a name="l00663"></a><span class="lineno"> 663</span>  }</div><div class="line"><a name="l00664"></a><span class="lineno"> 664</span> </div><div class="line"><a name="l00666"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a1469133c30fde6b28296e3ff6951e7a4"> 666</a></span>  <span class="keyword">static</span> <span class="keywordtype">size_t</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a1469133c30fde6b28296e3ff6951e7a4">get_workspace_size</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args) {</div><div class="line"><a name="l00667"></a><span class="lineno"> 667</span>  </div><div class="line"><a name="l00668"></a><span class="lineno"> 668</span>  <span class="keywordflow">return</span> UnderlyingOperator::get_workspace_size(to_underlying_arguments(args));</div><div class="line"><a name="l00669"></a><span class="lineno"> 669</span>  }</div><div class="line"><a name="l00670"></a><span class="lineno"> 670</span> </div><div class="line"><a name="l00672"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a7a14474e4238d2fac92ad71c6de087d8"> 672</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a7a14474e4238d2fac92ad71c6de087d8">initialize</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args, <span class="keywordtype">void</span> *workspace = <span class="keyword">nullptr</span>, cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00673"></a><span class="lineno"> 673</span> </div><div class="line"><a name="l00674"></a><span class="lineno"> 674</span>  <span class="keywordflow">return</span> underlying_operator_.<a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a53d79d1b434100da1e466e6378ec43ab">initialize</a>(to_underlying_arguments(args), workspace);</div><div class="line"><a name="l00675"></a><span class="lineno"> 675</span>  }</div><div class="line"><a name="l00676"></a><span class="lineno"> 676</span> </div><div class="line"><a name="l00678"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a2b6c5275c173d73cffe8e6b6b1ccf2c1"> 678</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a2b6c5275c173d73cffe8e6b6b1ccf2c1">update</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args, <span class="keywordtype">void</span> *workspace = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00679"></a><span class="lineno"> 679</span> </div><div class="line"><a name="l00680"></a><span class="lineno"> 680</span>  <span class="keywordflow">return</span> underlying_operator_.<a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aaaa871717d2fbe254a434160bc5acc65">update</a>(to_underlying_arguments(args), workspace);</div><div class="line"><a name="l00681"></a><span class="lineno"> 681</span>  }</div><div class="line"><a name="l00682"></a><span class="lineno"> 682</span> </div><div class="line"><a name="l00684"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a5f4f93ca97b358b4410f3d0b1e0a6387"> 684</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a5f4f93ca97b358b4410f3d0b1e0a6387">run</a>(cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00685"></a><span class="lineno"> 685</span> </div><div class="line"><a name="l00686"></a><span class="lineno"> 686</span>  <span class="keywordflow">return</span> underlying_operator_.<a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aef8c133e539ef91efc9dba9012118147">run</a>(stream);</div><div class="line"><a name="l00687"></a><span class="lineno"> 687</span>  }</div><div class="line"><a name="l00688"></a><span class="lineno"> 688</span> </div><div class="line"><a name="l00690"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a384db4125183e504fafc5a946b7ba757"> 690</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a384db4125183e504fafc5a946b7ba757">operator()</a>(cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00691"></a><span class="lineno"> 691</span>  <span class="keywordflow">return</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aef8c133e539ef91efc9dba9012118147">run</a>(stream);</div><div class="line"><a name="l00692"></a><span class="lineno"> 692</span>  }</div><div class="line"><a name="l00693"></a><span class="lineno"> 693</span> </div><div class="line"><a name="l00695"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a6115aa957b3ba8ad9e54b7efeefaacd1"> 695</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a6115aa957b3ba8ad9e54b7efeefaacd1">operator()</a>(</div><div class="line"><a name="l00696"></a><span class="lineno"> 696</span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">Arguments</a> <span class="keyword">const</span> &args, </div><div class="line"><a name="l00697"></a><span class="lineno"> 697</span>  <span class="keywordtype">void</span> *workspace = <span class="keyword">nullptr</span>, </div><div class="line"><a name="l00698"></a><span class="lineno"> 698</span>  cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00699"></a><span class="lineno"> 699</span>  </div><div class="line"><a name="l00700"></a><span class="lineno"> 700</span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> status = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a53d79d1b434100da1e466e6378ec43ab">initialize</a>(args, workspace);</div><div class="line"><a name="l00701"></a><span class="lineno"> 701</span>  </div><div class="line"><a name="l00702"></a><span class="lineno"> 702</span>  <span class="keywordflow">if</span> (status == <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>) {</div><div class="line"><a name="l00703"></a><span class="lineno"> 703</span>  status = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aef8c133e539ef91efc9dba9012118147">run</a>(stream);</div><div class="line"><a name="l00704"></a><span class="lineno"> 704</span>  }</div><div class="line"><a name="l00705"></a><span class="lineno"> 705</span> </div><div class="line"><a name="l00706"></a><span class="lineno"> 706</span>  <span class="keywordflow">return</span> status;</div><div class="line"><a name="l00707"></a><span class="lineno"> 707</span>  }</div><div class="line"><a name="l00708"></a><span class="lineno"> 708</span> };</div><div class="line"><a name="l00709"></a><span class="lineno"> 709</span> </div><div class="line"><a name="l00711"></a><span class="lineno"> 711</span> </div><div class="line"><a name="l00712"></a><span class="lineno"> 712</span> } <span class="comment">// namespace device</span></div><div class="line"><a name="l00713"></a><span class="lineno"> 713</span> } <span class="comment">// namespace gemm</span></div><div class="line"><a name="l00714"></a><span class="lineno"> 714</span> } <span class="comment">// namespace cutlass</span></div><div class="line"><a name="l00715"></a><span class="lineno"> 715</span> </div><div class="ttc" id="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemm_html"><div class="ttname"><a href="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemm.html">cutlass::gemm::kernel::DefaultGemm</a></div><div class="ttdef"><b>Definition:</b> default_gemm.h:116</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a5cd83a90660626f6c446d45881f0fc22"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a5cd83a90660626f6c446d45881f0fc22">cutlass::gemm::device::Gemm::kStages</a></div><div class="ttdeci">static int const kStages</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:238</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments_html_ad68cd06dadc163a13f5ed29e07d6535b"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#ad68cd06dadc163a13f5ed29e07d6535b">cutlass::gemm::device::Gemm::Arguments::problem_size</a></div><div class="ttdeci">GemmCoord problem_size</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:276</div></div>
|
|
<div class="ttc" id="namespacecutlass_html"><div class="ttname"><a href="namespacecutlass.html">cutlass</a></div><div class="ttdef"><b>Definition:</b> aligned_buffer.h:35</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709_html_aaef8450711318fa1a53fe3cb72b59263"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#aaef8450711318fa1a53fe3cb72b59263">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::Arguments::split_k_slices</a></div><div class="ttdeci">int split_k_slices</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:606</div></div>
|
|
<div class="ttc" id="namespacecutlass_html_ac5a88c5840a28a9e0206b9cc7812a18da58c16acfe39125423ff22c907487ba2c"><div class="ttname"><a href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da58c16acfe39125423ff22c907487ba2c">cutlass::Status::kErrorInvalidProblem</a></div><div class="ttdoc">Specified problem size is not supported by operator. </div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_adf1ec9ab942ed20321cb91f02154abf0"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#adf1ec9ab942ed20321cb91f02154abf0">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::ElementA</a></div><div class="ttdeci">ElementB ElementA</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:219</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709_html_a9bdaf3563983efcca649460be169b334"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a9bdaf3563983efcca649460be169b334">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::Arguments::ref_A</a></div><div class="ttdeci">TensorRef< ElementA const, LayoutA > ref_A</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:601</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a0fec423a58e8de8ff7b015e5167ac614"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a0fec423a58e8de8ff7b015e5167ac614">cutlass::gemm::device::Gemm::get_workspace_size</a></div><div class="ttdeci">static size_t get_workspace_size(Arguments const &args)</div><div class="ttdoc">Gets the workspace size. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:350</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a1eda40e6a86fb3ebeabed2f717e47ced"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a1eda40e6a86fb3ebeabed2f717e47ced">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::ThreadblockSwizzle</a></div><div class="ttdeci">ThreadblockSwizzle ThreadblockSwizzle</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:236</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709_html_ac6c397a181a52c0dbb39bf3710ee4658"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#ac6c397a181a52c0dbb39bf3710ee4658">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::Arguments::Arguments</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Arguments()</div><div class="ttdoc">Default ctor. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:614</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a40ad889da7ff420fd9f9000cd9f98e32"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a40ad889da7ff420fd9f9000cd9f98e32">cutlass::gemm::device::Gemm::can_implement</a></div><div class="ttdeci">static Status can_implement(Arguments const &args)</div><div class="ttdoc">Determines whether the GEMM can execute the given problem. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:328</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments_html_a4b5c3e7a0f2307ca90632fa476aac1aa"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a4b5c3e7a0f2307ca90632fa476aac1aa">cutlass::gemm::device::Gemm::Arguments::Arguments</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Arguments()</div><div class="ttdoc">Default ctor. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:290</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments_html_a975d1d55ab7d42eb02eb226e2bbb5812"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a975d1d55ab7d42eb02eb226e2bbb5812">cutlass::gemm::device::Gemm::Arguments::Arguments</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Arguments(GemmCoord problem_size_, TensorRef< ElementA const, LayoutA > ref_A_, TensorRef< ElementB const, LayoutB > ref_B_, TensorRef< ElementC const, LayoutC > ref_C_, TensorRef< ElementC, LayoutC > ref_D_, typename EpilogueOutputOp::Params epilogue_=typename EpilogueOutputOp::Params(), int split_k_slices=1)</div><div class="ttdoc">Constructs an Arguments structure. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:296</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1GemmCoord_html"><div class="ttname"><a href="structcutlass_1_1gemm_1_1GemmCoord.html">cutlass::gemm::GemmCoord</a></div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/gemm.h:94</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_a1469133c30fde6b28296e3ff6951e7a4"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a1469133c30fde6b28296e3ff6951e7a4">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::get_workspace_size</a></div><div class="ttdeci">static size_t get_workspace_size(Arguments const &args)</div><div class="ttdoc">Gets the workspace size. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:666</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html">cutlass::gemm::device::Gemm</a></div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:216</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1GemmCoord_html_a1b29d2cb15360ad5499216859ad5436a"><div class="ttname"><a href="structcutlass_1_1gemm_1_1GemmCoord.html#a1b29d2cb15360ad5499216859ad5436a">cutlass::gemm::GemmCoord::n</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Index const & n() const </div><div class="ttdoc">Returns the GEMM N coordinate. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/gemm.h:137</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_af3c79a0271e684d93c3dca5ad230f45f"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af3c79a0271e684d93c3dca5ad230f45f">cutlass::gemm::device::Gemm::kSplitKSerial</a></div><div class="ttdeci">static bool const kSplitKSerial</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:242</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_a2b6c5275c173d73cffe8e6b6b1ccf2c1"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a2b6c5275c173d73cffe8e6b6b1ccf2c1">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::update</a></div><div class="ttdeci">Status update(Arguments const &args, void *workspace=nullptr)</div><div class="ttdoc">Lightweight update given a subset of arguments. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:678</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments_html_a98c7946bf72f054b026bb6fd49175e90"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a98c7946bf72f054b026bb6fd49175e90">cutlass::gemm::device::Gemm::Arguments::epilogue</a></div><div class="ttdeci">EpilogueOutputOp::Params epilogue</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:281</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709_html_a426f402c08be99849a4477a07f010a5e"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a426f402c08be99849a4477a07f010a5e">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::Arguments::epilogue</a></div><div class="ttdeci">EpilogueOutputOp::Params epilogue</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:605</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments_html_a390abae1ca0e01a4b6e58f3724b48eed"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a390abae1ca0e01a4b6e58f3724b48eed">cutlass::gemm::device::Gemm::Arguments::ref_A</a></div><div class="ttdeci">TensorRef< ElementA const, LayoutA > ref_A</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:277</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_abcacf502806db50eb17a6d925aee16d5"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#abcacf502806db50eb17a6d925aee16d5">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::Gemm</a></div><div class="ttdeci">Gemm()</div><div class="ttdoc">Constructs the GEMM. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:644</div></div>
|
|
<div class="ttc" id="classcutlass_1_1layout_1_1ColumnMajor_html"><div class="ttname"><a href="classcutlass_1_1layout_1_1ColumnMajor.html">cutlass::layout::ColumnMajor</a></div><div class="ttdoc">Mapping function for column-major matrices. </div><div class="ttdef"><b>Definition:</b> layout/matrix.h:142</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_a662bcbcb6164c803ab490c86e69b9ee1"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a662bcbcb6164c803ab490c86e69b9ee1">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::can_implement</a></div><div class="ttdeci">static Status can_implement(Arguments const &args)</div><div class="ttdoc">Determines whether the GEMM can execute the given problem. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:660</div></div>
|
|
<div class="ttc" id="default__gemm_8h_html"><div class="ttname"><a href="default__gemm_8h.html">default_gemm.h</a></div><div class="ttdoc">Default kernel-level GEMM definitions combine threadblock-scoped matrix multiply-add with the appropr...</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_a384db4125183e504fafc5a946b7ba757"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a384db4125183e504fafc5a946b7ba757">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::operator()</a></div><div class="ttdeci">Status operator()(cudaStream_t stream=nullptr)</div><div class="ttdoc">Runs the kernel using initialized state. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:690</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_ac98211b8adf5c18b4b6c54c5d1cdbb1a"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ac98211b8adf5c18b4b6c54c5d1cdbb1a">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::InstructionShape</a></div><div class="ttdeci">InstructionShape InstructionShape</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:234</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a6f8a0ff6be313d9f2aac2de4259f65b4"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a6f8a0ff6be313d9f2aac2de4259f65b4">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::Operator</a></div><div class="ttdeci">Operator Operator</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:237</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_aaaa871717d2fbe254a434160bc5acc65"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aaaa871717d2fbe254a434160bc5acc65">cutlass::gemm::device::Gemm::update</a></div><div class="ttdeci">Status update(Arguments const &args, void *workspace=nullptr)</div><div class="ttdoc">Lightweight update given a subset of arguments. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:417</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments_html_ac9c3c764f72e29c3aea99a8f3998e6cd"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#ac9c3c764f72e29c3aea99a8f3998e6cd">cutlass::gemm::device::Gemm::Arguments::split_k_slices</a></div><div class="ttdeci">int split_k_slices</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:282</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_af4c31798e982518db6a8fe66e63962e6"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af4c31798e982518db6a8fe66e63962e6">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::ElementC</a></div><div class="ttdeci">ElementC ElementC</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:225</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_a5f4f93ca97b358b4410f3d0b1e0a6387"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a5f4f93ca97b358b4410f3d0b1e0a6387">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::run</a></div><div class="ttdeci">Status run(cudaStream_t stream=nullptr)</div><div class="ttdoc">Runs the kernel using initialized state. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:684</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a114b9c14e102d333ef3dcad7865a4efb"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a114b9c14e102d333ef3dcad7865a4efb">cutlass::gemm::device::Gemm::operator()</a></div><div class="ttdeci">Status operator()(cudaStream_t stream=nullptr)</div><div class="ttdoc">Runs the kernel using initialized state. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:471</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a359e8eaddc486fcd0e15b879b11247ad"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a359e8eaddc486fcd0e15b879b11247ad">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::OperatorClass</a></div><div class="ttdeci">OperatorClass OperatorClass</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:230</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_a7615ad046304360243729c29c65e878a"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a7615ad046304360243729c29c65e878a">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::UnderlyingArguments</a></div><div class="ttdeci">typename UnderlyingOperator::Arguments UnderlyingArguments</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:589</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_ad35f92eaf14c268f8bda2a0e48399a9a"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ad35f92eaf14c268f8bda2a0e48399a9a">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::ElementAccumulator</a></div><div class="ttdeci">ElementAccumulator ElementAccumulator</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:229</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a431ea9cb851566df881457a0f496b976"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a431ea9cb851566df881457a0f496b976">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::GemmKernel</a></div><div class="ttdeci">typename kernel::DefaultGemm< ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, kStages, kSplitKSerial, Operator, kIsBetaZero >::GemmKernel GemmKernel</div><div class="ttdoc">Define the kernel. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:267</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a30809c4843226401b9d2da5ecb8d838b"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a30809c4843226401b9d2da5ecb8d838b">cutlass::gemm::device::Gemm::kAlignmentB</a></div><div class="ttdeci">static int const kAlignmentB</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:240</div></div>
|
|
<div class="ttc" id="structcutlass_1_1layout_1_1LayoutTranspose_html"><div class="ttname"><a href="structcutlass_1_1layout_1_1LayoutTranspose.html">cutlass::layout::LayoutTranspose</a></div><div class="ttdoc">Defines transposes of matrix layouts. </div><div class="ttdef"><b>Definition:</b> layout/matrix.h:921</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_ad0f1a2669ecc62f24922613e9bad857f"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ad0f1a2669ecc62f24922613e9bad857f">cutlass::gemm::device::Gemm::kAlignmentA</a></div><div class="ttdeci">static int const kAlignmentA</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:239</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a01ef9ea6588f488e8d727e36bdec8ba8"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a01ef9ea6588f488e8d727e36bdec8ba8">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::ThreadblockShape</a></div><div class="ttdeci">ThreadblockShape ThreadblockShape</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:232</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments_html_abdba57a68d6982fffbb1cc3db34ef0f9"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#abdba57a68d6982fffbb1cc3db34ef0f9">cutlass::gemm::device::Gemm::Arguments::ref_C</a></div><div class="ttdeci">TensorRef< ElementC const, LayoutC > ref_C</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:279</div></div>
|
|
<div class="ttc" id="classcutlass_1_1TensorRef_html"><div class="ttname"><a href="classcutlass_1_1TensorRef.html">cutlass::TensorRef< ElementA const, LayoutA ></a></div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments_html_ae712c362f83fbd45679a6e989315d3dc"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#ae712c362f83fbd45679a6e989315d3dc">cutlass::gemm::device::Gemm::Arguments::ref_B</a></div><div class="ttdeci">TensorRef< ElementB const, LayoutB > ref_B</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:278</div></div>
|
|
<div class="ttc" id="namespacecutlass_html_ac5a88c5840a28a9e0206b9cc7812a18da4d3b5847a0b17037c6b69faf8b1d4d71"><div class="ttname"><a href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da4d3b5847a0b17037c6b69faf8b1d4d71">cutlass::Status::kErrorInternal</a></div><div class="ttdoc">An error within CUTLASS occurred. </div></div>
|
|
<div class="ttc" id="device__kernel_8h_html"><div class="ttname"><a href="device__kernel_8h.html">device_kernel.h</a></div><div class="ttdoc">Template for generic CUTLASS kernel. </div></div>
|
|
<div class="ttc" id="cutlass_8h_html_a28c2443a142676d3d71effdae1a986b1"><div class="ttname"><a href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="ttdeci">#define CUTLASS_HOST_DEVICE</div><div class="ttdef"><b>Definition:</b> cutlass.h:89</div></div>
|
|
<div class="ttc" id="numeric__types_8h_html"><div class="ttname"><a href="numeric__types_8h.html">numeric_types.h</a></div><div class="ttdoc">Top-level include for all CUTLASS numeric types. </div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709_html_acd02e86dfff866eade08415e0043ccc3"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#acd02e86dfff866eade08415e0043ccc3">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::Arguments::problem_size</a></div><div class="ttdeci">GemmCoord problem_size</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:600</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a35d9f803fdfbbd4608243881c04316dc"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a35d9f803fdfbbd4608243881c04316dc">cutlass::gemm::device::Gemm::operator()</a></div><div class="ttdeci">Status operator()(Arguments const &args, void *workspace=nullptr, cudaStream_t stream=nullptr)</div><div class="ttdoc">Runs the kernel using initialized state. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:476</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_aec04d65c6265eb5f63d703f2dd99cb3f"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aec04d65c6265eb5f63d703f2dd99cb3f">cutlass::gemm::device::Gemm::Gemm</a></div><div class="ttdeci">Gemm()</div><div class="ttdoc">Constructs the GEMM. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:325</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709_html_a331de1adfdcbea6d0137afe64a4f6f4c"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a331de1adfdcbea6d0137afe64a4f6f4c">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::Arguments::Arguments</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Arguments(GemmCoord problem_size_, TensorRef< ElementA const, LayoutA > ref_A_, TensorRef< ElementB const, LayoutB > ref_B_, TensorRef< ElementC const, LayoutC > ref_C_, TensorRef< ElementC, LayoutC > ref_D_, typename EpilogueOutputOp::Params epilogue_=typename EpilogueOutputOp::Params(), int split_k_slices=1)</div><div class="ttdoc">Constructs an Arguments structure. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:618</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a93acaff17a5916bc11040a9ddf04d9f2"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a93acaff17a5916bc11040a9ddf04d9f2">cutlass::gemm::device::Gemm::LayoutC</a></div><div class="ttdeci">LayoutC_ LayoutC</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:226</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_ab389fd311c5763587315c210f40a11f2"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ab389fd311c5763587315c210f40a11f2">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::LayoutB</a></div><div class="ttdeci">typename layout::LayoutTranspose< LayoutA >::type LayoutB</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:223</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709_html_a590b8da88ae9350042838451e3e37a22"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#a590b8da88ae9350042838451e3e37a22">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::Arguments::ref_C</a></div><div class="ttdeci">TensorRef< ElementC const, LayoutC > ref_C</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:603</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments_html"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html">cutlass::gemm::device::Gemm::Arguments</a></div><div class="ttdoc">Argument structure. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:270</div></div>
|
|
<div class="ttc" id="default__gemm__configuration_8h_html"><div class="ttname"><a href="default__gemm__configuration_8h.html">default_gemm_configuration.h</a></div><div class="ttdoc">Definitions for GEMM structures. </div></div>
|
|
<div class="ttc" id="classcutlass_1_1layout_1_1RowMajor_html"><div class="ttname"><a href="classcutlass_1_1layout_1_1RowMajor.html">cutlass::layout::RowMajor</a></div><div class="ttdoc">Mapping function for row-major matrices. </div><div class="ttdef"><b>Definition:</b> layout/matrix.h:50</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_a7a14474e4238d2fac92ad71c6de087d8"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a7a14474e4238d2fac92ad71c6de087d8">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::initialize</a></div><div class="ttdeci">Status initialize(Arguments const &args, void *workspace=nullptr, cudaStream_t stream=nullptr)</div><div class="ttdoc">Initializes GEMM state from arguments. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:672</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_aef8c133e539ef91efc9dba9012118147"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#aef8c133e539ef91efc9dba9012118147">cutlass::gemm::device::Gemm::run</a></div><div class="ttdeci">Status run(cudaStream_t stream=nullptr)</div><div class="ttdoc">Runs the kernel using initialized state. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:435</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_af34b612740d98d3066b2b14785f9e415"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#af34b612740d98d3066b2b14785f9e415">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::EpilogueOutputOp</a></div><div class="ttdeci">EpilogueOutputOp EpilogueOutputOp</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:235</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a53d79d1b434100da1e466e6378ec43ab"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a53d79d1b434100da1e466e6378ec43ab">cutlass::gemm::device::Gemm::initialize</a></div><div class="ttdeci">Status initialize(Arguments const &args, void *workspace=nullptr, cudaStream_t stream=nullptr)</div><div class="ttdoc">Initializes GEMM state from arguments. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:369</div></div>
|
|
<div class="ttc" id="namespacecutlass_html_ac5a88c5840a28a9e0206b9cc7812a18da6467125354474c5cff2f3c920abb7d3f"><div class="ttname"><a href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da6467125354474c5cff2f3c920abb7d3f">cutlass::Status::kErrorWorkspaceNull</a></div><div class="ttdoc">The given workspace is null when it is required to be non-null. </div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_adb5bd27efee1575710e6e4157e32449d"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#adb5bd27efee1575710e6e4157e32449d">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::WarpShape</a></div><div class="ttdeci">WarpShape WarpShape</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:233</div></div>
|
|
<div class="ttc" id="namespacecutlass_html_ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782"><div class="ttname"><a href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">cutlass::Status::kSuccess</a></div><div class="ttdoc">Operation was successful. </div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_ac85244d4a91ea8e8c1c3187fad1004db"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#ac85244d4a91ea8e8c1c3187fad1004db">cutlass::gemm::device::Gemm::kAlignmentC</a></div><div class="ttdeci">static int const kAlignmentC</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:241</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_aa9313915a6129f0c43b43ef3698b3ee4"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#aa9313915a6129f0c43b43ef3698b3ee4">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::to_underlying_arguments</a></div><div class="ttdeci">static UnderlyingArguments to_underlying_arguments(Arguments const &args)</div><div class="ttdoc">Helper to construct a transposed equivalent for the underying GEMM operator. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:647</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1GemmCoord_html_a93515a41db6c4b7e9101067f60d41b8c"><div class="ttname"><a href="structcutlass_1_1gemm_1_1GemmCoord.html#a93515a41db6c4b7e9101067f60d41b8c">cutlass::gemm::GemmCoord::m</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Index const & m() const </div><div class="ttdoc">Returns the GEMM M coordinate. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/gemm.h:129</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a7d04e3640e81787a28956c75acde997e"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a7d04e3640e81787a28956c75acde997e">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::ElementB</a></div><div class="ttdeci">ElementA ElementB</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:222</div></div>
|
|
<div class="ttc" id="gemm_2threadblock_2threadblock__swizzle_8h_html"><div class="ttname"><a href="gemm_2threadblock_2threadblock__swizzle_8h.html">threadblock_swizzle.h</a></div><div class="ttdoc">Implements several possible threadblock-swizzling functions mapping blockIdx to GEMM problems...</div></div>
|
|
<div class="ttc" id="arch_8h_html"><div class="ttname"><a href="arch_8h.html">arch.h</a></div><div class="ttdoc">Defines tags for architecture-specific configurations. </div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments_html_a66f9983db4a09ac0d90291c0f8723897"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html#a66f9983db4a09ac0d90291c0f8723897">cutlass::gemm::device::Gemm::Arguments::ref_D</a></div><div class="ttdeci">TensorRef< ElementC, LayoutC > ref_D</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:280</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709_html_ab1d4d5865786a415f87db1def1b029e7"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#ab1d4d5865786a415f87db1def1b029e7">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::Arguments::ref_D</a></div><div class="ttdeci">TensorRef< ElementC, LayoutC > ref_D</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:604</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a0303e85432228170eba7dc4b418c86b4"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a0303e85432228170eba7dc4b418c86b4">cutlass::gemm::device::Gemm::kIsBetaZero</a></div><div class="ttdeci">static bool const kIsBetaZero</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:243</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_a2bdbc5e737f9bfd1e09a7cfb30e60e29"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a2bdbc5e737f9bfd1e09a7cfb30e60e29">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::GemmKernel</a></div><div class="ttdeci">typename UnderlyingOperator::GemmKernel GemmKernel</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:590</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_afe7ab8c15e83c6cd59b6bcf3fe6e48c0"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#afe7ab8c15e83c6cd59b6bcf3fe6e48c0">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::ArchTag</a></div><div class="ttdeci">ArchTag ArchTag</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:231</div></div>
|
|
<div class="ttc" id="cutlass_8h_html"><div class="ttname"><a href="cutlass_8h.html">cutlass.h</a></div><div class="ttdoc">Basic include for CUTLASS. </div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c_html_a6115aa957b3ba8ad9e54b7efeefaacd1"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html#a6115aa957b3ba8ad9e54b7efeefaacd1">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::operator()</a></div><div class="ttdeci">Status operator()(Arguments const &args, void *workspace=nullptr, cudaStream_t stream=nullptr)</div><div class="ttdoc">Runs the kernel using initialized state. </div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:695</div></div>
|
|
<div class="ttc" id="namespacecutlass_html_ac5a88c5840a28a9e0206b9cc7812a18d"><div class="ttname"><a href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">cutlass::Status</a></div><div class="ttdeci">Status</div><div class="ttdoc">Status code returned by CUTLASS operations. </div><div class="ttdef"><b>Definition:</b> cutlass.h:39</div></div>
|
|
<div class="ttc" id="include_2cutlass_2gemm_2kernel_2gemm_8h_html"><div class="ttname"><a href="include_2cutlass_2gemm_2kernel_2gemm_8h.html">gemm.h</a></div><div class="ttdoc">Template for a pipelined GEMM kernel. Does not compute batching or support split-K. </div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709_html_ab77204c1010b17c6643d26a89f41c3d0"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html#ab77204c1010b17c6643d26a89f41c3d0">cutlass::gemm::device::Gemm< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero >::Arguments::ref_B</a></div><div class="ttdeci">TensorRef< ElementB const, LayoutB > ref_B</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:602</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1Gemm_html_a6bbefc44e1ed6822dc8a4e2edc2b276b"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1Gemm.html#a6bbefc44e1ed6822dc8a4e2edc2b276b">cutlass::gemm::device::Gemm< ElementB, typename layout::LayoutTranspose< LayoutB >::type, ElementA, typename layout::LayoutTranspose< LayoutA >::type, ElementC, layout::RowMajor, ElementAccumulator, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, Stages, kAlignmentB, kAlignmentA, SplitKSerial, Operator, kIsBetaZero >::LayoutA</a></div><div class="ttdeci">typename layout::LayoutTranspose< LayoutB >::type LayoutA</div><div class="ttdef"><b>Definition:</b> include/cutlass/gemm/device/gemm.h:220</div></div>
|
|
</div><!-- fragment --></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>
|