222 lines
180 KiB
HTML
222 lines
180 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_batched.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">device/gemm_batched.h</div> </div>
|
|
</div><!--header-->
|
|
<div class="contents">
|
|
<a href="device_2gemm__batched_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="kernel_2gemm__batched_8h.html">cutlass/gemm/kernel/gemm_batched.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::GemmBatchedIdentityThreadblockSwizzle,</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="keyword">typename</span> Operator_ = <span class="keyword">typename</span> DefaultGemmConfiguration<</div><div class="line"><a name="l00210"></a><span class="lineno"> 210</span>  OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_,</div><div class="line"><a name="l00211"></a><span class="lineno"> 211</span>  ElementAccumulator_>::Operator</div><div class="line"><a name="l00212"></a><span class="lineno"> 212</span> ></div><div class="line"><a name="l00213"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html"> 213</a></span> <span class="keyword">class </span><a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html">GemmBatched</a> {</div><div class="line"><a name="l00214"></a><span class="lineno"> 214</span>  <span class="keyword">public</span>:</div><div class="line"><a name="l00215"></a><span class="lineno"> 215</span> </div><div class="line"><a name="l00216"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a3fe2fcad97f15d63fa1a0214ef4861f2"> 216</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a3fe2fcad97f15d63fa1a0214ef4861f2">ElementA</a> = ElementA_;</div><div class="line"><a name="l00217"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a9f9d11529b28ced91c4b05c2530b7a70"> 217</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a9f9d11529b28ced91c4b05c2530b7a70">LayoutA</a> = LayoutA_;</div><div class="line"><a name="l00218"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a044074cbb894d8d184c72074ff3a3bf4"> 218</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="l00219"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac8bb1360bbc57bc63296cd48005c3c42"> 219</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac8bb1360bbc57bc63296cd48005c3c42">ElementB</a> = ElementB_;</div><div class="line"><a name="l00220"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#acb489f1bb4fed9e4314a6b6a3cbd04a9"> 220</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#acb489f1bb4fed9e4314a6b6a3cbd04a9">LayoutB</a> = LayoutB_;</div><div class="line"><a name="l00221"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a533fd90ffa10f464b1c1aa842c82bd26"> 221</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="l00222"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a008028eec1dd4b8c08128c6dfe44cce5"> 222</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a008028eec1dd4b8c08128c6dfe44cce5">ElementC</a> = ElementC_;</div><div class="line"><a name="l00223"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#af35efd1f40deeb9d8e295f700fa84dbd"> 223</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="l00224"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aeec2c03850540947cefb28cc90a293e7"> 224</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="l00225"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#af5092c2505a27f3e4160ff16046a1c33"> 225</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="l00226"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a6e4cff55a6834d43cfbc97df40609eea"> 226</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a6e4cff55a6834d43cfbc97df40609eea">ElementAccumulator</a> = ElementAccumulator_;</div><div class="line"><a name="l00227"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a24935d746d97b0c994c9a9ade820d2d0"> 227</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a24935d746d97b0c994c9a9ade820d2d0">OperatorClass</a> = OperatorClass_;</div><div class="line"><a name="l00228"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#add5d679e0acf0813a52c209d2448e81b"> 228</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#add5d679e0acf0813a52c209d2448e81b">ArchTag</a> = ArchTag_;</div><div class="line"><a name="l00229"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a74aece33b6fafe58db1b41a6b7b87729"> 229</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a74aece33b6fafe58db1b41a6b7b87729">ThreadblockShape</a> = ThreadblockShape_;</div><div class="line"><a name="l00230"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a7b8d2dbfa562869deb58c88583951b58"> 230</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a7b8d2dbfa562869deb58c88583951b58">WarpShape</a> = WarpShape_;</div><div class="line"><a name="l00231"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a0dbb6d5185f223bb8242fc47a3b77757"> 231</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a0dbb6d5185f223bb8242fc47a3b77757">InstructionShape</a> = InstructionShape_;</div><div class="line"><a name="l00232"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2d43ad49fa686ded524cc5f26b36c69"> 232</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2d43ad49fa686ded524cc5f26b36c69">EpilogueOutputOp</a> = EpilogueOutputOp_;</div><div class="line"><a name="l00233"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a1b685fd66f6dc2c572be067ef1396a89"> 233</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a1b685fd66f6dc2c572be067ef1396a89">ThreadblockSwizzle</a> = ThreadblockSwizzle_;</div><div class="line"><a name="l00234"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a73837bda9ba209e546f6d996ede1afad"> 234</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_1GemmBatched.html#a73837bda9ba209e546f6d996ede1afad">kStages</a> = Stages;</div><div class="line"><a name="l00235"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a8976ed5c5e404ee87deaea4455d0d960"> 235</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_1GemmBatched.html#a8976ed5c5e404ee87deaea4455d0d960">kAlignmentA</a> = AlignmentA;</div><div class="line"><a name="l00236"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a17d34b2884711522fafcfd7c7500955c"> 236</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_1GemmBatched.html#a17d34b2884711522fafcfd7c7500955c">kAlignmentB</a> = AlignmentB;</div><div class="line"><a name="l00237"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac1f190321a811fa91eec0096829b07ff"> 237</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_1GemmBatched.html#ac1f190321a811fa91eec0096829b07ff">kAlignmentC</a> = EpilogueOutputOp::kCount;</div><div class="line"><a name="l00238"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a52b7263c5c86e900bcca681d07f19101"> 238</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a52b7263c5c86e900bcca681d07f19101">Operator</a> = Operator_;</div><div class="line"><a name="l00239"></a><span class="lineno"> 239</span> </div><div class="line"><a name="l00241"></a><span class="lineno"> 241</span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a6acd50cfc477e95dbcf0d4fbba5df65c">DefaultGemmKernel</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="l00242"></a><span class="lineno"> 242</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a3fe2fcad97f15d63fa1a0214ef4861f2">ElementA</a>,</div><div class="line"><a name="l00243"></a><span class="lineno"> 243</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a9f9d11529b28ced91c4b05c2530b7a70">LayoutA</a>,</div><div class="line"><a name="l00244"></a><span class="lineno"> 244</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a8976ed5c5e404ee87deaea4455d0d960">kAlignmentA</a>,</div><div class="line"><a name="l00245"></a><span class="lineno"> 245</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac8bb1360bbc57bc63296cd48005c3c42">ElementB</a>,</div><div class="line"><a name="l00246"></a><span class="lineno"> 246</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#acb489f1bb4fed9e4314a6b6a3cbd04a9">LayoutB</a>,</div><div class="line"><a name="l00247"></a><span class="lineno"> 247</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a17d34b2884711522fafcfd7c7500955c">kAlignmentB</a>,</div><div class="line"><a name="l00248"></a><span class="lineno"> 248</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a008028eec1dd4b8c08128c6dfe44cce5">ElementC</a>,</div><div class="line"><a name="l00249"></a><span class="lineno"> 249</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#af35efd1f40deeb9d8e295f700fa84dbd">LayoutC</a>,</div><div class="line"><a name="l00250"></a><span class="lineno"> 250</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a6e4cff55a6834d43cfbc97df40609eea">ElementAccumulator</a>,</div><div class="line"><a name="l00251"></a><span class="lineno"> 251</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a24935d746d97b0c994c9a9ade820d2d0">OperatorClass</a>,</div><div class="line"><a name="l00252"></a><span class="lineno"> 252</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#add5d679e0acf0813a52c209d2448e81b">ArchTag</a>,</div><div class="line"><a name="l00253"></a><span class="lineno"> 253</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a74aece33b6fafe58db1b41a6b7b87729">ThreadblockShape</a>,</div><div class="line"><a name="l00254"></a><span class="lineno"> 254</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a7b8d2dbfa562869deb58c88583951b58">WarpShape</a>,</div><div class="line"><a name="l00255"></a><span class="lineno"> 255</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a0dbb6d5185f223bb8242fc47a3b77757">InstructionShape</a>,</div><div class="line"><a name="l00256"></a><span class="lineno"> 256</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2d43ad49fa686ded524cc5f26b36c69">EpilogueOutputOp</a>,</div><div class="line"><a name="l00257"></a><span class="lineno"> 257</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a1b685fd66f6dc2c572be067ef1396a89">ThreadblockSwizzle</a>,</div><div class="line"><a name="l00258"></a><span class="lineno"> 258</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a73837bda9ba209e546f6d996ede1afad">kStages</a>,</div><div class="line"><a name="l00259"></a><span class="lineno"> 259</span>  <span class="keyword">false</span>,</div><div class="line"><a name="l00260"></a><span class="lineno"> 260</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a52b7263c5c86e900bcca681d07f19101">Operator</a>,</div><div class="line"><a name="l00261"></a><span class="lineno"> 261</span>  <span class="keyword">false</span></div><div class="line"><a name="l00262"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a6acd50cfc477e95dbcf0d4fbba5df65c"> 262</a></span>  ><a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a35a2fb2e9ad63c316ac6fbb1cc8cf53a">::GemmKernel</a>;</div><div class="line"><a name="l00263"></a><span class="lineno"> 263</span> </div><div class="line"><a name="l00264"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a35a2fb2e9ad63c316ac6fbb1cc8cf53a"> 264</a></span>  <span class="keyword">using</span> <a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched.html">GemmKernel</a> = <a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched.html">kernel::GemmBatched<typename DefaultGemmKernel::Mma, typename DefaultGemmKernel::Epilogue, ThreadblockSwizzle></a>;</div><div class="line"><a name="l00265"></a><span class="lineno"> 265</span> </div><div class="line"><a name="l00267"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html"> 267</a></span>  <span class="keyword">struct </span><a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html">Arguments</a> {</div><div class="line"><a name="l00268"></a><span class="lineno"> 268</span> </div><div class="line"><a name="l00269"></a><span class="lineno"> 269</span>  <span class="comment">//</span></div><div class="line"><a name="l00270"></a><span class="lineno"> 270</span>  <span class="comment">// Data members</span></div><div class="line"><a name="l00271"></a><span class="lineno"> 271</span>  <span class="comment">//</span></div><div class="line"><a name="l00272"></a><span class="lineno"> 272</span> </div><div class="line"><a name="l00273"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#aab4fc258e38ebcf9b430a5dee6daba5e"> 273</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html">GemmCoord</a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#aab4fc258e38ebcf9b430a5dee6daba5e">problem_size</a>;</div><div class="line"><a name="l00274"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a55f32be45559dbf84dcc2db26784f625"> 274</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_1GemmBatched_1_1Arguments.html#a55f32be45559dbf84dcc2db26784f625">ref_A</a>;</div><div class="line"><a name="l00275"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a11ef91161a92459d72b56144cd6b4495"> 275</a></span>  int64_t <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a11ef91161a92459d72b56144cd6b4495">stride_A</a>;</div><div class="line"><a name="l00276"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a48844293c34b9c44fe57f577370664ea"> 276</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_1GemmBatched_1_1Arguments.html#a48844293c34b9c44fe57f577370664ea">ref_B</a>;</div><div class="line"><a name="l00277"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#aa867aa186538d34251d75ccc891453d7"> 277</a></span>  int64_t <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#aa867aa186538d34251d75ccc891453d7">stride_B</a>;</div><div class="line"><a name="l00278"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#ab0955b722ad4ea0217f725e34b3bcfbe"> 278</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_1GemmBatched_1_1Arguments.html#ab0955b722ad4ea0217f725e34b3bcfbe">ref_C</a>;</div><div class="line"><a name="l00279"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a3ce1385631b05430fa5dfc1e9a3671b8"> 279</a></span>  int64_t <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a3ce1385631b05430fa5dfc1e9a3671b8">stride_C</a>;</div><div class="line"><a name="l00280"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#ae4450f06a6975191d94026865e445578"> 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_1GemmBatched_1_1Arguments.html#ae4450f06a6975191d94026865e445578">ref_D</a>;</div><div class="line"><a name="l00281"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a05a1d9720fbb16a20b94049900b0d04f"> 281</a></span>  int64_t <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a05a1d9720fbb16a20b94049900b0d04f">stride_D</a>;</div><div class="line"><a name="l00282"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#aea0c4cd59daee8d3b497be411beb9b3a"> 282</a></span>  <span class="keyword">typename</span> EpilogueOutputOp::Params <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#aea0c4cd59daee8d3b497be411beb9b3a">epilogue</a>;</div><div class="line"><a name="l00283"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#ac99ca8f9d8a0053e647a6c99b018bda5"> 283</a></span>  <span class="keywordtype">int</span> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#ac99ca8f9d8a0053e647a6c99b018bda5">batch_count</a>;</div><div class="line"><a name="l00284"></a><span class="lineno"> 284</span> </div><div class="line"><a name="l00285"></a><span class="lineno"> 285</span>  <span class="comment">//</span></div><div class="line"><a name="l00286"></a><span class="lineno"> 286</span>  <span class="comment">// Methods</span></div><div class="line"><a name="l00287"></a><span class="lineno"> 287</span>  <span class="comment">//</span></div><div class="line"><a name="l00288"></a><span class="lineno"> 288</span> </div><div class="line"><a name="l00290"></a><span class="lineno"> 290</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00291"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5"> 291</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">Arguments</a>() { }</div><div class="line"><a name="l00292"></a><span class="lineno"> 292</span> </div><div class="line"><a name="l00294"></a><span class="lineno"> 294</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00295"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a0befb9945aadcba460f4d1ad73020e9c"> 295</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a0befb9945aadcba460f4d1ad73020e9c">Arguments</a>(</div><div class="line"><a name="l00296"></a><span class="lineno"> 296</span>  <a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html">GemmCoord</a> problem_size_,</div><div class="line"><a name="l00297"></a><span class="lineno"> 297</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementA const, LayoutA></a> ref_A_,</div><div class="line"><a name="l00298"></a><span class="lineno"> 298</span>  int64_t stride_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>  int64_t stride_B_,</div><div class="line"><a name="l00301"></a><span class="lineno"> 301</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC const, LayoutC></a> ref_C_,</div><div class="line"><a name="l00302"></a><span class="lineno"> 302</span>  int64_t stride_C_,</div><div class="line"><a name="l00303"></a><span class="lineno"> 303</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC, LayoutC></a> ref_D_,</div><div class="line"><a name="l00304"></a><span class="lineno"> 304</span>  int64_t stride_D_,</div><div class="line"><a name="l00305"></a><span class="lineno"> 305</span>  <span class="keyword">typename</span> EpilogueOutputOp::Params epilogue_,</div><div class="line"><a name="l00306"></a><span class="lineno"> 306</span>  <span class="keywordtype">int</span> batch_count_</div><div class="line"><a name="l00307"></a><span class="lineno"> 307</span>  ):</div><div class="line"><a name="l00308"></a><span class="lineno"> 308</span>  problem_size(problem_size_),</div><div class="line"><a name="l00309"></a><span class="lineno"> 309</span>  ref_A(ref_A_),</div><div class="line"><a name="l00310"></a><span class="lineno"> 310</span>  stride_A(stride_A_),</div><div class="line"><a name="l00311"></a><span class="lineno"> 311</span>  ref_B(ref_B_),</div><div class="line"><a name="l00312"></a><span class="lineno"> 312</span>  stride_B(stride_B_),</div><div class="line"><a name="l00313"></a><span class="lineno"> 313</span>  ref_C(ref_C_),</div><div class="line"><a name="l00314"></a><span class="lineno"> 314</span>  stride_C(stride_C_),</div><div class="line"><a name="l00315"></a><span class="lineno"> 315</span>  ref_D(ref_D_),</div><div class="line"><a name="l00316"></a><span class="lineno"> 316</span>  stride_D(stride_D_),</div><div class="line"><a name="l00317"></a><span class="lineno"> 317</span>  epilogue(epilogue_),</div><div class="line"><a name="l00318"></a><span class="lineno"> 318</span>  batch_count(batch_count_) { }</div><div class="line"><a name="l00319"></a><span class="lineno"> 319</span>  };</div><div class="line"><a name="l00320"></a><span class="lineno"> 320</span> </div><div class="line"><a name="l00321"></a><span class="lineno"> 321</span> <span class="keyword">private</span>:</div><div class="line"><a name="l00322"></a><span class="lineno"> 322</span> </div><div class="line"><a name="l00324"></a><span class="lineno"> 324</span>  <span class="keyword">typename</span> <a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html">GemmKernel::Params</a> params_;</div><div class="line"><a name="l00325"></a><span class="lineno"> 325</span> </div><div class="line"><a name="l00326"></a><span class="lineno"> 326</span> <span class="keyword">public</span>:</div><div class="line"><a name="l00327"></a><span class="lineno"> 327</span> </div><div class="line"><a name="l00329"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a72a26fb286181aa5ca1fb66d9b385f7f"> 329</a></span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a72a26fb286181aa5ca1fb66d9b385f7f">GemmBatched</a>() { }</div><div class="line"><a name="l00330"></a><span class="lineno"> 330</span> </div><div class="line"><a name="l00332"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#acb4d53fbea4366349574091d68594558"> 332</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_1GemmBatched.html#acb4d53fbea4366349574091d68594558">can_implement</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">Arguments</a> <span class="keyword">const</span> &args) {</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>  <span class="keywordflow">if</span> (!<a class="code" href="namespacecutlass.html#aa43b0a7d59635cb2d9ac96a077c988c3">TensorRef_aligned</a>(args.ref_A, kAlignmentA) || (args.stride_A % kAlignmentA)) {</div><div class="line"><a name="l00335"></a><span class="lineno"> 335</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18daa4867e1466f5d067dbec566abfe5a67a">Status::kErrorMisalignedOperand</a>;</div><div class="line"><a name="l00336"></a><span class="lineno"> 336</span>  }</div><div class="line"><a name="l00337"></a><span class="lineno"> 337</span> </div><div class="line"><a name="l00338"></a><span class="lineno"> 338</span>  <span class="keywordflow">if</span> (!<a class="code" href="namespacecutlass.html#aa43b0a7d59635cb2d9ac96a077c988c3">TensorRef_aligned</a>(args.ref_B, kAlignmentB) || (args.stride_B % kAlignmentB)) {</div><div class="line"><a name="l00339"></a><span class="lineno"> 339</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18daa4867e1466f5d067dbec566abfe5a67a">Status::kErrorMisalignedOperand</a>;</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> (!<a class="code" href="namespacecutlass.html#aa43b0a7d59635cb2d9ac96a077c988c3">TensorRef_aligned</a>(args.ref_C, kAlignmentC) || (args.stride_C % kAlignmentC)) {</div><div class="line"><a name="l00343"></a><span class="lineno"> 343</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18daa4867e1466f5d067dbec566abfe5a67a">Status::kErrorMisalignedOperand</a>;</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">if</span> (!<a class="code" href="namespacecutlass.html#aa43b0a7d59635cb2d9ac96a077c988c3">TensorRef_aligned</a>(args.ref_D, kAlignmentC) || (args.stride_D % kAlignmentC)) {</div><div class="line"><a name="l00347"></a><span class="lineno"> 347</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18daa4867e1466f5d067dbec566abfe5a67a">Status::kErrorMisalignedOperand</a>;</div><div class="line"><a name="l00348"></a><span class="lineno"> 348</span>  }</div><div class="line"><a name="l00349"></a><span class="lineno"> 349</span> </div><div class="line"><a name="l00350"></a><span class="lineno"> 350</span>  <span class="keywordflow">if</span> ((args.problem_size.m() % <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a8976ed5c5e404ee87deaea4455d0d960">kAlignmentA</a>) || (args.problem_size.k() % <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a8976ed5c5e404ee87deaea4455d0d960">kAlignmentA</a>) ||</div><div class="line"><a name="l00351"></a><span class="lineno"> 351</span>  (args.problem_size.n() % <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a17d34b2884711522fafcfd7c7500955c">kAlignmentB</a>) || (args.problem_size.k() % <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a17d34b2884711522fafcfd7c7500955c">kAlignmentB</a>) ||</div><div class="line"><a name="l00352"></a><span class="lineno"> 352</span>  (args.problem_size.m() % <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac1f190321a811fa91eec0096829b07ff">kAlignmentC</a>) || (args.problem_size.n() % <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac1f190321a811fa91eec0096829b07ff">kAlignmentC</a>)) {</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="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18daa4867e1466f5d067dbec566abfe5a67a">Status::kErrorMisalignedOperand</a>;</div><div class="line"><a name="l00355"></a><span class="lineno"> 355</span>  }</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>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>;</div><div class="line"><a name="l00358"></a><span class="lineno"> 358</span>  }</div><div class="line"><a name="l00359"></a><span class="lineno"> 359</span> </div><div class="line"><a name="l00361"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac2009bb52372115624aa5c4f75b720e5"> 361</a></span>  <span class="keyword">static</span> <span class="keywordtype">size_t</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac2009bb52372115624aa5c4f75b720e5">get_workspace_size</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">Arguments</a> <span class="keyword">const</span> &args) {</div><div class="line"><a name="l00362"></a><span class="lineno"> 362</span>  <span class="keywordflow">return</span> 0;</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="l00366"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2670ac441f48f6a0a2071c67c743ab8"> 366</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2670ac441f48f6a0a2071c67c743ab8">initialize</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">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="l00367"></a><span class="lineno"> 367</span> </div><div class="line"><a name="l00368"></a><span class="lineno"> 368</span>  <span class="comment">// Determine grid shape</span></div><div class="line"><a name="l00369"></a><span class="lineno"> 369</span>  ThreadblockSwizzle threadblock_swizzle;</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>  <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="l00372"></a><span class="lineno"> 372</span>  args.problem_size,</div><div class="line"><a name="l00373"></a><span class="lineno"> 373</span>  args.batch_count,</div><div class="line"><a name="l00374"></a><span class="lineno"> 374</span>  {ThreadblockShape::kM, ThreadblockShape::kN, ThreadblockShape::kK});</div><div class="line"><a name="l00375"></a><span class="lineno"> 375</span> </div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span>  <span class="comment">// Initialize the Params structure</span></div><div class="line"><a name="l00377"></a><span class="lineno"> 377</span>  params_ = <span class="keyword">typename</span> <a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html">GemmKernel::Params</a>{</div><div class="line"><a name="l00378"></a><span class="lineno"> 378</span>  args.<a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#a05909ba49e633c7eeb0707166c72a4ee">problem_size</a>,</div><div class="line"><a name="l00379"></a><span class="lineno"> 379</span>  grid_shape,</div><div class="line"><a name="l00380"></a><span class="lineno"> 380</span>  args.ref_A.non_const_ref(),</div><div class="line"><a name="l00381"></a><span class="lineno"> 381</span>  args.stride_A,</div><div class="line"><a name="l00382"></a><span class="lineno"> 382</span>  args.ref_B.non_const_ref(),</div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>  args.stride_B,</div><div class="line"><a name="l00384"></a><span class="lineno"> 384</span>  args.ref_C.non_const_ref(),</div><div class="line"><a name="l00385"></a><span class="lineno"> 385</span>  args.stride_C,</div><div class="line"><a name="l00386"></a><span class="lineno"> 386</span>  args.ref_D,</div><div class="line"><a name="l00387"></a><span class="lineno"> 387</span>  args.stride_D,</div><div class="line"><a name="l00388"></a><span class="lineno"> 388</span>  args.epilogue,</div><div class="line"><a name="l00389"></a><span class="lineno"> 389</span>  args.batch_count</div><div class="line"><a name="l00390"></a><span class="lineno"> 390</span>  };</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>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>;</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> </div><div class="line"><a name="l00396"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ab7c3e9a33a1c62513ec6eee3e2598df6"> 396</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ab7c3e9a33a1c62513ec6eee3e2598df6">update</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">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="l00397"></a><span class="lineno"> 397</span> </div><div class="line"><a name="l00398"></a><span class="lineno"> 398</span>  params_.<a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#ad1867c0875c10e6327c7fae16acd35a3">ref_A</a>.reset(args.ref_A.non_const_ref().data());</div><div class="line"><a name="l00399"></a><span class="lineno"> 399</span>  params_.<a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#ade55adc311c5561efe76f53ffd56d1f4">ref_B</a>.reset(args.ref_B.non_const_ref().data());</div><div class="line"><a name="l00400"></a><span class="lineno"> 400</span>  params_.<a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#a08ecd763b6785dfe872a6e517dc731e6">ref_C</a>.reset(args.ref_C.non_const_ref().data());</div><div class="line"><a name="l00401"></a><span class="lineno"> 401</span>  params_.<a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#a4f18093b18b0b6dd01a5df0a3813cd40">ref_D</a>.reset(args.ref_D.data()); </div><div class="line"><a name="l00402"></a><span class="lineno"> 402</span> </div><div class="line"><a name="l00403"></a><span class="lineno"> 403</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>;</div><div class="line"><a name="l00404"></a><span class="lineno"> 404</span>  }</div><div class="line"><a name="l00405"></a><span class="lineno"> 405</span> </div><div class="line"><a name="l00407"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ad7e9e393be872e401a5a777ceda529d9"> 407</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ad7e9e393be872e401a5a777ceda529d9">run</a>(cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00408"></a><span class="lineno"> 408</span> </div><div class="line"><a name="l00409"></a><span class="lineno"> 409</span>  ThreadblockSwizzle threadblock_swizzle;</div><div class="line"><a name="l00410"></a><span class="lineno"> 410</span> </div><div class="line"><a name="l00411"></a><span class="lineno"> 411</span>  dim3 grid = threadblock_swizzle.get_grid_shape(params_.<a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#a186903697a8ad274b6ac5416f7976c97">grid_tiled_shape</a>);</div><div class="line"><a name="l00412"></a><span class="lineno"> 412</span>  dim3 block(<a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched.html#aafddaefa35d27c76a89be8e692005615">GemmKernel::kThreadCount</a>, 1, 1);</div><div class="line"><a name="l00413"></a><span class="lineno"> 413</span> </div><div class="line"><a name="l00414"></a><span class="lineno"> 414</span>  cudaError_t result;</div><div class="line"><a name="l00415"></a><span class="lineno"> 415</span> </div><div class="line"><a name="l00416"></a><span class="lineno"> 416</span>  <span class="keywordtype">int</span> smem_size = int(<span class="keyword">sizeof</span>(<span class="keyword">typename</span> <a class="code" href="unioncutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1SharedStorage.html">GemmKernel::SharedStorage</a>));</div><div class="line"><a name="l00417"></a><span class="lineno"> 417</span>  <span class="keywordflow">if</span> (smem_size >= (48 << 10)) {</div><div class="line"><a name="l00418"></a><span class="lineno"> 418</span>  result = cudaFuncSetAttribute(Kernel<GemmKernel>,</div><div class="line"><a name="l00419"></a><span class="lineno"> 419</span>  cudaFuncAttributeMaxDynamicSharedMemorySize,</div><div class="line"><a name="l00420"></a><span class="lineno"> 420</span>  smem_size);</div><div class="line"><a name="l00421"></a><span class="lineno"> 421</span> </div><div class="line"><a name="l00422"></a><span class="lineno"> 422</span>  <span class="keywordflow">if</span> (result != cudaSuccess) {</div><div class="line"><a name="l00423"></a><span class="lineno"> 423</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da4d3b5847a0b17037c6b69faf8b1d4d71">Status::kErrorInternal</a>;</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> </div><div class="line"><a name="l00426"></a><span class="lineno"> 426</span>  result = cudaFuncSetAttribute(</div><div class="line"><a name="l00427"></a><span class="lineno"> 427</span>  Kernel<GemmKernel>,</div><div class="line"><a name="l00428"></a><span class="lineno"> 428</span>  cudaFuncAttributePreferredSharedMemoryCarveout, 100);</div><div class="line"><a name="l00429"></a><span class="lineno"> 429</span> </div><div class="line"><a name="l00430"></a><span class="lineno"> 430</span>  <span class="keywordflow">if</span> (result != cudaSuccess) {</div><div class="line"><a name="l00431"></a><span class="lineno"> 431</span>  <span class="keywordflow">return</span> <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da4d3b5847a0b17037c6b69faf8b1d4d71">Status::kErrorInternal</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="l00434"></a><span class="lineno"> 434</span> </div><div class="line"><a name="l00435"></a><span class="lineno"> 435</span>  cutlass::Kernel<GemmKernel><<<grid, block, smem_size, stream>>>(params_);</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>  result = cudaGetLastError();</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>  <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="l00440"></a><span class="lineno"> 440</span>  }</div><div class="line"><a name="l00441"></a><span class="lineno"> 441</span> </div><div class="line"><a name="l00443"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#afcdd646be7e79a60bac8dede563c56fa"> 443</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#afcdd646be7e79a60bac8dede563c56fa">operator()</a>(cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00444"></a><span class="lineno"> 444</span>  <span class="keywordflow">return</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ad7e9e393be872e401a5a777ceda529d9">run</a>(stream);</div><div class="line"><a name="l00445"></a><span class="lineno"> 445</span>  }</div><div class="line"><a name="l00446"></a><span class="lineno"> 446</span> </div><div class="line"><a name="l00448"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a45530b940ca86ce39cfc943da5713d80"> 448</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a45530b940ca86ce39cfc943da5713d80">operator()</a>(</div><div class="line"><a name="l00449"></a><span class="lineno"> 449</span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">Arguments</a> <span class="keyword">const</span> &args, </div><div class="line"><a name="l00450"></a><span class="lineno"> 450</span>  <span class="keywordtype">void</span> *workspace = <span class="keyword">nullptr</span>, </div><div class="line"><a name="l00451"></a><span class="lineno"> 451</span>  cudaStream_t stream = <span class="keyword">nullptr</span>) {</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>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> status = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2670ac441f48f6a0a2071c67c743ab8">initialize</a>(args, workspace);</div><div class="line"><a name="l00454"></a><span class="lineno"> 454</span>  </div><div class="line"><a name="l00455"></a><span class="lineno"> 455</span>  <span class="keywordflow">if</span> (status == <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>) {</div><div class="line"><a name="l00456"></a><span class="lineno"> 456</span>  status = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ad7e9e393be872e401a5a777ceda529d9">run</a>(stream);</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> </div><div class="line"><a name="l00459"></a><span class="lineno"> 459</span>  <span class="keywordflow">return</span> status;</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="l00464"></a><span class="lineno"> 464</span> </div><div class="line"><a name="l00466"></a><span class="lineno"> 466</span> <span class="keyword">template</span> <</div><div class="line"><a name="l00468"></a><span class="lineno"> 468</span>  <span class="keyword">typename</span> ElementA_,</div><div class="line"><a name="l00470"></a><span class="lineno"> 470</span>  <span class="keyword">typename</span> LayoutA_,</div><div class="line"><a name="l00472"></a><span class="lineno"> 472</span>  <span class="keyword">typename</span> ElementB_,</div><div class="line"><a name="l00474"></a><span class="lineno"> 474</span>  <span class="keyword">typename</span> LayoutB_,</div><div class="line"><a name="l00476"></a><span class="lineno"> 476</span>  <span class="keyword">typename</span> ElementC_,</div><div class="line"><a name="l00478"></a><span class="lineno"> 478</span>  <span class="keyword">typename</span> ElementAccumulator_,</div><div class="line"><a name="l00480"></a><span class="lineno"> 480</span>  <span class="keyword">typename</span> OperatorClass_,</div><div class="line"><a name="l00482"></a><span class="lineno"> 482</span>  <span class="keyword">typename</span> ArchTag_,</div><div class="line"><a name="l00484"></a><span class="lineno"> 484</span>  <span class="keyword">typename</span> ThreadblockShape_,</div><div class="line"><a name="l00486"></a><span class="lineno"> 486</span>  <span class="keyword">typename</span> WarpShape_,</div><div class="line"><a name="l00488"></a><span class="lineno"> 488</span>  <span class="keyword">typename</span> InstructionShape_,</div><div class="line"><a name="l00490"></a><span class="lineno"> 490</span>  <span class="keyword">typename</span> EpilogueOutputOp_,</div><div class="line"><a name="l00492"></a><span class="lineno"> 492</span>  <span class="keyword">typename</span> ThreadblockSwizzle_,</div><div class="line"><a name="l00494"></a><span class="lineno"> 494</span>  <span class="keywordtype">int</span> Stages,</div><div class="line"><a name="l00496"></a><span class="lineno"> 496</span>  <span class="keywordtype">int</span> AlignmentA,</div><div class="line"><a name="l00498"></a><span class="lineno"> 498</span>  <span class="keywordtype">int</span> AlignmentB,</div><div class="line"><a name="l00499"></a><span class="lineno"> 499</span>  <span class="keyword">typename</span> Operator_</div><div class="line"><a name="l00500"></a><span class="lineno"> 500</span> ></div><div class="line"><a name="l00501"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html"> 501</a></span> <span class="keyword">class </span><a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html">GemmBatched</a><</div><div class="line"><a name="l00502"></a><span class="lineno"> 502</span>  ElementA_,</div><div class="line"><a name="l00503"></a><span class="lineno"> 503</span>  LayoutA_,</div><div class="line"><a name="l00504"></a><span class="lineno"> 504</span>  ElementB_,</div><div class="line"><a name="l00505"></a><span class="lineno"> 505</span>  LayoutB_,</div><div class="line"><a name="l00506"></a><span class="lineno"> 506</span>  ElementC_,</div><div class="line"><a name="l00507"></a><span class="lineno"> 507</span>  layout::ColumnMajor,</div><div class="line"><a name="l00508"></a><span class="lineno"> 508</span>  ElementAccumulator_,</div><div class="line"><a name="l00509"></a><span class="lineno"> 509</span>  OperatorClass_,</div><div class="line"><a name="l00510"></a><span class="lineno"> 510</span>  ArchTag_,</div><div class="line"><a name="l00511"></a><span class="lineno"> 511</span>  ThreadblockShape_,</div><div class="line"><a name="l00512"></a><span class="lineno"> 512</span>  WarpShape_,</div><div class="line"><a name="l00513"></a><span class="lineno"> 513</span>  InstructionShape_,</div><div class="line"><a name="l00514"></a><span class="lineno"> 514</span>  EpilogueOutputOp_,</div><div class="line"><a name="l00515"></a><span class="lineno"> 515</span>  ThreadblockSwizzle_,</div><div class="line"><a name="l00516"></a><span class="lineno"> 516</span>  Stages,</div><div class="line"><a name="l00517"></a><span class="lineno"> 517</span>  AlignmentA,</div><div class="line"><a name="l00518"></a><span class="lineno"> 518</span>  AlignmentB,</div><div class="line"><a name="l00519"></a><span class="lineno"> 519</span>  Operator_</div><div class="line"><a name="l00520"></a><span class="lineno"> 520</span> > {</div><div class="line"><a name="l00521"></a><span class="lineno"> 521</span> <span class="keyword">public</span>:</div><div class="line"><a name="l00522"></a><span class="lineno"> 522</span> </div><div class="line"><a name="l00523"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a52b9261576b5633e901719f7c21d3369"> 523</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a52b9261576b5633e901719f7c21d3369">ElementA</a> = ElementA_;</div><div class="line"><a name="l00524"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#af623ca54d9554cdfafc09af7a22cdd62"> 524</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#af623ca54d9554cdfafc09af7a22cdd62">LayoutA</a> = LayoutA_;</div><div class="line"><a name="l00525"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a18266ad32200d3a72aba6e17a6297a3a"> 525</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="l00526"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3fd5c64783f88a7533801fef7d1375ad"> 526</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3fd5c64783f88a7533801fef7d1375ad">ElementB</a> = ElementB_;</div><div class="line"><a name="l00527"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a4aaaa6ca0e4b9f983fe37b4105fd058f"> 527</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a4aaaa6ca0e4b9f983fe37b4105fd058f">LayoutB</a> = LayoutB_;</div><div class="line"><a name="l00528"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a5595a5e74a0fb536794edf94cd5c7b7f"> 528</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="l00529"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#aef19ab5158e41856723852b3e307cc5d"> 529</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#aef19ab5158e41856723852b3e307cc5d">ElementC</a> = ElementC_;</div><div class="line"><a name="l00530"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#aed31a68c08cbfe9bf32d788be3f41679"> 530</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1layout_1_1ColumnMajor.html">LayoutC</a> = <a class="code" href="classcutlass_1_1layout_1_1ColumnMajor.html">layout::ColumnMajor</a>;</div><div class="line"><a name="l00531"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a04e1ec5b0634d45b9ae6811c0ea9f528"> 531</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="l00532"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#acd52c5c939493b3446af9682a2f7793c"> 532</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="l00533"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ae7f006ea8bc324d31de9dfbebc1b9327"> 533</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ae7f006ea8bc324d31de9dfbebc1b9327">ElementAccumulator</a> = ElementAccumulator_;</div><div class="line"><a name="l00534"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a37600c0bf3570bc4b21c26b2b64fc54a"> 534</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a37600c0bf3570bc4b21c26b2b64fc54a">OperatorClass</a> = OperatorClass_;</div><div class="line"><a name="l00535"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a681b145a9701109f9d72059bb874895b"> 535</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a681b145a9701109f9d72059bb874895b">ArchTag</a> = ArchTag_;</div><div class="line"><a name="l00536"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a657e50fb03ea4d16f7b904920d9aa000"> 536</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a657e50fb03ea4d16f7b904920d9aa000">ThreadblockShape</a> = ThreadblockShape_;</div><div class="line"><a name="l00537"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3760f803bd2b31b3fdf47741caa950fa"> 537</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3760f803bd2b31b3fdf47741caa950fa">WarpShape</a> = WarpShape_;</div><div class="line"><a name="l00538"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ae073edad6dd4447d7f99c94f4cd0c1c8"> 538</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ae073edad6dd4447d7f99c94f4cd0c1c8">InstructionShape</a> = InstructionShape_;</div><div class="line"><a name="l00539"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a494be150d3b809a4ecf66df682481905"> 539</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a494be150d3b809a4ecf66df682481905">EpilogueOutputOp</a> = EpilogueOutputOp_;</div><div class="line"><a name="l00540"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#af8b282788223086b80fbb097b22459ec"> 540</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#af8b282788223086b80fbb097b22459ec">ThreadblockSwizzle</a> = ThreadblockSwizzle_;</div><div class="line"><a name="l00541"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ab7f6a87909a3c2d45de71367a0d6eae3"> 541</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_1GemmBatched.html#a73837bda9ba209e546f6d996ede1afad">kStages</a> = Stages;</div><div class="line"><a name="l00542"></a><span class="lineno"> 542</span> </div><div class="line"><a name="l00543"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a4b924723475dcef72e0130ce1bb43956"> 543</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_1GemmBatched.html#a8976ed5c5e404ee87deaea4455d0d960">kAlignmentA</a> = AlignmentA;</div><div class="line"><a name="l00544"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a8f5d41976058b08562aa1819687d79a2"> 544</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_1GemmBatched.html#a17d34b2884711522fafcfd7c7500955c">kAlignmentB</a> = AlignmentB;</div><div class="line"><a name="l00545"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a79d27ed8dc23cc975f287ec0f041ddf9"> 545</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_1GemmBatched.html#ac1f190321a811fa91eec0096829b07ff">kAlignmentC</a> = EpilogueOutputOp::kCount;</div><div class="line"><a name="l00546"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a5a77d26d895197ff5224dac759e05766"> 546</a></span>  <span class="keyword">static</span> <span class="keywordtype">bool</span> <span class="keyword">const</span> kSplitKSerial = <span class="keyword">false</span>;</div><div class="line"><a name="l00547"></a><span class="lineno"> 547</span> </div><div class="line"><a name="l00548"></a><span class="lineno"> 548</span>  <span class="comment">//</span></div><div class="line"><a name="l00549"></a><span class="lineno"> 549</span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html">UnderlyingOperator</a> = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html">GemmBatched</a>< </div><div class="line"><a name="l00550"></a><span class="lineno"> 550</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac8bb1360bbc57bc63296cd48005c3c42">ElementB</a>,</div><div class="line"><a name="l00551"></a><span class="lineno"> 551</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="l00552"></a><span class="lineno"> 552</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a3fe2fcad97f15d63fa1a0214ef4861f2">ElementA</a>,</div><div class="line"><a name="l00553"></a><span class="lineno"> 553</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="l00554"></a><span class="lineno"> 554</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a008028eec1dd4b8c08128c6dfe44cce5">ElementC</a>,</div><div class="line"><a name="l00555"></a><span class="lineno"> 555</span>  <a class="code" href="classcutlass_1_1layout_1_1RowMajor.html">layout::RowMajor</a>, </div><div class="line"><a name="l00556"></a><span class="lineno"> 556</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a6e4cff55a6834d43cfbc97df40609eea">ElementAccumulator</a>,</div><div class="line"><a name="l00557"></a><span class="lineno"> 557</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a24935d746d97b0c994c9a9ade820d2d0">OperatorClass</a>,</div><div class="line"><a name="l00558"></a><span class="lineno"> 558</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#add5d679e0acf0813a52c209d2448e81b">ArchTag</a>,</div><div class="line"><a name="l00559"></a><span class="lineno"> 559</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a74aece33b6fafe58db1b41a6b7b87729">ThreadblockShape</a>,</div><div class="line"><a name="l00560"></a><span class="lineno"> 560</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a7b8d2dbfa562869deb58c88583951b58">WarpShape</a>,</div><div class="line"><a name="l00561"></a><span class="lineno"> 561</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a0dbb6d5185f223bb8242fc47a3b77757">InstructionShape</a>,</div><div class="line"><a name="l00562"></a><span class="lineno"> 562</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2d43ad49fa686ded524cc5f26b36c69">EpilogueOutputOp</a>,</div><div class="line"><a name="l00563"></a><span class="lineno"> 563</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a1b685fd66f6dc2c572be067ef1396a89">ThreadblockSwizzle</a>,</div><div class="line"><a name="l00564"></a><span class="lineno"> 564</span>  Stages,</div><div class="line"><a name="l00565"></a><span class="lineno"> 565</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a17d34b2884711522fafcfd7c7500955c">kAlignmentB</a>,</div><div class="line"><a name="l00566"></a><span class="lineno"> 566</span>  kAlignmentA</div><div class="line"><a name="l00567"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a55141da9e85b0c3556e531a2a6c19126"> 567</a></span>  >;</div><div class="line"><a name="l00568"></a><span class="lineno"> 568</span> </div><div class="line"><a name="l00569"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ab5f57fac13e42a08d351ac48c2cc9992"> 569</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ab5f57fac13e42a08d351ac48c2cc9992">UnderlyingArguments</a> = <span class="keyword">typename</span> UnderlyingOperator::Arguments;</div><div class="line"><a name="l00570"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3947c9b192bec2fad631334f31632353"> 570</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3947c9b192bec2fad631334f31632353">GemmKernel</a> = <span class="keyword">typename</span> <a class="code" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched.html">UnderlyingOperator::GemmKernel</a>;</div><div class="line"><a name="l00571"></a><span class="lineno"> 571</span> </div><div class="line"><a name="l00573"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html"> 573</a></span>  <span class="keyword">struct </span><a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">Arguments</a> {</div><div class="line"><a name="l00574"></a><span class="lineno"> 574</span> </div><div class="line"><a name="l00575"></a><span class="lineno"> 575</span>  <span class="comment">//</span></div><div class="line"><a name="l00576"></a><span class="lineno"> 576</span>  <span class="comment">// Data members</span></div><div class="line"><a name="l00577"></a><span class="lineno"> 577</span>  <span class="comment">//</span></div><div class="line"><a name="l00578"></a><span class="lineno"> 578</span> </div><div class="line"><a name="l00579"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ad0469cc3e961d21e212d026bccf6fe1a"> 579</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html">GemmCoord</a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ad0469cc3e961d21e212d026bccf6fe1a">problem_size</a>;</div><div class="line"><a name="l00580"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a1727630fc0525724df28a75ccf2580b9"> 580</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_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a1727630fc0525724df28a75ccf2580b9">ref_A</a>;</div><div class="line"><a name="l00581"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ac8830c9ed0e0a8bd7aa2aa4382550a2f"> 581</a></span>  int64_t <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ac8830c9ed0e0a8bd7aa2aa4382550a2f">stride_A</a>;</div><div class="line"><a name="l00582"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ad7d2b82b83d7503b9f920ce3bdcdffa5"> 582</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_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ad7d2b82b83d7503b9f920ce3bdcdffa5">ref_B</a>;</div><div class="line"><a name="l00583"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a302101a4e5c00c843b3c525ddb94c117"> 583</a></span>  int64_t <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a302101a4e5c00c843b3c525ddb94c117">stride_B</a>;</div><div class="line"><a name="l00584"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#aa9e30e41627595590421d8b53941b2b2"> 584</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_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#aa9e30e41627595590421d8b53941b2b2">ref_C</a>;</div><div class="line"><a name="l00585"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a9f8a044d7b7439192dfe2bf488558ed3"> 585</a></span>  int64_t <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a9f8a044d7b7439192dfe2bf488558ed3">stride_C</a>;</div><div class="line"><a name="l00586"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a17c4e381e91229a8ef15b18ee5ec073d"> 586</a></span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC, LayoutC></a> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a17c4e381e91229a8ef15b18ee5ec073d">ref_D</a>;</div><div class="line"><a name="l00587"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ac181dba327e605b6cde9de5c7f176e7c"> 587</a></span>  int64_t <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ac181dba327e605b6cde9de5c7f176e7c">stride_D</a>;</div><div class="line"><a name="l00588"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#af9c2fa1e0cc0456197c2cc0840c89982"> 588</a></span>  <span class="keyword">typename</span> EpilogueOutputOp::Params <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#af9c2fa1e0cc0456197c2cc0840c89982">epilogue</a>;</div><div class="line"><a name="l00589"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#adb66f3083f56c15578b139b7935452b5"> 589</a></span>  <span class="keywordtype">int</span> <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#adb66f3083f56c15578b139b7935452b5">batch_count</a>;</div><div class="line"><a name="l00590"></a><span class="lineno"> 590</span> </div><div class="line"><a name="l00591"></a><span class="lineno"> 591</span>  <span class="comment">//</span></div><div class="line"><a name="l00592"></a><span class="lineno"> 592</span>  <span class="comment">// Methods</span></div><div class="line"><a name="l00593"></a><span class="lineno"> 593</span>  <span class="comment">//</span></div><div class="line"><a name="l00594"></a><span class="lineno"> 594</span> </div><div class="line"><a name="l00596"></a><span class="lineno"> 596</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00597"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ae86daa985279c77e57e682b64a68d330"> 597</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ae86daa985279c77e57e682b64a68d330">Arguments</a>() { }</div><div class="line"><a name="l00598"></a><span class="lineno"> 598</span> </div><div class="line"><a name="l00600"></a><span class="lineno"> 600</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00601"></a><span class="lineno"><a class="line" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a2129a4dccbd73f8c0f26b08ce5a5cb28"> 601</a></span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a2129a4dccbd73f8c0f26b08ce5a5cb28">Arguments</a>(</div><div class="line"><a name="l00602"></a><span class="lineno"> 602</span>  <a class="code" href="structcutlass_1_1gemm_1_1GemmCoord.html">GemmCoord</a> problem_size_,</div><div class="line"><a name="l00603"></a><span class="lineno"> 603</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementA const, LayoutA></a> ref_A_,</div><div class="line"><a name="l00604"></a><span class="lineno"> 604</span>  int64_t stride_A_,</div><div class="line"><a name="l00605"></a><span class="lineno"> 605</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementB const, LayoutB></a> ref_B_,</div><div class="line"><a name="l00606"></a><span class="lineno"> 606</span>  int64_t stride_B_,</div><div class="line"><a name="l00607"></a><span class="lineno"> 607</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC const, LayoutC></a> ref_C_,</div><div class="line"><a name="l00608"></a><span class="lineno"> 608</span>  int64_t stride_C_,</div><div class="line"><a name="l00609"></a><span class="lineno"> 609</span>  <a class="code" href="classcutlass_1_1TensorRef.html">TensorRef<ElementC, LayoutC></a> ref_D_,</div><div class="line"><a name="l00610"></a><span class="lineno"> 610</span>  int64_t stride_D_,</div><div class="line"><a name="l00611"></a><span class="lineno"> 611</span>  <span class="keyword">typename</span> EpilogueOutputOp::Params epilogue_,</div><div class="line"><a name="l00612"></a><span class="lineno"> 612</span>  <span class="keywordtype">int</span> batch_count_</div><div class="line"><a name="l00613"></a><span class="lineno"> 613</span>  ):</div><div class="line"><a name="l00614"></a><span class="lineno"> 614</span>  problem_size(problem_size_),</div><div class="line"><a name="l00615"></a><span class="lineno"> 615</span>  ref_A(ref_A_),</div><div class="line"><a name="l00616"></a><span class="lineno"> 616</span>  stride_A(stride_A_),</div><div class="line"><a name="l00617"></a><span class="lineno"> 617</span>  ref_B(ref_B_),</div><div class="line"><a name="l00618"></a><span class="lineno"> 618</span>  stride_B(stride_B_),</div><div class="line"><a name="l00619"></a><span class="lineno"> 619</span>  ref_C(ref_C_),</div><div class="line"><a name="l00620"></a><span class="lineno"> 620</span>  stride_C(stride_C_),</div><div class="line"><a name="l00621"></a><span class="lineno"> 621</span>  ref_D(ref_D_),</div><div class="line"><a name="l00622"></a><span class="lineno"> 622</span>  stride_D(stride_D_),</div><div class="line"><a name="l00623"></a><span class="lineno"> 623</span>  epilogue(epilogue_),</div><div class="line"><a name="l00624"></a><span class="lineno"> 624</span>  batch_count(batch_count_) { }</div><div class="line"><a name="l00625"></a><span class="lineno"> 625</span>  };</div><div class="line"><a name="l00626"></a><span class="lineno"> 626</span> </div><div class="line"><a name="l00627"></a><span class="lineno"> 627</span> <span class="keyword">private</span>:</div><div class="line"><a name="l00628"></a><span class="lineno"> 628</span> </div><div class="line"><a name="l00629"></a><span class="lineno"> 629</span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html">UnderlyingOperator</a> underlying_operator_;</div><div class="line"><a name="l00630"></a><span class="lineno"> 630</span> </div><div class="line"><a name="l00631"></a><span class="lineno"> 631</span> <span class="keyword">public</span>:</div><div class="line"><a name="l00632"></a><span class="lineno"> 632</span> </div><div class="line"><a name="l00634"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a75922fd7bcd77fbc714cd87681f692bf"> 634</a></span>  <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a75922fd7bcd77fbc714cd87681f692bf">GemmBatched</a>() { }</div><div class="line"><a name="l00635"></a><span class="lineno"> 635</span> </div><div class="line"><a name="l00637"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ac4ef1ac1e0876aaee5bff50dc09fe8a9"> 637</a></span>  <span class="keyword">static</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ab5f57fac13e42a08d351ac48c2cc9992">UnderlyingArguments</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ac4ef1ac1e0876aaee5bff50dc09fe8a9">to_underlying_arguments</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">Arguments</a> <span class="keyword">const</span> &args) {</div><div class="line"><a name="l00638"></a><span class="lineno"> 638</span>  <span class="keywordflow">return</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ab5f57fac13e42a08d351ac48c2cc9992">UnderlyingArguments</a>(</div><div class="line"><a name="l00639"></a><span class="lineno"> 639</span>  {args.problem_size.n(), args.problem_size.m(), args.problem_size.k()},</div><div class="line"><a name="l00640"></a><span class="lineno"> 640</span>  {args.ref_B.data(), args.ref_B.stride(0)},</div><div class="line"><a name="l00641"></a><span class="lineno"> 641</span>  args.stride_B,</div><div class="line"><a name="l00642"></a><span class="lineno"> 642</span>  {args.ref_A.data(), args.ref_A.stride(0)},</div><div class="line"><a name="l00643"></a><span class="lineno"> 643</span>  args.stride_A,</div><div class="line"><a name="l00644"></a><span class="lineno"> 644</span>  {args.ref_C.data(), args.ref_C.stride(0)},</div><div class="line"><a name="l00645"></a><span class="lineno"> 645</span>  args.stride_C,</div><div class="line"><a name="l00646"></a><span class="lineno"> 646</span>  {args.ref_D.data(), args.ref_D.stride(0)},</div><div class="line"><a name="l00647"></a><span class="lineno"> 647</span>  args.stride_D,</div><div class="line"><a name="l00648"></a><span class="lineno"> 648</span>  args.epilogue,</div><div class="line"><a name="l00649"></a><span class="lineno"> 649</span>  args.batch_count</div><div class="line"><a name="l00650"></a><span class="lineno"> 650</span>  );</div><div class="line"><a name="l00651"></a><span class="lineno"> 651</span>  }</div><div class="line"><a name="l00652"></a><span class="lineno"> 652</span> </div><div class="line"><a name="l00654"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#abbd82c0f989a9d07e5e222db96386701"> 654</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_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#abbd82c0f989a9d07e5e222db96386701">can_implement</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">Arguments</a> <span class="keyword">const</span> &args) {</div><div class="line"><a name="l00655"></a><span class="lineno"> 655</span> </div><div class="line"><a name="l00656"></a><span class="lineno"> 656</span>  <span class="keywordflow">return</span> UnderlyingOperator::can_implement(to_underlying_arguments(args));</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_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3687659e826ba7f38bb060ad6020a739"> 660</a></span>  <span class="keyword">static</span> <span class="keywordtype">size_t</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3687659e826ba7f38bb060ad6020a739">get_workspace_size</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">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::get_workspace_size(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_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a428d8b1c4ac36040145a59d8e4cff3d2"> 666</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a428d8b1c4ac36040145a59d8e4cff3d2">initialize</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">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="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> underlying_operator_.<a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2670ac441f48f6a0a2071c67c743ab8">initialize</a>(to_underlying_arguments(args), workspace);</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_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a9f0c7054068175c1891e4820857603c3"> 672</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a9f0c7054068175c1891e4820857603c3">update</a>(<a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">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="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_1GemmBatched.html#ab7c3e9a33a1c62513ec6eee3e2598df6">update</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_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#abcae3d15f1ec2ee7ae93690c82fbee8a"> 678</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#abcae3d15f1ec2ee7ae93690c82fbee8a">run</a>(cudaStream_t stream = <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_1GemmBatched.html#ad7e9e393be872e401a5a777ceda529d9">run</a>(stream);</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_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a00805989734182945f982cab23a5dca8"> 684</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a00805989734182945f982cab23a5dca8">operator()</a>(cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00685"></a><span class="lineno"> 685</span>  <span class="keywordflow">return</span> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ad7e9e393be872e401a5a777ceda529d9">run</a>(stream);</div><div class="line"><a name="l00686"></a><span class="lineno"> 686</span>  }</div><div class="line"><a name="l00687"></a><span class="lineno"> 687</span> </div><div class="line"><a name="l00689"></a><span class="lineno"><a class="line" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a53ca4db66d0d2c96d9036d8eb7c6072b"> 689</a></span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a53ca4db66d0d2c96d9036d8eb7c6072b">operator()</a>(</div><div class="line"><a name="l00690"></a><span class="lineno"> 690</span>  <a class="code" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">Arguments</a> <span class="keyword">const</span> &args, </div><div class="line"><a name="l00691"></a><span class="lineno"> 691</span>  <span class="keywordtype">void</span> *workspace = <span class="keyword">nullptr</span>, </div><div class="line"><a name="l00692"></a><span class="lineno"> 692</span>  cudaStream_t stream = <span class="keyword">nullptr</span>) {</div><div class="line"><a name="l00693"></a><span class="lineno"> 693</span>  </div><div class="line"><a name="l00694"></a><span class="lineno"> 694</span>  <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18d">Status</a> status = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2670ac441f48f6a0a2071c67c743ab8">initialize</a>(args, workspace);</div><div class="line"><a name="l00695"></a><span class="lineno"> 695</span>  </div><div class="line"><a name="l00696"></a><span class="lineno"> 696</span>  <span class="keywordflow">if</span> (status == <a class="code" href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18da8c632159fa131f09d04f94e3cbcd8782">Status::kSuccess</a>) {</div><div class="line"><a name="l00697"></a><span class="lineno"> 697</span>  status = <a class="code" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ad7e9e393be872e401a5a777ceda529d9">run</a>(stream);</div><div class="line"><a name="l00698"></a><span class="lineno"> 698</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>  <span class="keywordflow">return</span> status;</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> </div><div class="line"><a name="l00703"></a><span class="lineno"> 703</span> };</div><div class="line"><a name="l00704"></a><span class="lineno"> 704</span> </div><div class="line"><a name="l00706"></a><span class="lineno"> 706</span> </div><div class="line"><a name="l00707"></a><span class="lineno"> 707</span> } <span class="comment">// namespace device</span></div><div class="line"><a name="l00708"></a><span class="lineno"> 708</span> } <span class="comment">// namespace gemm</span></div><div class="line"><a name="l00709"></a><span class="lineno"> 709</span> } <span class="comment">// namespace cutlass</span></div><div class="line"><a name="l00710"></a><span class="lineno"> 710</span> </div><div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_aef19ab5158e41856723852b3e307cc5d"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#aef19ab5158e41856723852b3e307cc5d">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::ElementC</a></div><div class="ttdeci">ElementC_ ElementC</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:529</div></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_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_af8b282788223086b80fbb097b22459ec"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#af8b282788223086b80fbb097b22459ec">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::ThreadblockSwizzle</a></div><div class="ttdeci">ThreadblockSwizzle_ ThreadblockSwizzle</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:540</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a17d34b2884711522fafcfd7c7500955c"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a17d34b2884711522fafcfd7c7500955c">cutlass::gemm::device::GemmBatched::kAlignmentB</a></div><div class="ttdeci">static int const kAlignmentB</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:236</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_a1727630fc0525724df28a75ccf2580b9"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a1727630fc0525724df28a75ccf2580b9">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::ref_A</a></div><div class="ttdeci">TensorRef< ElementA const, LayoutA > ref_A</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:580</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_ae4450f06a6975191d94026865e445578"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#ae4450f06a6975191d94026865e445578">cutlass::gemm::device::GemmBatched::Arguments::ref_D</a></div><div class="ttdeci">TensorRef< ElementC, LayoutC > ref_D</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:280</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_aab4fc258e38ebcf9b430a5dee6daba5e"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#aab4fc258e38ebcf9b430a5dee6daba5e">cutlass::gemm::device::GemmBatched::Arguments::problem_size</a></div><div class="ttdeci">GemmCoord problem_size</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:273</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="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a3fd5c64783f88a7533801fef7d1375ad"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3fd5c64783f88a7533801fef7d1375ad">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::ElementB</a></div><div class="ttdeci">ElementB_ ElementB</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:526</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a53ca4db66d0d2c96d9036d8eb7c6072b"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a53ca4db66d0d2c96d9036d8eb7c6072b">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::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> device/gemm_batched.h:689</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_a05a1d9720fbb16a20b94049900b0d04f"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a05a1d9720fbb16a20b94049900b0d04f">cutlass::gemm::device::GemmBatched::Arguments::stride_D</a></div><div class="ttdeci">int64_t stride_D</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:281</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params_html_a4f18093b18b0b6dd01a5df0a3813cd40"><div class="ttname"><a href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#a4f18093b18b0b6dd01a5df0a3813cd40">cutlass::gemm::kernel::GemmBatched::Params::ref_D</a></div><div class="ttdeci">Epilogue::OutputTileIterator::TensorRef ref_D</div><div class="ttdef"><b>Definition:</b> kernel/gemm_batched.h:74</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_a0befb9945aadcba460f4d1ad73020e9c"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a0befb9945aadcba460f4d1ad73020e9c">cutlass::gemm::device::GemmBatched::Arguments::Arguments</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Arguments(GemmCoord problem_size_, TensorRef< ElementA const, LayoutA > ref_A_, int64_t stride_A_, TensorRef< ElementB const, LayoutB > ref_B_, int64_t stride_B_, TensorRef< ElementC const, LayoutC > ref_C_, int64_t stride_C_, TensorRef< ElementC, LayoutC > ref_D_, int64_t stride_D_, typename EpilogueOutputOp::Params epilogue_, int batch_count_)</div><div class="ttdoc">Constructs an Arguments structure. </div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:295</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a4aaaa6ca0e4b9f983fe37b4105fd058f"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a4aaaa6ca0e4b9f983fe37b4105fd058f">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::LayoutB</a></div><div class="ttdeci">LayoutB_ LayoutB</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:527</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a6acd50cfc477e95dbcf0d4fbba5df65c"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a6acd50cfc477e95dbcf0d4fbba5df65c">cutlass::gemm::device::GemmBatched< 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 >::DefaultGemmKernel</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, false, Operator, false >::GemmKernel DefaultGemmKernel</div><div class="ttdoc">Define the kernel. </div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:262</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_abbd82c0f989a9d07e5e222db96386701"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#abbd82c0f989a9d07e5e222db96386701">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::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> device/gemm_batched.h:654</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a657e50fb03ea4d16f7b904920d9aa000"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a657e50fb03ea4d16f7b904920d9aa000">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::ThreadblockShape</a></div><div class="ttdeci">ThreadblockShape_ ThreadblockShape</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:536</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_afcdd646be7e79a60bac8dede563c56fa"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#afcdd646be7e79a60bac8dede563c56fa">cutlass::gemm::device::GemmBatched::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> device/gemm_batched.h:443</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="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html">cutlass::gemm::device::GemmBatched::Arguments</a></div><div class="ttdoc">Argument structure. </div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:267</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a52b7263c5c86e900bcca681d07f19101"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a52b7263c5c86e900bcca681d07f19101">cutlass::gemm::device::GemmBatched< 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 >::Operator</a></div><div class="ttdeci">typename DefaultGemmConfiguration< OperatorClass, ArchTag, ElementB, ElementA, ElementC,ElementAccumulator >::Operator Operator</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:238</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params_html_ade55adc311c5561efe76f53ffd56d1f4"><div class="ttname"><a href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#ade55adc311c5561efe76f53ffd56d1f4">cutlass::gemm::kernel::GemmBatched::Params::ref_B</a></div><div class="ttdeci">Mma::IteratorB::TensorRef ref_B</div><div class="ttdef"><b>Definition:</b> kernel/gemm_batched.h:68</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_a9f8a044d7b7439192dfe2bf488558ed3"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a9f8a044d7b7439192dfe2bf488558ed3">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::stride_C</a></div><div class="ttdeci">int64_t stride_C</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:585</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a37600c0bf3570bc4b21c26b2b64fc54a"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a37600c0bf3570bc4b21c26b2b64fc54a">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::OperatorClass</a></div><div class="ttdeci">OperatorClass_ OperatorClass</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:534</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_acb489f1bb4fed9e4314a6b6a3cbd04a9"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#acb489f1bb4fed9e4314a6b6a3cbd04a9">cutlass::gemm::device::GemmBatched< 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 >::LayoutB</a></div><div class="ttdeci">typename layout::LayoutTranspose< LayoutA >::type LayoutB</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:220</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a681b145a9701109f9d72059bb874895b"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a681b145a9701109f9d72059bb874895b">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::ArchTag</a></div><div class="ttdeci">ArchTag_ ArchTag</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:535</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a3947c9b192bec2fad631334f31632353"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3947c9b192bec2fad631334f31632353">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::GemmKernel</a></div><div class="ttdeci">typename UnderlyingOperator::GemmKernel GemmKernel</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:570</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_ac2009bb52372115624aa5c4f75b720e5"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac2009bb52372115624aa5c4f75b720e5">cutlass::gemm::device::GemmBatched::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> device/gemm_batched.h:361</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_ac181dba327e605b6cde9de5c7f176e7c"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ac181dba327e605b6cde9de5c7f176e7c">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::stride_D</a></div><div class="ttdeci">int64_t stride_D</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:587</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a3760f803bd2b31b3fdf47741caa950fa"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3760f803bd2b31b3fdf47741caa950fa">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::WarpShape</a></div><div class="ttdeci">WarpShape_ WarpShape</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:537</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_af9c2fa1e0cc0456197c2cc0840c89982"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#af9c2fa1e0cc0456197c2cc0840c89982">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::epilogue</a></div><div class="ttdeci">EpilogueOutputOp::Params epilogue</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:588</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a9f0c7054068175c1891e4820857603c3"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a9f0c7054068175c1891e4820857603c3">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::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> device/gemm_batched.h:672</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_a11ef91161a92459d72b56144cd6b4495"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a11ef91161a92459d72b56144cd6b4495">cutlass::gemm::device::GemmBatched::Arguments::stride_A</a></div><div class="ttdeci">int64_t stride_A</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:275</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_ae073edad6dd4447d7f99c94f4cd0c1c8"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ae073edad6dd4447d7f99c94f4cd0c1c8">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::InstructionShape</a></div><div class="ttdeci">InstructionShape_ InstructionShape</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:538</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a428d8b1c4ac36040145a59d8e4cff3d2"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a428d8b1c4ac36040145a59d8e4cff3d2">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::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> device/gemm_batched.h:666</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_abcae3d15f1ec2ee7ae93690c82fbee8a"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#abcae3d15f1ec2ee7ae93690c82fbee8a">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::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> device/gemm_batched.h:678</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_ae7f006ea8bc324d31de9dfbebc1b9327"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ae7f006ea8bc324d31de9dfbebc1b9327">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::ElementAccumulator</a></div><div class="ttdeci">ElementAccumulator_ ElementAccumulator</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:533</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params_html_a08ecd763b6785dfe872a6e517dc731e6"><div class="ttname"><a href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#a08ecd763b6785dfe872a6e517dc731e6">cutlass::gemm::kernel::GemmBatched::Params::ref_C</a></div><div class="ttdeci">Epilogue::OutputTileIterator::TensorRef ref_C</div><div class="ttdef"><b>Definition:</b> kernel/gemm_batched.h:71</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_1GemmBatched_html_a24935d746d97b0c994c9a9ade820d2d0"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a24935d746d97b0c994c9a9ade820d2d0">cutlass::gemm::device::GemmBatched< 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 >::OperatorClass</a></div><div class="ttdeci">OperatorClass OperatorClass</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:227</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_a302101a4e5c00c843b3c525ddb94c117"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a302101a4e5c00c843b3c525ddb94c117">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::stride_B</a></div><div class="ttdeci">int64_t stride_B</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:583</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_1GemmBatched_html_ac1f190321a811fa91eec0096829b07ff"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac1f190321a811fa91eec0096829b07ff">cutlass::gemm::device::GemmBatched::kAlignmentC</a></div><div class="ttdeci">static int const kAlignmentC</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:237</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_a17c4e381e91229a8ef15b18ee5ec073d"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a17c4e381e91229a8ef15b18ee5ec073d">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::ref_D</a></div><div class="ttdeci">TensorRef< ElementC, LayoutC > ref_D</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:586</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a74aece33b6fafe58db1b41a6b7b87729"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a74aece33b6fafe58db1b41a6b7b87729">cutlass::gemm::device::GemmBatched< 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 >::ThreadblockShape</a></div><div class="ttdeci">ThreadblockShape ThreadblockShape</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:229</div></div>
|
|
<div class="ttc" id="unioncutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1SharedStorage_html"><div class="ttname"><a href="unioncutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1SharedStorage.html">cutlass::gemm::kernel::GemmBatched::SharedStorage</a></div><div class="ttdoc">Shared memory storage structure. </div><div class="ttdef"><b>Definition:</b> kernel/gemm_batched.h:124</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_a72bd469f15b44e492cf84658b5f09ad5"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a72bd469f15b44e492cf84658b5f09ad5">cutlass::gemm::device::GemmBatched::Arguments::Arguments</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Arguments()</div><div class="ttdoc">Default ctor. </div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:291</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params_html_a186903697a8ad274b6ac5416f7976c97"><div class="ttname"><a href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#a186903697a8ad274b6ac5416f7976c97">cutlass::gemm::kernel::GemmBatched::Params::grid_tiled_shape</a></div><div class="ttdeci">cutlass::gemm::GemmCoord grid_tiled_shape</div><div class="ttdef"><b>Definition:</b> kernel/gemm_batched.h:63</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a0dbb6d5185f223bb8242fc47a3b77757"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a0dbb6d5185f223bb8242fc47a3b77757">cutlass::gemm::device::GemmBatched< 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 >::InstructionShape</a></div><div class="ttdeci">InstructionShape InstructionShape</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:231</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a75922fd7bcd77fbc714cd87681f692bf"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a75922fd7bcd77fbc714cd87681f692bf">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::GemmBatched</a></div><div class="ttdeci">GemmBatched()</div><div class="ttdoc">Constructs the GEMM. </div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:634</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_ad0469cc3e961d21e212d026bccf6fe1a"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ad0469cc3e961d21e212d026bccf6fe1a">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::problem_size</a></div><div class="ttdeci">GemmCoord problem_size</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:579</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a73837bda9ba209e546f6d996ede1afad"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a73837bda9ba209e546f6d996ede1afad">cutlass::gemm::device::GemmBatched::kStages</a></div><div class="ttdeci">static int const kStages</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:234</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_ab7c3e9a33a1c62513ec6eee3e2598df6"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ab7c3e9a33a1c62513ec6eee3e2598df6">cutlass::gemm::device::GemmBatched::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> device/gemm_batched.h:396</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_html_aafddaefa35d27c76a89be8e692005615"><div class="ttname"><a href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched.html#aafddaefa35d27c76a89be8e692005615">cutlass::gemm::kernel::GemmBatched::kThreadCount</a></div><div class="ttdeci">static int const kThreadCount</div><div class="ttdef"><b>Definition:</b> kernel/gemm_batched.h:58</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_acb4d53fbea4366349574091d68594558"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#acb4d53fbea4366349574091d68594558">cutlass::gemm::device::GemmBatched::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> device/gemm_batched.h:332</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_a3ce1385631b05430fa5dfc1e9a3671b8"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a3ce1385631b05430fa5dfc1e9a3671b8">cutlass::gemm::device::GemmBatched::Arguments::stride_C</a></div><div class="ttdeci">int64_t stride_C</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:279</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params_html"><div class="ttname"><a href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html">cutlass::gemm::kernel::GemmBatched::Params</a></div><div class="ttdoc">Parameters structure. </div><div class="ttdef"><b>Definition:</b> kernel/gemm_batched.h:61</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="namespacecutlass_html_ac5a88c5840a28a9e0206b9cc7812a18daa4867e1466f5d067dbec566abfe5a67a"><div class="ttname"><a href="namespacecutlass.html#ac5a88c5840a28a9e0206b9cc7812a18daa4867e1466f5d067dbec566abfe5a67a">cutlass::Status::kErrorMisalignedOperand</a></div><div class="ttdoc">operands fail alignment requirements. </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="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a008028eec1dd4b8c08128c6dfe44cce5"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a008028eec1dd4b8c08128c6dfe44cce5">cutlass::gemm::device::GemmBatched< 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 >::ElementC</a></div><div class="ttdeci">ElementC ElementC</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:222</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_ab5f57fac13e42a08d351ac48c2cc9992"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ab5f57fac13e42a08d351ac48c2cc9992">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::UnderlyingArguments</a></div><div class="ttdeci">typename UnderlyingOperator::Arguments UnderlyingArguments</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:569</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_ac4ef1ac1e0876aaee5bff50dc09fe8a9"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#ac4ef1ac1e0876aaee5bff50dc09fe8a9">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::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> device/gemm_batched.h:637</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="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_a48844293c34b9c44fe57f577370664ea"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a48844293c34b9c44fe57f577370664ea">cutlass::gemm::device::GemmBatched::Arguments::ref_B</a></div><div class="ttdeci">TensorRef< ElementB const, LayoutB > ref_B</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:276</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a8976ed5c5e404ee87deaea4455d0d960"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a8976ed5c5e404ee87deaea4455d0d960">cutlass::gemm::device::GemmBatched::kAlignmentA</a></div><div class="ttdeci">static int const kAlignmentA</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:235</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="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_aa2670ac441f48f6a0a2071c67c743ab8"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2670ac441f48f6a0a2071c67c743ab8">cutlass::gemm::device::GemmBatched::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> device/gemm_batched.h:366</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a1b685fd66f6dc2c572be067ef1396a89"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a1b685fd66f6dc2c572be067ef1396a89">cutlass::gemm::device::GemmBatched< 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 >::ThreadblockSwizzle</a></div><div class="ttdeci">ThreadblockSwizzle ThreadblockSwizzle</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:233</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a494be150d3b809a4ecf66df682481905"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a494be150d3b809a4ecf66df682481905">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::EpilogueOutputOp</a></div><div class="ttdeci">EpilogueOutputOp_ EpilogueOutputOp</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:539</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="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a72a26fb286181aa5ca1fb66d9b385f7f"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a72a26fb286181aa5ca1fb66d9b385f7f">cutlass::gemm::device::GemmBatched::GemmBatched</a></div><div class="ttdeci">GemmBatched()</div><div class="ttdoc">Constructs the GEMM. </div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:329</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="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_aa2d43ad49fa686ded524cc5f26b36c69"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#aa2d43ad49fa686ded524cc5f26b36c69">cutlass::gemm::device::GemmBatched< 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 >::EpilogueOutputOp</a></div><div class="ttdeci">EpilogueOutputOp EpilogueOutputOp</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:232</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_ac99ca8f9d8a0053e647a6c99b018bda5"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#ac99ca8f9d8a0053e647a6c99b018bda5">cutlass::gemm::device::GemmBatched::Arguments::batch_count</a></div><div class="ttdeci">int batch_count</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:283</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a00805989734182945f982cab23a5dca8"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a00805989734182945f982cab23a5dca8">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::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> device/gemm_batched.h:684</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="kernel_2gemm__batched_8h_html"><div class="ttname"><a href="kernel_2gemm__batched_8h.html">gemm_batched.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="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a52b9261576b5633e901719f7c21d3369"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a52b9261576b5633e901719f7c21d3369">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::ElementA</a></div><div class="ttdeci">ElementA_ ElementA</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:523</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a9f9d11529b28ced91c4b05c2530b7a70"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a9f9d11529b28ced91c4b05c2530b7a70">cutlass::gemm::device::GemmBatched< 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 >::LayoutA</a></div><div class="ttdeci">typename layout::LayoutTranspose< LayoutB >::type LayoutA</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:217</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params_html_a05909ba49e633c7eeb0707166c72a4ee"><div class="ttname"><a href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#a05909ba49e633c7eeb0707166c72a4ee">cutlass::gemm::kernel::GemmBatched::Params::problem_size</a></div><div class="ttdeci">cutlass::gemm::GemmCoord problem_size</div><div class="ttdef"><b>Definition:</b> kernel/gemm_batched.h:62</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_1GemmBatched_html_ac8bb1360bbc57bc63296cd48005c3c42"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ac8bb1360bbc57bc63296cd48005c3c42">cutlass::gemm::device::GemmBatched< 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 >::ElementB</a></div><div class="ttdeci">ElementA ElementB</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:219</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_a3687659e826ba7f38bb060ad6020a739"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#a3687659e826ba7f38bb060ad6020a739">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::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> device/gemm_batched.h:660</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a3fe2fcad97f15d63fa1a0214ef4861f2"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a3fe2fcad97f15d63fa1a0214ef4861f2">cutlass::gemm::device::GemmBatched< 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 >::ElementA</a></div><div class="ttdeci">ElementB ElementA</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:216</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_ab0955b722ad4ea0217f725e34b3bcfbe"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#ab0955b722ad4ea0217f725e34b3bcfbe">cutlass::gemm::device::GemmBatched::Arguments::ref_C</a></div><div class="ttdeci">TensorRef< ElementC const, LayoutC > ref_C</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:278</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_a2129a4dccbd73f8c0f26b08ce5a5cb28"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#a2129a4dccbd73f8c0f26b08ce5a5cb28">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::Arguments</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Arguments(GemmCoord problem_size_, TensorRef< ElementA const, LayoutA > ref_A_, int64_t stride_A_, TensorRef< ElementB const, LayoutB > ref_B_, int64_t stride_B_, TensorRef< ElementC const, LayoutC > ref_C_, int64_t stride_C_, TensorRef< ElementC, LayoutC > ref_D_, int64_t stride_D_, typename EpilogueOutputOp::Params epilogue_, int batch_count_)</div><div class="ttdoc">Constructs an Arguments structure. </div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:601</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_ae86daa985279c77e57e682b64a68d330"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ae86daa985279c77e57e682b64a68d330">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::Arguments</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Arguments()</div><div class="ttdoc">Default ctor. </div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:597</div></div>
|
|
<div class="ttc" id="namespacecutlass_html_aa43b0a7d59635cb2d9ac96a077c988c3"><div class="ttname"><a href="namespacecutlass.html#aa43b0a7d59635cb2d9ac96a077c988c3">cutlass::TensorRef_aligned</a></div><div class="ttdeci">bool TensorRef_aligned(TensorRef< Element, Layout > const &ref, int alignment)</div><div class="ttdef"><b>Definition:</b> tensor_ref.h:382</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_ad7d2b82b83d7503b9f920ce3bdcdffa5"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ad7d2b82b83d7503b9f920ce3bdcdffa5">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::ref_B</a></div><div class="ttdeci">TensorRef< ElementB const, LayoutB > ref_B</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:582</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="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params_html_ad1867c0875c10e6327c7fae16acd35a3"><div class="ttname"><a href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html#ad1867c0875c10e6327c7fae16acd35a3">cutlass::gemm::kernel::GemmBatched::Params::ref_A</a></div><div class="ttdeci">Mma::IteratorA::TensorRef ref_A</div><div class="ttdef"><b>Definition:</b> kernel/gemm_batched.h:65</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_ad7e9e393be872e401a5a777ceda529d9"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#ad7e9e393be872e401a5a777ceda529d9">cutlass::gemm::device::GemmBatched::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> device/gemm_batched.h:407</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_1kernel_1_1GemmBatched_html"><div class="ttname"><a href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched.html">cutlass::gemm::kernel::GemmBatched</a></div><div class="ttdef"><b>Definition:</b> kernel/gemm_batched.h:49</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a6e4cff55a6834d43cfbc97df40609eea"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a6e4cff55a6834d43cfbc97df40609eea">cutlass::gemm::device::GemmBatched< 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 >::ElementAccumulator</a></div><div class="ttdeci">ElementAccumulator ElementAccumulator</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:226</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_aa867aa186538d34251d75ccc891453d7"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#aa867aa186538d34251d75ccc891453d7">cutlass::gemm::device::GemmBatched::Arguments::stride_B</a></div><div class="ttdeci">int64_t stride_B</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:277</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a7b8d2dbfa562869deb58c88583951b58"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a7b8d2dbfa562869deb58c88583951b58">cutlass::gemm::device::GemmBatched< 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 >::WarpShape</a></div><div class="ttdeci">WarpShape WarpShape</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:230</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_add5d679e0acf0813a52c209d2448e81b"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#add5d679e0acf0813a52c209d2448e81b">cutlass::gemm::device::GemmBatched< 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 >::ArchTag</a></div><div class="ttdeci">ArchTag ArchTag</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:228</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd_html_af623ca54d9554cdfafc09af7a22cdd62"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html#af623ca54d9554cdfafc09af7a22cdd62">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::LayoutA</a></div><div class="ttdeci">LayoutA_ LayoutA</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:524</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a45530b940ca86ce39cfc943da5713d80"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a45530b940ca86ce39cfc943da5713d80">cutlass::gemm::device::GemmBatched::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> device/gemm_batched.h:448</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_aa9e30e41627595590421d8b53941b2b2"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#aa9e30e41627595590421d8b53941b2b2">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::ref_C</a></div><div class="ttdeci">TensorRef< ElementC const, LayoutC > ref_C</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:584</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_aea0c4cd59daee8d3b497be411beb9b3a"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#aea0c4cd59daee8d3b497be411beb9b3a">cutlass::gemm::device::GemmBatched::Arguments::epilogue</a></div><div class="ttdeci">EpilogueOutputOp::Params epilogue</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:282</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="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_adb66f3083f56c15578b139b7935452b5"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#adb66f3083f56c15578b139b7935452b5">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::batch_count</a></div><div class="ttdeci">int batch_count</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:589</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_a35a2fb2e9ad63c316ac6fbb1cc8cf53a"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#a35a2fb2e9ad63c316ac6fbb1cc8cf53a">cutlass::gemm::device::GemmBatched::GemmKernel</a></div><div class="ttdeci">kernel::GemmBatched< typename DefaultGemmKernel::Mma, typename DefaultGemmKernel::Epilogue, ThreadblockSwizzle > GemmKernel</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:264</div></div>
|
|
<div class="ttc" id="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html">cutlass::gemm::device::GemmBatched</a></div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:213</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="classcutlass_1_1gemm_1_1device_1_1GemmBatched_html_af35efd1f40deeb9d8e295f700fa84dbd"><div class="ttname"><a href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html#af35efd1f40deeb9d8e295f700fa84dbd">cutlass::gemm::device::GemmBatched::LayoutC</a></div><div class="ttdeci">LayoutC_ LayoutC</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:223</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments_html_a55f32be45559dbf84dcc2db26784f625"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html#a55f32be45559dbf84dcc2db26784f625">cutlass::gemm::device::GemmBatched::Arguments::ref_A</a></div><div class="ttdeci">TensorRef< ElementA const, LayoutA > ref_A</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:274</div></div>
|
|
<div class="ttc" id="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5_html_ac8830c9ed0e0a8bd7aa2aa4382550a2f"><div class="ttname"><a href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html#ac8830c9ed0e0a8bd7aa2aa4382550a2f">cutlass::gemm::device::GemmBatched< ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ >::Arguments::stride_A</a></div><div class="ttdeci">int64_t stride_A</div><div class="ttdef"><b>Definition:</b> device/gemm_batched.h:581</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>
|