forked from NVIDIA/cutlass
-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy patharray__subbyte_8h_source.html
181 lines (179 loc) · 123 KB
/
array__subbyte_8h_source.html
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
<!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: array_subbyte.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> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="headertitle">
<div class="title">array_subbyte.h</div> </div>
</div><!--header-->
<div class="contents">
<a href="array__subbyte_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="l00030"></a><span class="lineno"> 30</span> <span class="preprocessor">#pragma once</span></div><div class="line"><a name="l00031"></a><span class="lineno"> 31</span> </div><div class="line"><a name="l00032"></a><span class="lineno"> 32</span> <span class="preprocessor">#include "<a class="code" href="cutlass_8h.html">cutlass/cutlass.h</a>"</span></div><div class="line"><a name="l00033"></a><span class="lineno"> 33</span> <span class="preprocessor">#include "<a class="code" href="array_8h.html">cutlass/array.h</a>"</span></div><div class="line"><a name="l00034"></a><span class="lineno"> 34</span> <span class="preprocessor">#include "<a class="code" href="platform_8h.html">cutlass/platform/platform.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="keyword">namespace </span><a class="code" href="namespacecutlass.html">cutlass</a> {</div><div class="line"><a name="l00037"></a><span class="lineno"> 37</span> </div><div class="line"><a name="l00039"></a><span class="lineno"> 39</span> </div><div class="line"><a name="l00041"></a><span class="lineno"> 41</span> <span class="keyword">template</span> <</div><div class="line"><a name="l00042"></a><span class="lineno"> 42</span>  <span class="keyword">typename</span> T,</div><div class="line"><a name="l00043"></a><span class="lineno"> 43</span>  <span class="keywordtype">int</span> N</div><div class="line"><a name="l00044"></a><span class="lineno"> 44</span> ></div><div class="line"><a name="l00045"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html"> 45</a></span> <span class="keyword">class </span>Array<T, N, false> {</div><div class="line"><a name="l00046"></a><span class="lineno"> 46</span> <span class="keyword">public</span>:</div><div class="line"><a name="l00047"></a><span class="lineno"> 47</span> </div><div class="line"><a name="l00048"></a><span class="lineno"> 48</span>  <a class="code" href="platform_8h.html#adde4c9ea91b753491851361a4198c009">static_assert</a>(<a class="code" href="structcutlass_1_1sizeof__bits.html">sizeof_bits<T>::value</a> * N >= 8,</div><div class="line"><a name="l00049"></a><span class="lineno"> 49</span>  <span class="stringliteral">"Array<> specialized for sub-byte types assume the actual stored element size is 1 byte"</span>);</div><div class="line"><a name="l00050"></a><span class="lineno"> 50</span> </div><div class="line"><a name="l00051"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a45932cad6b905c9ab72889c53112d529"> 51</a></span>  <span class="keyword">static</span> <span class="keywordtype">int</span> <span class="keyword">const</span> kSizeBits = <a class="code" href="structcutlass_1_1sizeof__bits.html">sizeof_bits<T>::value</a> * N;</div><div class="line"><a name="l00052"></a><span class="lineno"> 52</span> </div><div class="line"><a name="l00054"></a><span class="lineno"> 54</span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> = <span class="keyword">typename</span> <a class="code" href="structcutlass_1_1platform_1_1conditional.html">platform::conditional</a><</div><div class="line"><a name="l00055"></a><span class="lineno"> 55</span>  ((kSizeBits % 32) != 0),</div><div class="line"><a name="l00056"></a><span class="lineno"> 56</span>  <span class="keyword">typename</span> <a class="code" href="structcutlass_1_1platform_1_1conditional.html">platform::conditional</a><</div><div class="line"><a name="l00057"></a><span class="lineno"> 57</span>  ((kSizeBits % 16) != 0),</div><div class="line"><a name="l00058"></a><span class="lineno"> 58</span>  uint8_t,</div><div class="line"><a name="l00059"></a><span class="lineno"> 59</span>  uint16_t</div><div class="line"><a name="l00060"></a><span class="lineno"> 60</span>  >::type,</div><div class="line"><a name="l00061"></a><span class="lineno"> 61</span>  uint32_t</div><div class="line"><a name="l00062"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82"> 62</a></span>  >::type;</div><div class="line"><a name="l00063"></a><span class="lineno"> 63</span> </div><div class="line"><a name="l00065"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a700940b7ec4aa2c10506b8109b58b709"> 65</a></span>  <span class="keyword">using</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a700940b7ec4aa2c10506b8109b58b709">Element</a> = T;</div><div class="line"><a name="l00066"></a><span class="lineno"> 66</span> </div><div class="line"><a name="l00068"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a4a6f489743eb03c5c97fe6bb3ed2fa22"> 68</a></span>  <span class="keyword">static</span> <span class="keywordtype">int</span> <span class="keyword">const</span> kElementsPerStoredItem = (<span class="keyword">sizeof</span>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a>) * 8) / <a class="code" href="structcutlass_1_1sizeof__bits.html">sizeof_bits<T>::value</a>;</div><div class="line"><a name="l00069"></a><span class="lineno"> 69</span> </div><div class="line"><a name="l00071"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#afbe4f574d87e61bf18ac5b9f5a6ea8aa"> 71</a></span>  <span class="keyword">static</span> <span class="keywordtype">size_t</span> <span class="keyword">const</span> kStorageElements = N / kElementsPerStoredItem;</div><div class="line"><a name="l00072"></a><span class="lineno"> 72</span> </div><div class="line"><a name="l00074"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a56c28da772c3cf49799eeef4ee1eb981"> 74</a></span>  <span class="keyword">static</span> <span class="keywordtype">size_t</span> <span class="keyword">const</span> kElements = N;</div><div class="line"><a name="l00075"></a><span class="lineno"> 75</span> </div><div class="line"><a name="l00077"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a6981c3aa259d3a1cc4818e29fa1d1423"> 77</a></span>  <span class="keyword">static</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> <span class="keyword">const</span> kMask = ((<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a>(1) << <a class="code" href="structcutlass_1_1sizeof__bits.html">sizeof_bits<T>::value</a>) - 1);</div><div class="line"><a name="l00078"></a><span class="lineno"> 78</span> </div><div class="line"><a name="l00079"></a><span class="lineno"> 79</span>  <span class="comment">//</span></div><div class="line"><a name="l00080"></a><span class="lineno"> 80</span>  <span class="comment">// C++ standard members with pointer types removed</span></div><div class="line"><a name="l00081"></a><span class="lineno"> 81</span>  <span class="comment">//</span></div><div class="line"><a name="l00082"></a><span class="lineno"> 82</span> </div><div class="line"><a name="l00083"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ac1a07d3bbf76e850a948c8efe864acdb"> 83</a></span>  <span class="keyword">typedef</span> T <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ac1a07d3bbf76e850a948c8efe864acdb">value_type</a>;</div><div class="line"><a name="l00084"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a714f3275de8a7f9d14f8b04aed45988d"> 84</a></span>  <span class="keyword">typedef</span> <span class="keywordtype">size_t</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a714f3275de8a7f9d14f8b04aed45988d">size_type</a>;</div><div class="line"><a name="l00085"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#af8dd11bf19216707ab3340b66833c9c9"> 85</a></span>  <span class="keyword">typedef</span> ptrdiff_t <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#af8dd11bf19216707ab3340b66833c9c9">difference_type</a>;</div><div class="line"><a name="l00086"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2a77712281a0ddbf880a4f6fb9aa2ea3"> 86</a></span>  <span class="keyword">typedef</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ac1a07d3bbf76e850a948c8efe864acdb">value_type</a> *<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2a77712281a0ddbf880a4f6fb9aa2ea3">pointer</a>;</div><div class="line"><a name="l00087"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a8a90423fc5483b3ee1d31f377321e9e0"> 87</a></span>  <span class="keyword">typedef</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ac1a07d3bbf76e850a948c8efe864acdb">value_type</a> <span class="keyword">const</span> *<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a8a90423fc5483b3ee1d31f377321e9e0">const_pointer</a>;</div><div class="line"><a name="l00088"></a><span class="lineno"> 88</span> </div><div class="line"><a name="l00089"></a><span class="lineno"> 89</span>  <span class="comment">//</span></div><div class="line"><a name="l00090"></a><span class="lineno"> 90</span>  <span class="comment">// References</span></div><div class="line"><a name="l00091"></a><span class="lineno"> 91</span>  <span class="comment">//</span></div><div class="line"><a name="l00092"></a><span class="lineno"> 92</span> </div><div class="line"><a name="l00094"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html"> 94</a></span>  <span class="keyword">class </span>reference {</div><div class="line"><a name="l00096"></a><span class="lineno"> 96</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> *ptr_;</div><div class="line"><a name="l00097"></a><span class="lineno"> 97</span> </div><div class="line"><a name="l00099"></a><span class="lineno"> 99</span>  <span class="keywordtype">int</span> idx_;</div><div class="line"><a name="l00100"></a><span class="lineno"> 100</span> </div><div class="line"><a name="l00101"></a><span class="lineno"> 101</span>  <span class="keyword">public</span>:</div><div class="line"><a name="l00102"></a><span class="lineno"> 102</span> </div><div class="line"><a name="l00104"></a><span class="lineno"> 104</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00105"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#a257c25bee7fa54ff1d492bc0697b05cc"> 105</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#a257c25bee7fa54ff1d492bc0697b05cc">reference</a>(): ptr_(<a class="code" href="platform_8h.html#ab979d9d4b4923f7c54d6caa6e1a61936">nullptr</a>), idx_(0) { }</div><div class="line"><a name="l00106"></a><span class="lineno"> 106</span> </div><div class="line"><a name="l00108"></a><span class="lineno"> 108</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00109"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#a2c2ac2556e27f48703a9bc1c4e6ed2aa"> 109</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#a2c2ac2556e27f48703a9bc1c4e6ed2aa">reference</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> *ptr, <span class="keywordtype">int</span> idx = 0): ptr_(ptr), idx_(idx) { }</div><div class="line"><a name="l00110"></a><span class="lineno"> 110</span> </div><div class="line"><a name="l00112"></a><span class="lineno"> 112</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00113"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#a629bfbf64481de5252896b45721254ad"> 113</a></span>  reference &<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#a629bfbf64481de5252896b45721254ad">operator=</a>(T x) {</div><div class="line"><a name="l00114"></a><span class="lineno"> 114</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> item = (<span class="keyword">reinterpret_cast<</span><a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> <span class="keyword">const </span>&<span class="keyword">></span>(x) & kMask);</div><div class="line"><a name="l00115"></a><span class="lineno"> 115</span> </div><div class="line"><a name="l00116"></a><span class="lineno"> 116</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> kUpdateMask = <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a>(~(kMask << (idx_ * <a class="code" href="structcutlass_1_1sizeof__bits.html">sizeof_bits<T>::value</a>)));</div><div class="line"><a name="l00117"></a><span class="lineno"> 117</span>  *ptr_ = <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a>(((*ptr_ & kUpdateMask) | (item << idx_ * sizeof_bits<T>::value)));</div><div class="line"><a name="l00118"></a><span class="lineno"> 118</span> </div><div class="line"><a name="l00119"></a><span class="lineno"> 119</span>  <span class="keywordflow">return</span> *<span class="keyword">this</span>;</div><div class="line"><a name="l00120"></a><span class="lineno"> 120</span>  }</div><div class="line"><a name="l00121"></a><span class="lineno"> 121</span> </div><div class="line"><a name="l00122"></a><span class="lineno"> 122</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00123"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#a3bb74e5ee555773803b39cc478af5069"> 123</a></span>  T <span class="keyword">get</span>() <span class="keyword">const</span> {</div><div class="line"><a name="l00124"></a><span class="lineno"> 124</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> item = <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a>((*ptr_ >> (idx_ * <a class="code" href="structcutlass_1_1sizeof__bits.html">sizeof_bits<T>::value</a>)) & kMask);</div><div class="line"><a name="l00125"></a><span class="lineno"> 125</span>  <span class="keywordflow">return</span> <span class="keyword">reinterpret_cast<</span>T <span class="keyword">const </span>&<span class="keyword">></span>(item);</div><div class="line"><a name="l00126"></a><span class="lineno"> 126</span>  }</div><div class="line"><a name="l00127"></a><span class="lineno"> 127</span> </div><div class="line"><a name="l00129"></a><span class="lineno"> 129</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00130"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#ac51b14ff76b80e7f7bf4142a1af01d82"> 130</a></span>  <span class="keyword">operator</span> T()<span class="keyword"> const </span>{</div><div class="line"><a name="l00131"></a><span class="lineno"> 131</span>  <span class="keywordflow">return</span> <span class="keyword">get</span>();</div><div class="line"><a name="l00132"></a><span class="lineno"> 132</span>  }</div><div class="line"><a name="l00133"></a><span class="lineno"> 133</span> </div><div class="line"><a name="l00135"></a><span class="lineno"> 135</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00136"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#ad50d57afc8e33ce406f31dc15564cb3a"> 136</a></span>  <span class="keyword">explicit</span> <span class="keyword">operator</span> int()<span class="keyword"> const </span>{</div><div class="line"><a name="l00137"></a><span class="lineno"> 137</span>  <span class="keywordflow">return</span> int(<span class="keyword">get</span>());</div><div class="line"><a name="l00138"></a><span class="lineno"> 138</span>  }</div><div class="line"><a name="l00139"></a><span class="lineno"> 139</span> </div><div class="line"><a name="l00141"></a><span class="lineno"> 141</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00142"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#ac800b24861e676e21c7d6201338175bc"> 142</a></span>  <span class="keyword">explicit</span> <span class="keyword">operator</span> float()<span class="keyword"> const </span>{</div><div class="line"><a name="l00143"></a><span class="lineno"> 143</span>  <span class="keywordflow">return</span> float(<span class="keyword">get</span>());</div><div class="line"><a name="l00144"></a><span class="lineno"> 144</span>  }</div><div class="line"><a name="l00145"></a><span class="lineno"> 145</span>  };</div><div class="line"><a name="l00146"></a><span class="lineno"> 146</span> </div><div class="line"><a name="l00148"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html"> 148</a></span>  <span class="keyword">class </span>const_reference {</div><div class="line"><a name="l00149"></a><span class="lineno"> 149</span> </div><div class="line"><a name="l00151"></a><span class="lineno"> 151</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> <span class="keyword">const</span> *ptr_;</div><div class="line"><a name="l00152"></a><span class="lineno"> 152</span> </div><div class="line"><a name="l00154"></a><span class="lineno"> 154</span>  <span class="keywordtype">int</span> idx_;</div><div class="line"><a name="l00155"></a><span class="lineno"> 155</span> </div><div class="line"><a name="l00156"></a><span class="lineno"> 156</span>  <span class="keyword">public</span>:</div><div class="line"><a name="l00157"></a><span class="lineno"> 157</span> </div><div class="line"><a name="l00159"></a><span class="lineno"> 159</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00160"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#abf1841f0ac863891efcf23bd5ac57847"> 160</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#abf1841f0ac863891efcf23bd5ac57847">const_reference</a>(): ptr_(<a class="code" href="platform_8h.html#ab979d9d4b4923f7c54d6caa6e1a61936">nullptr</a>), idx_(0) { }</div><div class="line"><a name="l00161"></a><span class="lineno"> 161</span> </div><div class="line"><a name="l00163"></a><span class="lineno"> 163</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00164"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#ac9e3b9e2f5797efbc47e3415aa204079"> 164</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#ac9e3b9e2f5797efbc47e3415aa204079">const_reference</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> <span class="keyword">const</span> *ptr, <span class="keywordtype">int</span> idx = 0): ptr_(ptr), idx_(idx) { }</div><div class="line"><a name="l00165"></a><span class="lineno"> 165</span> </div><div class="line"><a name="l00166"></a><span class="lineno"> 166</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00167"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#a37a90c6f1edcc3d7a916211aa7520cc1"> 167</a></span>  <span class="keyword">const</span> T <span class="keyword">get</span>() <span class="keyword">const</span> {</div><div class="line"><a name="l00168"></a><span class="lineno"> 168</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> item = (*ptr_ >> (idx_ * <a class="code" href="structcutlass_1_1sizeof__bits.html">sizeof_bits<T>::value</a>)) & kMask;</div><div class="line"><a name="l00169"></a><span class="lineno"> 169</span>  <span class="keywordflow">return</span> <span class="keyword">reinterpret_cast<</span>T <span class="keyword">const </span>&<span class="keyword">></span>(item);</div><div class="line"><a name="l00170"></a><span class="lineno"> 170</span>  }</div><div class="line"><a name="l00171"></a><span class="lineno"> 171</span> </div><div class="line"><a name="l00173"></a><span class="lineno"> 173</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00174"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#a7c5f7d59a22d89a7dd5c923d9bcebd97"> 174</a></span>  <span class="keyword">operator</span> T()<span class="keyword"> const </span>{</div><div class="line"><a name="l00175"></a><span class="lineno"> 175</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> item = <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a>(*ptr_ >> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a>(idx_ * <a class="code" href="structcutlass_1_1sizeof__bits.html">sizeof_bits<T>::value</a>)) & kMask);</div><div class="line"><a name="l00176"></a><span class="lineno"> 176</span>  <span class="keywordflow">return</span> <span class="keyword">reinterpret_cast<</span>T <span class="keyword">const </span>&<span class="keyword">></span>(item);</div><div class="line"><a name="l00177"></a><span class="lineno"> 177</span>  }</div><div class="line"><a name="l00178"></a><span class="lineno"> 178</span> </div><div class="line"><a name="l00180"></a><span class="lineno"> 180</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00181"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#a61648afeb4e15881fb001611c37df1ec"> 181</a></span>  <span class="keyword">explicit</span> <span class="keyword">operator</span> int()<span class="keyword"> const </span>{</div><div class="line"><a name="l00182"></a><span class="lineno"> 182</span>  <span class="keywordflow">return</span> int(<span class="keyword">get</span>());</div><div class="line"><a name="l00183"></a><span class="lineno"> 183</span>  }</div><div class="line"><a name="l00184"></a><span class="lineno"> 184</span> </div><div class="line"><a name="l00186"></a><span class="lineno"> 186</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00187"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#afa022bf34a7086c43b5bd45b40c2b25f"> 187</a></span>  <span class="keyword">explicit</span> <span class="keyword">operator</span> float()<span class="keyword"> const </span>{</div><div class="line"><a name="l00188"></a><span class="lineno"> 188</span>  <span class="keywordflow">return</span> float(<span class="keyword">get</span>());</div><div class="line"><a name="l00189"></a><span class="lineno"> 189</span>  }</div><div class="line"><a name="l00190"></a><span class="lineno"> 190</span>  };</div><div class="line"><a name="l00191"></a><span class="lineno"> 191</span> </div><div class="line"><a name="l00192"></a><span class="lineno"> 192</span>  <span class="comment">//</span></div><div class="line"><a name="l00193"></a><span class="lineno"> 193</span>  <span class="comment">// Iterators</span></div><div class="line"><a name="l00194"></a><span class="lineno"> 194</span>  <span class="comment">//</span></div><div class="line"><a name="l00195"></a><span class="lineno"> 195</span> </div><div class="line"><a name="l00197"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html"> 197</a></span>  <span class="keyword">class </span>iterator {</div><div class="line"><a name="l00198"></a><span class="lineno"> 198</span> </div><div class="line"><a name="l00200"></a><span class="lineno"> 200</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> *ptr_;</div><div class="line"><a name="l00201"></a><span class="lineno"> 201</span> </div><div class="line"><a name="l00203"></a><span class="lineno"> 203</span>  <span class="keywordtype">int</span> idx_;</div><div class="line"><a name="l00204"></a><span class="lineno"> 204</span> </div><div class="line"><a name="l00205"></a><span class="lineno"> 205</span>  <span class="keyword">public</span>:</div><div class="line"><a name="l00206"></a><span class="lineno"> 206</span> </div><div class="line"><a name="l00207"></a><span class="lineno"> 207</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00208"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#adb69680f23a0ba9bbe107900fa537228"> 208</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#adb69680f23a0ba9bbe107900fa537228">iterator</a>(): ptr_(<a class="code" href="platform_8h.html#ab979d9d4b4923f7c54d6caa6e1a61936">nullptr</a>), idx_(0) { }</div><div class="line"><a name="l00209"></a><span class="lineno"> 209</span> </div><div class="line"><a name="l00210"></a><span class="lineno"> 210</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00211"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#af7a5f107d79655c43e2f2a42d05a6014"> 211</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#af7a5f107d79655c43e2f2a42d05a6014">iterator</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> *ptr, <span class="keywordtype">int</span> idx = 0): ptr_(ptr), idx_(idx) { }</div><div class="line"><a name="l00212"></a><span class="lineno"> 212</span> </div><div class="line"><a name="l00213"></a><span class="lineno"> 213</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00214"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#a3fa26abda72f9714e39af23bcb5f97df"> 214</a></span>  iterator &<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#a3fa26abda72f9714e39af23bcb5f97df">operator++</a>() {</div><div class="line"><a name="l00215"></a><span class="lineno"> 215</span>  ++idx_;</div><div class="line"><a name="l00216"></a><span class="lineno"> 216</span>  <span class="keywordflow">if</span> (idx_ == kElementsPerStoredItem) {</div><div class="line"><a name="l00217"></a><span class="lineno"> 217</span>  ++ptr_;</div><div class="line"><a name="l00218"></a><span class="lineno"> 218</span>  idx_ = 0;</div><div class="line"><a name="l00219"></a><span class="lineno"> 219</span>  }</div><div class="line"><a name="l00220"></a><span class="lineno"> 220</span>  <span class="keywordflow">return</span> *<span class="keyword">this</span>;</div><div class="line"><a name="l00221"></a><span class="lineno"> 221</span>  }</div><div class="line"><a name="l00222"></a><span class="lineno"> 222</span> </div><div class="line"><a name="l00223"></a><span class="lineno"> 223</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00224"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#a16e57c02d414c3a5591e289c3fd01a22"> 224</a></span>  iterator &<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#a16e57c02d414c3a5591e289c3fd01a22">operator--</a>() {</div><div class="line"><a name="l00225"></a><span class="lineno"> 225</span>  <span class="keywordflow">if</span> (!idx_) {</div><div class="line"><a name="l00226"></a><span class="lineno"> 226</span>  --ptr_;</div><div class="line"><a name="l00227"></a><span class="lineno"> 227</span>  idx_ = kElementsPerStoredItem - 1;</div><div class="line"><a name="l00228"></a><span class="lineno"> 228</span>  }</div><div class="line"><a name="l00229"></a><span class="lineno"> 229</span>  <span class="keywordflow">else</span> {</div><div class="line"><a name="l00230"></a><span class="lineno"> 230</span>  --idx_;</div><div class="line"><a name="l00231"></a><span class="lineno"> 231</span>  }</div><div class="line"><a name="l00232"></a><span class="lineno"> 232</span>  <span class="keywordflow">return</span> *<span class="keyword">this</span>;</div><div class="line"><a name="l00233"></a><span class="lineno"> 233</span>  }</div><div class="line"><a name="l00234"></a><span class="lineno"> 234</span> </div><div class="line"><a name="l00235"></a><span class="lineno"> 235</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00236"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#a3b5e8ff9cb4e7875a6cc26403400d7c3"> 236</a></span>  iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#a3b5e8ff9cb4e7875a6cc26403400d7c3">operator++</a>(<span class="keywordtype">int</span>) {</div><div class="line"><a name="l00237"></a><span class="lineno"> 237</span>  iterator ret(*<span class="keyword">this</span>);</div><div class="line"><a name="l00238"></a><span class="lineno"> 238</span>  ++idx_;</div><div class="line"><a name="l00239"></a><span class="lineno"> 239</span>  <span class="keywordflow">if</span> (idx_ == kElementsPerStoredItem) {</div><div class="line"><a name="l00240"></a><span class="lineno"> 240</span>  ++ptr_;</div><div class="line"><a name="l00241"></a><span class="lineno"> 241</span>  idx_ = 0;</div><div class="line"><a name="l00242"></a><span class="lineno"> 242</span>  }</div><div class="line"><a name="l00243"></a><span class="lineno"> 243</span>  <span class="keywordflow">return</span> ret;</div><div class="line"><a name="l00244"></a><span class="lineno"> 244</span>  }</div><div class="line"><a name="l00245"></a><span class="lineno"> 245</span> </div><div class="line"><a name="l00246"></a><span class="lineno"> 246</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00247"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ae0fd752d82e67eb74ace86cfdaa69020"> 247</a></span>  iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ae0fd752d82e67eb74ace86cfdaa69020">operator--</a>(<span class="keywordtype">int</span>) {</div><div class="line"><a name="l00248"></a><span class="lineno"> 248</span>  iterator ret(*<span class="keyword">this</span>);</div><div class="line"><a name="l00249"></a><span class="lineno"> 249</span>  <span class="keywordflow">if</span> (!idx_) {</div><div class="line"><a name="l00250"></a><span class="lineno"> 250</span>  --ptr_;</div><div class="line"><a name="l00251"></a><span class="lineno"> 251</span>  idx_ = kElementsPerStoredItem - 1;</div><div class="line"><a name="l00252"></a><span class="lineno"> 252</span>  }</div><div class="line"><a name="l00253"></a><span class="lineno"> 253</span>  <span class="keywordflow">else</span> {</div><div class="line"><a name="l00254"></a><span class="lineno"> 254</span>  --idx_;</div><div class="line"><a name="l00255"></a><span class="lineno"> 255</span>  }</div><div class="line"><a name="l00256"></a><span class="lineno"> 256</span>  <span class="keywordflow">return</span> ret;</div><div class="line"><a name="l00257"></a><span class="lineno"> 257</span>  }</div><div class="line"><a name="l00258"></a><span class="lineno"> 258</span> </div><div class="line"><a name="l00259"></a><span class="lineno"> 259</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00260"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ac43741ba9bcacd11dfb91fe02c57bef5"> 260</a></span>  reference <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ac43741ba9bcacd11dfb91fe02c57bef5">operator*</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00261"></a><span class="lineno"> 261</span>  <span class="keywordflow">return</span> reference(ptr_, idx_);</div><div class="line"><a name="l00262"></a><span class="lineno"> 262</span>  }</div><div class="line"><a name="l00263"></a><span class="lineno"> 263</span> </div><div class="line"><a name="l00264"></a><span class="lineno"> 264</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00265"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ac6173c654b3cc22cb357e5ad847dffc9"> 265</a></span>  <span class="keywordtype">bool</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ac6173c654b3cc22cb357e5ad847dffc9">operator==</a>(iterator <span class="keyword">const</span> &other)<span class="keyword"> const </span>{</div><div class="line"><a name="l00266"></a><span class="lineno"> 266</span>  <span class="keywordflow">return</span> ptr_ == other.ptr_ && idx_ == other.idx_;</div><div class="line"><a name="l00267"></a><span class="lineno"> 267</span>  }</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>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00270"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ae72e0d7919ac6d40e1d4f8ce5458af1e"> 270</a></span>  <span class="keywordtype">bool</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ae72e0d7919ac6d40e1d4f8ce5458af1e">operator!=</a>(iterator <span class="keyword">const</span> &other)<span class="keyword"> const </span>{</div><div class="line"><a name="l00271"></a><span class="lineno"> 271</span>  <span class="keywordflow">return</span> !(*<span class="keyword">this</span> == other);</div><div class="line"><a name="l00272"></a><span class="lineno"> 272</span>  }</div><div class="line"><a name="l00273"></a><span class="lineno"> 273</span>  };</div><div class="line"><a name="l00274"></a><span class="lineno"> 274</span> </div><div class="line"><a name="l00276"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html"> 276</a></span>  <span class="keyword">class </span>const_iterator {</div><div class="line"><a name="l00277"></a><span class="lineno"> 277</span> </div><div class="line"><a name="l00279"></a><span class="lineno"> 279</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> <span class="keyword">const</span> *ptr_;</div><div class="line"><a name="l00280"></a><span class="lineno"> 280</span> </div><div class="line"><a name="l00282"></a><span class="lineno"> 282</span>  <span class="keywordtype">int</span> idx_;</div><div class="line"><a name="l00283"></a><span class="lineno"> 283</span> </div><div class="line"><a name="l00284"></a><span class="lineno"> 284</span>  <span class="keyword">public</span>:</div><div class="line"><a name="l00285"></a><span class="lineno"> 285</span> </div><div class="line"><a name="l00286"></a><span class="lineno"> 286</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00287"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a2baacc6de7180213621a2d6b2328ca7d"> 287</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a2baacc6de7180213621a2d6b2328ca7d">const_iterator</a>(): ptr_(<a class="code" href="platform_8h.html#ab979d9d4b4923f7c54d6caa6e1a61936">nullptr</a>), idx_(0) { }</div><div class="line"><a name="l00288"></a><span class="lineno"> 288</span> </div><div class="line"><a name="l00289"></a><span class="lineno"> 289</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00290"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a273a0ea9cf66fac0787e90339fd49371"> 290</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a273a0ea9cf66fac0787e90339fd49371">const_iterator</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> <span class="keyword">const</span> *ptr, <span class="keywordtype">int</span> idx = 0): ptr_(ptr), idx_(idx) { }</div><div class="line"><a name="l00291"></a><span class="lineno"> 291</span> </div><div class="line"><a name="l00292"></a><span class="lineno"> 292</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00293"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#adcdcdf49b5d8e3ed801e2555c4f02b99"> 293</a></span>  iterator &<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#adcdcdf49b5d8e3ed801e2555c4f02b99">operator++</a>() {</div><div class="line"><a name="l00294"></a><span class="lineno"> 294</span>  ++idx_;</div><div class="line"><a name="l00295"></a><span class="lineno"> 295</span>  <span class="keywordflow">if</span> (idx_ == kElementsPerStoredItem) {</div><div class="line"><a name="l00296"></a><span class="lineno"> 296</span>  ++ptr_;</div><div class="line"><a name="l00297"></a><span class="lineno"> 297</span>  idx_ = 0;</div><div class="line"><a name="l00298"></a><span class="lineno"> 298</span>  }</div><div class="line"><a name="l00299"></a><span class="lineno"> 299</span>  <span class="keywordflow">return</span> *<span class="keyword">this</span>;</div><div class="line"><a name="l00300"></a><span class="lineno"> 300</span>  }</div><div class="line"><a name="l00301"></a><span class="lineno"> 301</span> </div><div class="line"><a name="l00302"></a><span class="lineno"> 302</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00303"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#aa2c9f9bb9601208bd784bdc821b62f3a"> 303</a></span>  iterator &<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#aa2c9f9bb9601208bd784bdc821b62f3a">operator--</a>() {</div><div class="line"><a name="l00304"></a><span class="lineno"> 304</span>  <span class="keywordflow">if</span> (!idx_) {</div><div class="line"><a name="l00305"></a><span class="lineno"> 305</span>  --ptr_;</div><div class="line"><a name="l00306"></a><span class="lineno"> 306</span>  idx_ = kElementsPerStoredItem - 1;</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>  <span class="keywordflow">else</span> {</div><div class="line"><a name="l00309"></a><span class="lineno"> 309</span>  --idx_;</div><div class="line"><a name="l00310"></a><span class="lineno"> 310</span>  }</div><div class="line"><a name="l00311"></a><span class="lineno"> 311</span>  <span class="keywordflow">return</span> *<span class="keyword">this</span>;</div><div class="line"><a name="l00312"></a><span class="lineno"> 312</span>  }</div><div class="line"><a name="l00313"></a><span class="lineno"> 313</span> </div><div class="line"><a name="l00314"></a><span class="lineno"> 314</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00315"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a4094d6ae6bb6ade0f850ce96870bbc37"> 315</a></span>  iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a4094d6ae6bb6ade0f850ce96870bbc37">operator++</a>(<span class="keywordtype">int</span>) {</div><div class="line"><a name="l00316"></a><span class="lineno"> 316</span>  iterator ret(*<span class="keyword">this</span>);</div><div class="line"><a name="l00317"></a><span class="lineno"> 317</span>  ++idx_;</div><div class="line"><a name="l00318"></a><span class="lineno"> 318</span>  <span class="keywordflow">if</span> (idx_ == kElementsPerStoredItem) {</div><div class="line"><a name="l00319"></a><span class="lineno"> 319</span>  ++ptr_;</div><div class="line"><a name="l00320"></a><span class="lineno"> 320</span>  idx_ = 0;</div><div class="line"><a name="l00321"></a><span class="lineno"> 321</span>  }</div><div class="line"><a name="l00322"></a><span class="lineno"> 322</span>  <span class="keywordflow">return</span> ret;</div><div class="line"><a name="l00323"></a><span class="lineno"> 323</span>  }</div><div class="line"><a name="l00324"></a><span class="lineno"> 324</span> </div><div class="line"><a name="l00325"></a><span class="lineno"> 325</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00326"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a3eebbf306ba37383e98360c0aa882e34"> 326</a></span>  iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a3eebbf306ba37383e98360c0aa882e34">operator--</a>(<span class="keywordtype">int</span>) {</div><div class="line"><a name="l00327"></a><span class="lineno"> 327</span>  iterator ret(*<span class="keyword">this</span>);</div><div class="line"><a name="l00328"></a><span class="lineno"> 328</span>  <span class="keywordflow">if</span> (!idx_) {</div><div class="line"><a name="l00329"></a><span class="lineno"> 329</span>  --ptr_;</div><div class="line"><a name="l00330"></a><span class="lineno"> 330</span>  idx_ = kElementsPerStoredItem - 1;</div><div class="line"><a name="l00331"></a><span class="lineno"> 331</span>  }</div><div class="line"><a name="l00332"></a><span class="lineno"> 332</span>  <span class="keywordflow">else</span> {</div><div class="line"><a name="l00333"></a><span class="lineno"> 333</span>  --idx_;</div><div class="line"><a name="l00334"></a><span class="lineno"> 334</span>  }</div><div class="line"><a name="l00335"></a><span class="lineno"> 335</span>  <span class="keywordflow">return</span> ret;</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>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00339"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a36aa6aa70a9536a7d2750d83d53f39f3"> 339</a></span>  const_reference <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a36aa6aa70a9536a7d2750d83d53f39f3">operator*</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00340"></a><span class="lineno"> 340</span>  <span class="keywordflow">return</span> const_reference(ptr_, idx_);</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> </div><div class="line"><a name="l00343"></a><span class="lineno"> 343</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00344"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a42dd93a0f0df4ec86de4880fa9cc5843"> 344</a></span>  <span class="keywordtype">bool</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a42dd93a0f0df4ec86de4880fa9cc5843">operator==</a>(iterator <span class="keyword">const</span> &other)<span class="keyword"> const </span>{</div><div class="line"><a name="l00345"></a><span class="lineno"> 345</span>  <span class="keywordflow">return</span> ptr_ == other.ptr_ && idx_ == other.idx_;</div><div class="line"><a name="l00346"></a><span class="lineno"> 346</span>  }</div><div class="line"><a name="l00347"></a><span class="lineno"> 347</span> </div><div class="line"><a name="l00348"></a><span class="lineno"> 348</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00349"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#ad8a6c87e370a53e7ff783ee4ad3d1198"> 349</a></span>  <span class="keywordtype">bool</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#ad8a6c87e370a53e7ff783ee4ad3d1198">operator!=</a>(iterator <span class="keyword">const</span> &other)<span class="keyword"> const </span>{</div><div class="line"><a name="l00350"></a><span class="lineno"> 350</span>  <span class="keywordflow">return</span> !(*<span class="keyword">this</span> == other);</div><div class="line"><a name="l00351"></a><span class="lineno"> 351</span>  }</div><div class="line"><a name="l00352"></a><span class="lineno"> 352</span>  };</div><div class="line"><a name="l00353"></a><span class="lineno"> 353</span> </div><div class="line"><a name="l00355"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator.html"> 355</a></span>  <span class="keyword">class </span>reverse_iterator {</div><div class="line"><a name="l00356"></a><span class="lineno"> 356</span> </div><div class="line"><a name="l00358"></a><span class="lineno"> 358</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> *ptr_;</div><div class="line"><a name="l00359"></a><span class="lineno"> 359</span> </div><div class="line"><a name="l00361"></a><span class="lineno"> 361</span>  <span class="keywordtype">int</span> idx_;</div><div class="line"><a name="l00362"></a><span class="lineno"> 362</span> </div><div class="line"><a name="l00363"></a><span class="lineno"> 363</span>  <span class="keyword">public</span>:</div><div class="line"><a name="l00364"></a><span class="lineno"> 364</span> </div><div class="line"><a name="l00365"></a><span class="lineno"> 365</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00366"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator.html#a539eda60222f630592b9914b51307ea1"> 366</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator.html#a539eda60222f630592b9914b51307ea1">reverse_iterator</a>(): ptr_(<a class="code" href="platform_8h.html#ab979d9d4b4923f7c54d6caa6e1a61936">nullptr</a>), idx_(0) { }</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>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00369"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator.html#a939c336c7c727748d9efcd5efa066a88"> 369</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator.html#a939c336c7c727748d9efcd5efa066a88">reverse_iterator</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> *ptr, <span class="keywordtype">int</span> idx = 0): ptr_(ptr), idx_(idx) { }</div><div class="line"><a name="l00370"></a><span class="lineno"> 370</span> </div><div class="line"><a name="l00371"></a><span class="lineno"> 371</span>  <span class="comment">// TODO</span></div><div class="line"><a name="l00372"></a><span class="lineno"> 372</span>  };</div><div class="line"><a name="l00373"></a><span class="lineno"> 373</span> </div><div class="line"><a name="l00375"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator.html"> 375</a></span>  <span class="keyword">class </span>const_reverse_iterator {</div><div class="line"><a name="l00376"></a><span class="lineno"> 376</span> </div><div class="line"><a name="l00378"></a><span class="lineno"> 378</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> <span class="keyword">const</span> *ptr_;</div><div class="line"><a name="l00379"></a><span class="lineno"> 379</span> </div><div class="line"><a name="l00381"></a><span class="lineno"> 381</span>  <span class="keywordtype">int</span> idx_;</div><div class="line"><a name="l00382"></a><span class="lineno"> 382</span> </div><div class="line"><a name="l00383"></a><span class="lineno"> 383</span>  <span class="keyword">public</span>:</div><div class="line"><a name="l00384"></a><span class="lineno"> 384</span> </div><div class="line"><a name="l00385"></a><span class="lineno"> 385</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00386"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator.html#aae7705a26ea52ebd18d5f5809d816ee2"> 386</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator.html#aae7705a26ea52ebd18d5f5809d816ee2">const_reverse_iterator</a>(): ptr_(<a class="code" href="platform_8h.html#ab979d9d4b4923f7c54d6caa6e1a61936">nullptr</a>), idx_(0) { }</div><div class="line"><a name="l00387"></a><span class="lineno"> 387</span> </div><div class="line"><a name="l00388"></a><span class="lineno"> 388</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00389"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator.html#a4bef88847b70f6bca81dd46bd883373b"> 389</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator.html#a4bef88847b70f6bca81dd46bd883373b">const_reverse_iterator</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> <span class="keyword">const</span> *ptr, <span class="keywordtype">int</span> idx = 0): ptr_(ptr), idx_(idx) { }</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>  <span class="comment">// TODO</span></div><div class="line"><a name="l00392"></a><span class="lineno"> 392</span>  };</div><div class="line"><a name="l00393"></a><span class="lineno"> 393</span> </div><div class="line"><a name="l00394"></a><span class="lineno"> 394</span> <span class="keyword">private</span>:</div><div class="line"><a name="l00395"></a><span class="lineno"> 395</span> </div><div class="line"><a name="l00397"></a><span class="lineno"> 397</span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> storage[kStorageElements];</div><div class="line"><a name="l00398"></a><span class="lineno"> 398</span> </div><div class="line"><a name="l00399"></a><span class="lineno"> 399</span> <span class="keyword">public</span>:</div><div class="line"><a name="l00400"></a><span class="lineno"> 400</span> </div><div class="line"><a name="l00401"></a><span class="lineno"> 401</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00402"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ac37d0c85dd6246ff7e08d12903f49c4d"> 402</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ac37d0c85dd6246ff7e08d12903f49c4d">Array</a>() { }</div><div class="line"><a name="l00403"></a><span class="lineno"> 403</span> </div><div class="line"><a name="l00404"></a><span class="lineno"> 404</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00405"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a5d4667c3c9ebf3322ba94d43421e2577"> 405</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a5d4667c3c9ebf3322ba94d43421e2577">Array</a>(Array <span class="keyword">const</span> &x) {</div><div class="line"><a name="l00406"></a><span class="lineno"> 406</span>  <a class="code" href="cutlass_8h.html#a4b1c9f25ab6eaa25e1f2258dd63e6ce4">CUTLASS_PRAGMA_UNROLL</a></div><div class="line"><a name="l00407"></a><span class="lineno"> 407</span>  <span class="keywordflow">for</span> (<span class="keywordtype">int</span> i = 0; i < int(kStorageElements); ++i) {</div><div class="line"><a name="l00408"></a><span class="lineno"> 408</span>  storage[i] = x.storage[i];</div><div class="line"><a name="l00409"></a><span class="lineno"> 409</span>  }</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> </div><div class="line"><a name="l00413"></a><span class="lineno"> 413</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00414"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a5b84c4dc5257f31108a0598915f03f94"> 414</a></span>  <span class="keywordtype">void</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a5b84c4dc5257f31108a0598915f03f94">clear</a>() {</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>  <a class="code" href="cutlass_8h.html#a4b1c9f25ab6eaa25e1f2258dd63e6ce4">CUTLASS_PRAGMA_UNROLL</a></div><div class="line"><a name="l00417"></a><span class="lineno"> 417</span>  <span class="keywordflow">for</span> (<span class="keywordtype">int</span> i = 0; i < int(kStorageElements); ++i) {</div><div class="line"><a name="l00418"></a><span class="lineno"> 418</span>  storage[i] = <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a>(0);</div><div class="line"><a name="l00419"></a><span class="lineno"> 419</span>  }</div><div class="line"><a name="l00420"></a><span class="lineno"> 420</span>  }</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>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00423"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a6268f2bbbdfc671cf7066ea0ee1bb46f"> 423</a></span>  reference <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a6268f2bbbdfc671cf7066ea0ee1bb46f">at</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a714f3275de8a7f9d14f8b04aed45988d">size_type</a> pos) {</div><div class="line"><a name="l00424"></a><span class="lineno"> 424</span>  <span class="keywordflow">return</span> reference(storage + pos / kElementsPerStoredItem, pos % kElementsPerStoredItem);</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> </div><div class="line"><a name="l00427"></a><span class="lineno"> 427</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00428"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a0443a4af7c9594492bfb8a84bbd12a52"> 428</a></span>  const_reference <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a0443a4af7c9594492bfb8a84bbd12a52">at</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a714f3275de8a7f9d14f8b04aed45988d">size_type</a> pos)<span class="keyword"> const </span>{</div><div class="line"><a name="l00429"></a><span class="lineno"> 429</span>  <span class="keywordflow">return</span> const_reference(storage + pos / kElementsPerStoredItem, pos % kElementsPerStoredItem);</div><div class="line"><a name="l00430"></a><span class="lineno"> 430</span>  }</div><div class="line"><a name="l00431"></a><span class="lineno"> 431</span> </div><div class="line"><a name="l00432"></a><span class="lineno"> 432</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00433"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#aeaeeb7bddb6824adc6feb5ab912d65dc"> 433</a></span>  reference <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#aeaeeb7bddb6824adc6feb5ab912d65dc">operator[]</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a714f3275de8a7f9d14f8b04aed45988d">size_type</a> pos) {</div><div class="line"><a name="l00434"></a><span class="lineno"> 434</span>  <span class="keywordflow">return</span> at(pos);</div><div class="line"><a name="l00435"></a><span class="lineno"> 435</span>  }</div><div class="line"><a name="l00436"></a><span class="lineno"> 436</span> </div><div class="line"><a name="l00437"></a><span class="lineno"> 437</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00438"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a35db1c6ac0d42a486eb3a0a0eee95c80"> 438</a></span>  const_reference <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a35db1c6ac0d42a486eb3a0a0eee95c80">operator[]</a>(<a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a714f3275de8a7f9d14f8b04aed45988d">size_type</a> pos)<span class="keyword"> const </span>{</div><div class="line"><a name="l00439"></a><span class="lineno"> 439</span>  <span class="keywordflow">return</span> at(pos);</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="l00442"></a><span class="lineno"> 442</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00443"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#aa89dd0781c0a81421589182a5402df8b"> 443</a></span>  reference <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#aa89dd0781c0a81421589182a5402df8b">front</a>() {</div><div class="line"><a name="l00444"></a><span class="lineno"> 444</span>  <span class="keywordflow">return</span> at(0);</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="l00447"></a><span class="lineno"> 447</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00448"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ab7ebd33505e48ab3beb6b551e8b762e5"> 448</a></span>  const_reference <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ab7ebd33505e48ab3beb6b551e8b762e5">front</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00449"></a><span class="lineno"> 449</span>  <span class="keywordflow">return</span> at(0);</div><div class="line"><a name="l00450"></a><span class="lineno"> 450</span>  }</div><div class="line"><a name="l00451"></a><span class="lineno"> 451</span> </div><div class="line"><a name="l00452"></a><span class="lineno"> 452</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00453"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a693677ee48012a4d013d55741d38764e"> 453</a></span>  reference <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a693677ee48012a4d013d55741d38764e">back</a>() {</div><div class="line"><a name="l00454"></a><span class="lineno"> 454</span>  <span class="keywordflow">return</span> reference(storage + kStorageElements - 1, kElementsPerStoredItem - 1);</div><div class="line"><a name="l00455"></a><span class="lineno"> 455</span>  }</div><div class="line"><a name="l00456"></a><span class="lineno"> 456</span> </div><div class="line"><a name="l00457"></a><span class="lineno"> 457</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00458"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2c1665d0eff4c1788b0a5a3bfa3bc63e"> 458</a></span>  const_reference <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2c1665d0eff4c1788b0a5a3bfa3bc63e">back</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00459"></a><span class="lineno"> 459</span>  <span class="keywordflow">return</span> const_reference(storage + kStorageElements - 1, kElementsPerStoredItem - 1);</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>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00463"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a1949c8a8c81dc2743328a56ff19fc933"> 463</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2a77712281a0ddbf880a4f6fb9aa2ea3">pointer</a> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a1949c8a8c81dc2743328a56ff19fc933">data</a>() {</div><div class="line"><a name="l00464"></a><span class="lineno"> 464</span>  <span class="keywordflow">return</span> <span class="keyword">reinterpret_cast<</span><a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2a77712281a0ddbf880a4f6fb9aa2ea3">pointer</a><span class="keyword">></span>(storage);</div><div class="line"><a name="l00465"></a><span class="lineno"> 465</span>  }</div><div class="line"><a name="l00466"></a><span class="lineno"> 466</span> </div><div class="line"><a name="l00467"></a><span class="lineno"> 467</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00468"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ab617ed6c9cc6336baf1030713d6dfbbb"> 468</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a8a90423fc5483b3ee1d31f377321e9e0">const_pointer</a> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ab617ed6c9cc6336baf1030713d6dfbbb">data</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00469"></a><span class="lineno"> 469</span>  <span class="keywordflow">return</span> <span class="keyword">reinterpret_cast<</span><a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a8a90423fc5483b3ee1d31f377321e9e0">const_pointer</a><span class="keyword">></span>(storage);</div><div class="line"><a name="l00470"></a><span class="lineno"> 470</span>  }</div><div class="line"><a name="l00471"></a><span class="lineno"> 471</span>  </div><div class="line"><a name="l00472"></a><span class="lineno"> 472</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00473"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a66e2465301e46afebf9e56c4060fb3cb"> 473</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> * <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a66e2465301e46afebf9e56c4060fb3cb">raw_data</a>() {</div><div class="line"><a name="l00474"></a><span class="lineno"> 474</span>  <span class="keywordflow">return</span> storage;</div><div class="line"><a name="l00475"></a><span class="lineno"> 475</span>  }</div><div class="line"><a name="l00476"></a><span class="lineno"> 476</span> </div><div class="line"><a name="l00477"></a><span class="lineno"> 477</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00478"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a16e55f7c4ae1700ae09c2bce137d06ae"> 478</a></span>  <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">Storage</a> <span class="keyword">const</span> * <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a16e55f7c4ae1700ae09c2bce137d06ae">raw_data</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00479"></a><span class="lineno"> 479</span>  <span class="keywordflow">return</span> storage;</div><div class="line"><a name="l00480"></a><span class="lineno"> 480</span>  }</div><div class="line"><a name="l00481"></a><span class="lineno"> 481</span> </div><div class="line"><a name="l00482"></a><span class="lineno"> 482</span> </div><div class="line"><a name="l00483"></a><span class="lineno"> 483</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00484"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a40829269d53d097b5b7bfce32e4afcc4"> 484</a></span>  <a class="code" href="platform_8h.html#a72f0657181cca64b44eb186b707eb380">constexpr</a> <span class="keywordtype">bool</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a40829269d53d097b5b7bfce32e4afcc4">empty</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00485"></a><span class="lineno"> 485</span>  <span class="keywordflow">return</span> !kElements;</div><div class="line"><a name="l00486"></a><span class="lineno"> 486</span>  }</div><div class="line"><a name="l00487"></a><span class="lineno"> 487</span> </div><div class="line"><a name="l00488"></a><span class="lineno"> 488</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00489"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ae1b48e77c8381a8059a09a791d6b8d37"> 489</a></span>  <a class="code" href="platform_8h.html#a72f0657181cca64b44eb186b707eb380">constexpr</a> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a714f3275de8a7f9d14f8b04aed45988d">size_type</a> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ae1b48e77c8381a8059a09a791d6b8d37">size</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00490"></a><span class="lineno"> 490</span>  <span class="keywordflow">return</span> kElements;</div><div class="line"><a name="l00491"></a><span class="lineno"> 491</span>  }</div><div class="line"><a name="l00492"></a><span class="lineno"> 492</span> </div><div class="line"><a name="l00493"></a><span class="lineno"> 493</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00494"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a8f982c95366ce4fda90e35281adfe63c"> 494</a></span>  <a class="code" href="platform_8h.html#a72f0657181cca64b44eb186b707eb380">constexpr</a> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a714f3275de8a7f9d14f8b04aed45988d">size_type</a> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a8f982c95366ce4fda90e35281adfe63c">max_size</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00495"></a><span class="lineno"> 495</span>  <span class="keywordflow">return</span> kElements;</div><div class="line"><a name="l00496"></a><span class="lineno"> 496</span>  }</div><div class="line"><a name="l00497"></a><span class="lineno"> 497</span> </div><div class="line"><a name="l00498"></a><span class="lineno"> 498</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00499"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a1c31d3673a48b2ed275bd56714fbcfbe"> 499</a></span>  <span class="keywordtype">void</span> <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a1c31d3673a48b2ed275bd56714fbcfbe">fill</a>(T <span class="keyword">const</span> &value) {</div><div class="line"><a name="l00500"></a><span class="lineno"> 500</span>  <span class="comment">// TODO</span></div><div class="line"><a name="l00501"></a><span class="lineno"> 501</span>  }</div><div class="line"><a name="l00502"></a><span class="lineno"> 502</span> </div><div class="line"><a name="l00503"></a><span class="lineno"> 503</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00504"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a6e9dbf4a486f07dc72dd5140a7628971"> 504</a></span>  iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a6e9dbf4a486f07dc72dd5140a7628971">begin</a>() {</div><div class="line"><a name="l00505"></a><span class="lineno"> 505</span>  <span class="keywordflow">return</span> iterator(storage);</div><div class="line"><a name="l00506"></a><span class="lineno"> 506</span>  }</div><div class="line"><a name="l00507"></a><span class="lineno"> 507</span> </div><div class="line"><a name="l00508"></a><span class="lineno"> 508</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00509"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a86a56cc907c8566068034ef8294cf7c2"> 509</a></span>  const_iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a86a56cc907c8566068034ef8294cf7c2">cbegin</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00510"></a><span class="lineno"> 510</span>  <span class="keywordflow">return</span> const_iterator(storage);</div><div class="line"><a name="l00511"></a><span class="lineno"> 511</span>  }</div><div class="line"><a name="l00512"></a><span class="lineno"> 512</span> </div><div class="line"><a name="l00513"></a><span class="lineno"> 513</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00514"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a80258d6b5e43ae529cd726f0d4292619"> 514</a></span>  iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a80258d6b5e43ae529cd726f0d4292619">end</a>() {</div><div class="line"><a name="l00515"></a><span class="lineno"> 515</span>  <span class="keywordflow">return</span> iterator(storage + kStorageElements);</div><div class="line"><a name="l00516"></a><span class="lineno"> 516</span>  }</div><div class="line"><a name="l00517"></a><span class="lineno"> 517</span> </div><div class="line"><a name="l00518"></a><span class="lineno"> 518</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00519"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ae6106b72ee9035389afb313801561b16"> 519</a></span>  const_iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ae6106b72ee9035389afb313801561b16">cend</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00520"></a><span class="lineno"> 520</span>  <span class="keywordflow">return</span> const_iterator(storage + kStorageElements);</div><div class="line"><a name="l00521"></a><span class="lineno"> 521</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"> 523</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00524"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2098c88aed61f9b27bac37a083130336"> 524</a></span>  reverse_iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2098c88aed61f9b27bac37a083130336">rbegin</a>() {</div><div class="line"><a name="l00525"></a><span class="lineno"> 525</span>  <span class="keywordflow">return</span> reverse_iterator(storage + kStorageElements);</div><div class="line"><a name="l00526"></a><span class="lineno"> 526</span>  }</div><div class="line"><a name="l00527"></a><span class="lineno"> 527</span> </div><div class="line"><a name="l00528"></a><span class="lineno"> 528</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00529"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a01b9f76c6052dc2467095b91c1ebe34e"> 529</a></span>  const_reverse_iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a01b9f76c6052dc2467095b91c1ebe34e">crbegin</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00530"></a><span class="lineno"> 530</span>  <span class="keywordflow">return</span> const_reverse_iterator(storage + kStorageElements);</div><div class="line"><a name="l00531"></a><span class="lineno"> 531</span>  }</div><div class="line"><a name="l00532"></a><span class="lineno"> 532</span> </div><div class="line"><a name="l00533"></a><span class="lineno"> 533</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00534"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a39c08a75c7cc22fcd296e6c9fefe754e"> 534</a></span>  reverse_iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a39c08a75c7cc22fcd296e6c9fefe754e">rend</a>() {</div><div class="line"><a name="l00535"></a><span class="lineno"> 535</span>  <span class="keywordflow">return</span> reverse_iterator(storage);</div><div class="line"><a name="l00536"></a><span class="lineno"> 536</span>  }</div><div class="line"><a name="l00537"></a><span class="lineno"> 537</span> </div><div class="line"><a name="l00538"></a><span class="lineno"> 538</span>  <a class="code" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a></div><div class="line"><a name="l00539"></a><span class="lineno"><a class="line" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#abbc436f18649c1578ef95eb501872094"> 539</a></span>  const_reverse_iterator <a class="code" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#abbc436f18649c1578ef95eb501872094">crend</a>()<span class="keyword"> const </span>{</div><div class="line"><a name="l00540"></a><span class="lineno"> 540</span>  <span class="keywordflow">return</span> const_reverse_iterator(storage);</div><div class="line"><a name="l00541"></a><span class="lineno"> 541</span>  }</div><div class="line"><a name="l00542"></a><span class="lineno"> 542</span> </div><div class="line"><a name="l00543"></a><span class="lineno"> 543</span>  <span class="comment">//</span></div><div class="line"><a name="l00544"></a><span class="lineno"> 544</span>  <span class="comment">// Comparison operators</span></div><div class="line"><a name="l00545"></a><span class="lineno"> 545</span>  <span class="comment">//</span></div><div class="line"><a name="l00546"></a><span class="lineno"> 546</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> </div><div class="line"><a name="l00550"></a><span class="lineno"> 550</span> </div><div class="line"><a name="l00551"></a><span class="lineno"> 551</span> } <span class="comment">// namespace cutlass</span></div><div class="line"><a name="l00552"></a><span class="lineno"> 552</span> </div><div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference_html_ac9e3b9e2f5797efbc47e3415aa204079"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#ac9e3b9e2f5797efbc47e3415aa204079">cutlass::Array< T, N, false >::const_reference::const_reference</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reference(Storage const *ptr, int idx=0)</div><div class="ttdoc">Ctor. </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:164</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a0443a4af7c9594492bfb8a84bbd12a52"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a0443a4af7c9594492bfb8a84bbd12a52">cutlass::Array< T, N, false >::at</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reference at(size_type pos) const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:428</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a2c1665d0eff4c1788b0a5a3bfa3bc63e"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2c1665d0eff4c1788b0a5a3bfa3bc63e">cutlass::Array< T, N, false >::back</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reference back() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:458</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_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator_html_a539eda60222f630592b9914b51307ea1"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator.html#a539eda60222f630592b9914b51307ea1">cutlass::Array< T, N, false >::reverse_iterator::reverse_iterator</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reverse_iterator()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:366</div></div>
<div class="ttc" id="platform_8h_html_a72f0657181cca64b44eb186b707eb380"><div class="ttname"><a href="platform_8h.html#a72f0657181cca64b44eb186b707eb380">constexpr</a></div><div class="ttdeci">#define constexpr</div><div class="ttdef"><b>Definition:</b> platform.h:137</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator_html_a3eebbf306ba37383e98360c0aa882e34"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a3eebbf306ba37383e98360c0aa882e34">cutlass::Array< T, N, false >::const_iterator::operator--</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator operator--(int)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:326</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a5d4667c3c9ebf3322ba94d43421e2577"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a5d4667c3c9ebf3322ba94d43421e2577">cutlass::Array< T, N, false >::Array</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Array(Array const &x)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:405</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_aeaeeb7bddb6824adc6feb5ab912d65dc"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#aeaeeb7bddb6824adc6feb5ab912d65dc">cutlass::Array< T, N, false >::operator[]</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reference operator[](size_type pos)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:433</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_abbc436f18649c1578ef95eb501872094"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#abbc436f18649c1578ef95eb501872094">cutlass::Array< T, N, false >::crend</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reverse_iterator crend() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:539</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a2098c88aed61f9b27bac37a083130336"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2098c88aed61f9b27bac37a083130336">cutlass::Array< T, N, false >::rbegin</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reverse_iterator rbegin()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:524</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator_html_ac6173c654b3cc22cb357e5ad847dffc9"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ac6173c654b3cc22cb357e5ad847dffc9">cutlass::Array< T, N, false >::iterator::operator==</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE bool operator==(iterator const &other) const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:265</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a714f3275de8a7f9d14f8b04aed45988d"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a714f3275de8a7f9d14f8b04aed45988d">cutlass::Array< T, N, false >::size_type</a></div><div class="ttdeci">size_t size_type</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:84</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator_html_a3fa26abda72f9714e39af23bcb5f97df"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#a3fa26abda72f9714e39af23bcb5f97df">cutlass::Array< T, N, false >::iterator::operator++</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator & operator++()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:214</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a01b9f76c6052dc2467095b91c1ebe34e"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a01b9f76c6052dc2467095b91c1ebe34e">cutlass::Array< T, N, false >::crbegin</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reverse_iterator crbegin() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:529</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a39c08a75c7cc22fcd296e6c9fefe754e"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a39c08a75c7cc22fcd296e6c9fefe754e">cutlass::Array< T, N, false >::rend</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reverse_iterator rend()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:534</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a40829269d53d097b5b7bfce32e4afcc4"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a40829269d53d097b5b7bfce32e4afcc4">cutlass::Array< T, N, false >::empty</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE constexpr bool empty() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:484</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator_html_a273a0ea9cf66fac0787e90339fd49371"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a273a0ea9cf66fac0787e90339fd49371">cutlass::Array< T, N, false >::const_iterator::const_iterator</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_iterator(Storage const *ptr, int idx=0)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:290</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_af8dd11bf19216707ab3340b66833c9c9"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#af8dd11bf19216707ab3340b66833c9c9">cutlass::Array< T, N, false >::difference_type</a></div><div class="ttdeci">ptrdiff_t difference_type</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:85</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_ab617ed6c9cc6336baf1030713d6dfbbb"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ab617ed6c9cc6336baf1030713d6dfbbb">cutlass::Array< T, N, false >::data</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_pointer data() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:468</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator_html_ac43741ba9bcacd11dfb91fe02c57bef5"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ac43741ba9bcacd11dfb91fe02c57bef5">cutlass::Array< T, N, false >::iterator::operator*</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reference operator*() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:260</div></div>
<div class="ttc" id="platform_8h_html"><div class="ttname"><a href="platform_8h.html">platform.h</a></div><div class="ttdoc">C++ features that may be otherwise unimplemented for CUDA device functions. </div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference_html_a257c25bee7fa54ff1d492bc0697b05cc"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#a257c25bee7fa54ff1d492bc0697b05cc">cutlass::Array< T, N, false >::reference::reference</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reference()</div><div class="ttdoc">Default ctor. </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:105</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_ae6106b72ee9035389afb313801561b16"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ae6106b72ee9035389afb313801561b16">cutlass::Array< T, N, false >::cend</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_iterator cend() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:519</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference_html_a2c2ac2556e27f48703a9bc1c4e6ed2aa"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#a2c2ac2556e27f48703a9bc1c4e6ed2aa">cutlass::Array< T, N, false >::reference::reference</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reference(Storage *ptr, int idx=0)</div><div class="ttdoc">Ctor. </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:109</div></div>
<div class="ttc" id="array_8h_html"><div class="ttname"><a href="array_8h.html">array.h</a></div><div class="ttdoc">Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...</div></div>
<div class="ttc" id="cutlass_8h_html_a4b1c9f25ab6eaa25e1f2258dd63e6ce4"><div class="ttname"><a href="cutlass_8h.html#a4b1c9f25ab6eaa25e1f2258dd63e6ce4">CUTLASS_PRAGMA_UNROLL</a></div><div class="ttdeci">#define CUTLASS_PRAGMA_UNROLL</div><div class="ttdef"><b>Definition:</b> cutlass.h:110</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a66e2465301e46afebf9e56c4060fb3cb"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a66e2465301e46afebf9e56c4060fb3cb">cutlass::Array< T, N, false >::raw_data</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Storage * raw_data()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:473</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference_html_abf1841f0ac863891efcf23bd5ac57847"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#abf1841f0ac863891efcf23bd5ac57847">cutlass::Array< T, N, false >::const_reference::const_reference</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reference()</div><div class="ttdoc">Default ctor. </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:160</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a693677ee48012a4d013d55741d38764e"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a693677ee48012a4d013d55741d38764e">cutlass::Array< T, N, false >::back</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reference back()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:453</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a6e9dbf4a486f07dc72dd5140a7628971"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a6e9dbf4a486f07dc72dd5140a7628971">cutlass::Array< T, N, false >::begin</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator begin()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:504</div></div>
<div class="ttc" id="structcutlass_1_1sizeof__bits_html"><div class="ttname"><a href="structcutlass_1_1sizeof__bits.html">cutlass::sizeof_bits</a></div><div class="ttdoc">Defines the size of an element in bits. </div><div class="ttdef"><b>Definition:</b> numeric_types.h:42</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a1c31d3673a48b2ed275bd56714fbcfbe"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a1c31d3673a48b2ed275bd56714fbcfbe">cutlass::Array< T, N, false >::fill</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE void fill(T const &value)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:499</div></div>
<div class="ttc" id="platform_8h_html_ab979d9d4b4923f7c54d6caa6e1a61936"><div class="ttname"><a href="platform_8h.html#ab979d9d4b4923f7c54d6caa6e1a61936">nullptr</a></div><div class="ttdeci">#define nullptr</div><div class="ttdoc">nullptr </div><div class="ttdef"><b>Definition:</b> platform.h:144</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator_html_aa2c9f9bb9601208bd784bdc821b62f3a"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#aa2c9f9bb9601208bd784bdc821b62f3a">cutlass::Array< T, N, false >::const_iterator::operator--</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator & operator--()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:303</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_ac37d0c85dd6246ff7e08d12903f49c4d"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ac37d0c85dd6246ff7e08d12903f49c4d">cutlass::Array< T, N, false >::Array</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Array()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:402</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a700940b7ec4aa2c10506b8109b58b709"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a700940b7ec4aa2c10506b8109b58b709">cutlass::Array< T, N, false >::Element</a></div><div class="ttdeci">T Element</div><div class="ttdoc">Element type. </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:65</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator_html_a36aa6aa70a9536a7d2750d83d53f39f3"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a36aa6aa70a9536a7d2750d83d53f39f3">cutlass::Array< T, N, false >::const_iterator::operator*</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reference operator*() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:339</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator_html_ae0fd752d82e67eb74ace86cfdaa69020"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ae0fd752d82e67eb74ace86cfdaa69020">cutlass::Array< T, N, false >::iterator::operator--</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator operator--(int)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:247</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference_html_a629bfbf64481de5252896b45721254ad"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#a629bfbf64481de5252896b45721254ad">cutlass::Array< T, N, false >::reference::operator=</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reference & operator=(T x)</div><div class="ttdoc">Assignment. </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:113</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a80258d6b5e43ae529cd726f0d4292619"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a80258d6b5e43ae529cd726f0d4292619">cutlass::Array< T, N, false >::end</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator end()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:514</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator_html_a939c336c7c727748d9efcd5efa066a88"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator.html#a939c336c7c727748d9efcd5efa066a88">cutlass::Array< T, N, false >::reverse_iterator::reverse_iterator</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reverse_iterator(Storage *ptr, int idx=0)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:369</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a1949c8a8c81dc2743328a56ff19fc933"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a1949c8a8c81dc2743328a56ff19fc933">cutlass::Array< T, N, false >::data</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE pointer data()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:463</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="structcutlass_1_1platform_1_1conditional_html"><div class="ttname"><a href="structcutlass_1_1platform_1_1conditional.html">cutlass::platform::conditional</a></div><div class="ttdoc">std::conditional (true specialization) </div><div class="ttdef"><b>Definition:</b> platform.h:325</div></div>
<div class="ttc" id="platform_8h_html_adde4c9ea91b753491851361a4198c009"><div class="ttname"><a href="platform_8h.html#adde4c9ea91b753491851361a4198c009">static_assert</a></div><div class="ttdeci">#define static_assert(__e, __m)</div><div class="ttdef"><b>Definition:</b> platform.h:153</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_aa89dd0781c0a81421589182a5402df8b"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#aa89dd0781c0a81421589182a5402df8b">cutlass::Array< T, N, false >::front</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reference front()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:443</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator_html_a4bef88847b70f6bca81dd46bd883373b"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator.html#a4bef88847b70f6bca81dd46bd883373b">cutlass::Array< T, N, false >::const_reverse_iterator::const_reverse_iterator</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reverse_iterator(Storage const *ptr, int idx=0)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:389</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a86a56cc907c8566068034ef8294cf7c2"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a86a56cc907c8566068034ef8294cf7c2">cutlass::Array< T, N, false >::cbegin</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_iterator cbegin() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:509</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator_html_a2baacc6de7180213621a2d6b2328ca7d"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a2baacc6de7180213621a2d6b2328ca7d">cutlass::Array< T, N, false >::const_iterator::const_iterator</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_iterator()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:287</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a6268f2bbbdfc671cf7066ea0ee1bb46f"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a6268f2bbbdfc671cf7066ea0ee1bb46f">cutlass::Array< T, N, false >::at</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE reference at(size_type pos)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:423</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator_html_aae7705a26ea52ebd18d5f5809d816ee2"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator.html#aae7705a26ea52ebd18d5f5809d816ee2">cutlass::Array< T, N, false >::const_reverse_iterator::const_reverse_iterator</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reverse_iterator()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:386</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator_html_a4094d6ae6bb6ade0f850ce96870bbc37"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a4094d6ae6bb6ade0f850ce96870bbc37">cutlass::Array< T, N, false >::const_iterator::operator++</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator operator++(int)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:315</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a16e55f7c4ae1700ae09c2bce137d06ae"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a16e55f7c4ae1700ae09c2bce137d06ae">cutlass::Array< T, N, false >::raw_data</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE Storage const * raw_data() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:478</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a8a90423fc5483b3ee1d31f377321e9e0"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a8a90423fc5483b3ee1d31f377321e9e0">cutlass::Array< T, N, false >::const_pointer</a></div><div class="ttdeci">value_type const * const_pointer</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:87</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_ae1b48e77c8381a8059a09a791d6b8d37"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ae1b48e77c8381a8059a09a791d6b8d37">cutlass::Array< T, N, false >::size</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE constexpr size_type size() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:489</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a8f982c95366ce4fda90e35281adfe63c"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a8f982c95366ce4fda90e35281adfe63c">cutlass::Array< T, N, false >::max_size</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE constexpr size_type max_size() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:494</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a5b84c4dc5257f31108a0598915f03f94"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a5b84c4dc5257f31108a0598915f03f94">cutlass::Array< T, N, false >::clear</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE void clear()</div><div class="ttdoc">Efficient clear method. </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:414</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_ab7ebd33505e48ab3beb6b551e8b762e5"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ab7ebd33505e48ab3beb6b551e8b762e5">cutlass::Array< T, N, false >::front</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reference front() const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:448</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator_html_adb69680f23a0ba9bbe107900fa537228"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#adb69680f23a0ba9bbe107900fa537228">cutlass::Array< T, N, false >::iterator::iterator</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:208</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator_html_af7a5f107d79655c43e2f2a42d05a6014"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#af7a5f107d79655c43e2f2a42d05a6014">cutlass::Array< T, N, false >::iterator::iterator</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator(Storage *ptr, int idx=0)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:211</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a878e152905d602bcdb98e0e6acd8bd82"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a878e152905d602bcdb98e0e6acd8bd82">cutlass::Array< T, N, false >::Storage</a></div><div class="ttdeci">typename platform::conditional< ((kSizeBits%32)!=0), typename platform::conditional< ((kSizeBits%16)!=0), uint8_t, uint16_t >::type, uint32_t >::type Storage</div><div class="ttdoc">Storage type. </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:62</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator_html_a3b5e8ff9cb4e7875a6cc26403400d7c3"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#a3b5e8ff9cb4e7875a6cc26403400d7c3">cutlass::Array< T, N, false >::iterator::operator++</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator operator++(int)</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:236</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator_html_a16e57c02d414c3a5591e289c3fd01a22"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#a16e57c02d414c3a5591e289c3fd01a22">cutlass::Array< T, N, false >::iterator::operator--</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator & operator--()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:224</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator_html_ad8a6c87e370a53e7ff783ee4ad3d1198"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#ad8a6c87e370a53e7ff783ee4ad3d1198">cutlass::Array< T, N, false >::const_iterator::operator!=</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE bool operator!=(iterator const &other) const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:349</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator_html_ae72e0d7919ac6d40e1d4f8ce5458af1e"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#ae72e0d7919ac6d40e1d4f8ce5458af1e">cutlass::Array< T, N, false >::iterator::operator!=</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE bool operator!=(iterator const &other) const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:270</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_ac1a07d3bbf76e850a948c8efe864acdb"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#ac1a07d3bbf76e850a948c8efe864acdb">cutlass::Array< T, N, false >::value_type</a></div><div class="ttdeci">T value_type</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:83</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a2a77712281a0ddbf880a4f6fb9aa2ea3"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a2a77712281a0ddbf880a4f6fb9aa2ea3">cutlass::Array< T, N, false >::pointer</a></div><div class="ttdeci">value_type * pointer</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:86</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator_html_a42dd93a0f0df4ec86de4880fa9cc5843"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#a42dd93a0f0df4ec86de4880fa9cc5843">cutlass::Array< T, N, false >::const_iterator::operator==</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE bool operator==(iterator const &other) const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:344</div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator_html_adcdcdf49b5d8e3ed801e2555c4f02b99"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#adcdcdf49b5d8e3ed801e2555c4f02b99">cutlass::Array< T, N, false >::const_iterator::operator++</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE iterator & operator++()</div><div class="ttdef"><b>Definition:</b> array_subbyte.h:293</div></div>
<div class="ttc" id="cutlass_8h_html"><div class="ttname"><a href="cutlass_8h.html">cutlass.h</a></div><div class="ttdoc">Basic include for CUTLASS. </div></div>
<div class="ttc" id="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_html_a35db1c6ac0d42a486eb3a0a0eee95c80"><div class="ttname"><a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#a35db1c6ac0d42a486eb3a0a0eee95c80">cutlass::Array< T, N, false >::operator[]</a></div><div class="ttdeci">CUTLASS_HOST_DEVICE const_reference operator[](size_type pos) const </div><div class="ttdef"><b>Definition:</b> array_subbyte.h:438</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>