|
178 | 178 | <div class="line"><a name="l00081"></a><span class="lineno"> 81</span>  </div> |
179 | 179 | <div class="line"><a name="l00082"></a><span class="lineno"> 82</span>  <span class="keywordflow">return</span> __longlong_as_double(old);</div> |
180 | 180 | <div class="line"><a name="l00083"></a><span class="lineno"> 83</span>  }</div> |
181 | | -<div class="line"><a name="l00084"></a><span class="lineno"> 84</span>  }</div> |
182 | | -<div class="line"><a name="l00085"></a><span class="lineno"> 85</span> }</div> |
183 | | -<div class="line"><a name="l00086"></a><span class="lineno"> 86</span>  </div> |
184 | | -<div class="line"><a name="l00087"></a><span class="lineno"> 87</span>  </div> |
185 | | -<div class="line"><a name="l00088"></a><span class="lineno"> 88</span> <span class="preprocessor">#endif</span></div> |
| 181 | +<div class="line"><a name="l00084"></a><span class="lineno"> 84</span>  </div> |
| 182 | +<div class="line"><a name="l00085"></a><span class="lineno"> 85</span>  </div> |
| 183 | +<div class="line"><a name="l00086"></a><span class="lineno"> 86</span>  </div> |
| 184 | +<div class="line"><a name="l00087"></a><span class="lineno"> 87</span>  <span class="keyword">static</span> __device__ <span class="keywordtype">float</span> atomicMul(<span class="keywordtype">float</span> *address, <span class="keywordtype">float</span> val) </div> |
| 185 | +<div class="line"><a name="l00088"></a><span class="lineno"> 88</span>  {</div> |
| 186 | +<div class="line"><a name="l00089"></a><span class="lineno"> 89</span>  <span class="keywordtype">int</span> *address_int = (<span class="keywordtype">int</span>*) address;</div> |
| 187 | +<div class="line"><a name="l00090"></a><span class="lineno"> 90</span>  <span class="keywordtype">int</span> old = *address_int;</div> |
| 188 | +<div class="line"><a name="l00091"></a><span class="lineno"> 91</span>  <span class="keywordtype">int</span> assumed;</div> |
| 189 | +<div class="line"><a name="l00092"></a><span class="lineno"> 92</span>  </div> |
| 190 | +<div class="line"><a name="l00093"></a><span class="lineno"> 93</span>  <span class="keywordflow">do</span></div> |
| 191 | +<div class="line"><a name="l00094"></a><span class="lineno"> 94</span>  {</div> |
| 192 | +<div class="line"><a name="l00095"></a><span class="lineno"> 95</span>  assumed = old;</div> |
| 193 | +<div class="line"><a name="l00096"></a><span class="lineno"> 96</span>  old = atomicCAS(address_int, assumed,</div> |
| 194 | +<div class="line"><a name="l00097"></a><span class="lineno"> 97</span>  __float_as_int(val * __float_as_int(assumed)));</div> |
| 195 | +<div class="line"><a name="l00098"></a><span class="lineno"> 98</span>  } <span class="keywordflow">while</span> (old != assumed);</div> |
| 196 | +<div class="line"><a name="l00099"></a><span class="lineno"> 99</span>  </div> |
| 197 | +<div class="line"><a name="l00100"></a><span class="lineno"> 100</span>  <span class="keywordflow">return</span> __int_as_float(old);</div> |
| 198 | +<div class="line"><a name="l00101"></a><span class="lineno"> 101</span>  }</div> |
| 199 | +<div class="line"><a name="l00102"></a><span class="lineno"> 102</span>  </div> |
| 200 | +<div class="line"><a name="l00103"></a><span class="lineno"> 103</span>  <span class="keyword">static</span> __device__ <span class="keywordtype">double</span> atomicMul(<span class="keywordtype">double</span> *address, <span class="keywordtype">double</span> val) </div> |
| 201 | +<div class="line"><a name="l00104"></a><span class="lineno"> 104</span>  {</div> |
| 202 | +<div class="line"><a name="l00105"></a><span class="lineno"> 105</span>  <span class="keywordtype">unsigned</span> <span class="keywordtype">long</span> <span class="keywordtype">long</span> *address_int = (<span class="keywordtype">unsigned</span> <span class="keywordtype">long</span> <span class="keywordtype">long</span>*) address;</div> |
| 203 | +<div class="line"><a name="l00106"></a><span class="lineno"> 106</span>  <span class="keywordtype">unsigned</span> <span class="keywordtype">long</span> <span class="keywordtype">long</span> old = *address_int;</div> |
| 204 | +<div class="line"><a name="l00107"></a><span class="lineno"> 107</span>  <span class="keywordtype">unsigned</span> <span class="keywordtype">long</span> <span class="keywordtype">long</span> assumed;</div> |
| 205 | +<div class="line"><a name="l00108"></a><span class="lineno"> 108</span>  </div> |
| 206 | +<div class="line"><a name="l00109"></a><span class="lineno"> 109</span>  <span class="keywordflow">do</span></div> |
| 207 | +<div class="line"><a name="l00110"></a><span class="lineno"> 110</span>  {</div> |
| 208 | +<div class="line"><a name="l00111"></a><span class="lineno"> 111</span>  assumed = old;</div> |
| 209 | +<div class="line"><a name="l00112"></a><span class="lineno"> 112</span>  old = atomicCAS(address_int, assumed,</div> |
| 210 | +<div class="line"><a name="l00113"></a><span class="lineno"> 113</span>  __double_as_longlong(val * __double_as_longlong(assumed)));</div> |
| 211 | +<div class="line"><a name="l00114"></a><span class="lineno"> 114</span>  } <span class="keywordflow">while</span> (old != assumed);</div> |
| 212 | +<div class="line"><a name="l00115"></a><span class="lineno"> 115</span>  </div> |
| 213 | +<div class="line"><a name="l00116"></a><span class="lineno"> 116</span>  <span class="keywordflow">return</span> __int_as_float(old);</div> |
| 214 | +<div class="line"><a name="l00117"></a><span class="lineno"> 117</span>  }</div> |
| 215 | +<div class="line"><a name="l00118"></a><span class="lineno"> 118</span>  }</div> |
| 216 | +<div class="line"><a name="l00119"></a><span class="lineno"> 119</span> }</div> |
| 217 | +<div class="line"><a name="l00120"></a><span class="lineno"> 120</span>  </div> |
| 218 | +<div class="line"><a name="l00121"></a><span class="lineno"> 121</span>  </div> |
| 219 | +<div class="line"><a name="l00122"></a><span class="lineno"> 122</span> <span class="preprocessor">#endif</span></div> |
186 | 220 | </div><!-- fragment --></div><!-- contents --> |
187 | 221 | </div><!-- doc-content --> |
188 | 222 | <div class="ttc" id="anamespacefml_html"><div class="ttname"><a href="namespacefml.html">fml</a></div><div class="ttdoc">Core namespace.</div><div class="ttdef"><b>Definition:</b> linalgutils.hh:15</div></div> |
|
0 commit comments