Kokkos Core Kernels Package  Version of the Day
Kokkos_Atomic.hpp
Go to the documentation of this file.
1 /*
2 //@HEADER
3 // ************************************************************************
4 //
5 // Kokkos v. 2.0
6 // Copyright (2014) Sandia Corporation
7 //
8 // Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
9 // the U.S. Government retains certain rights in this software.
10 //
11 // Redistribution and use in source and binary forms, with or without
12 // modification, are permitted provided that the following conditions are
13 // met:
14 //
15 // 1. Redistributions of source code must retain the above copyright
16 // notice, this list of conditions and the following disclaimer.
17 //
18 // 2. Redistributions in binary form must reproduce the above copyright
19 // notice, this list of conditions and the following disclaimer in the
20 // documentation and/or other materials provided with the distribution.
21 //
22 // 3. Neither the name of the Corporation nor the names of the
23 // contributors may be used to endorse or promote products derived from
24 // this software without specific prior written permission.
25 //
26 // THIS SOFTWARE IS PROVIDED BY SANDIA CORPORATION "AS IS" AND ANY
27 // EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
28 // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
29 // PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL SANDIA CORPORATION OR THE
30 // CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
31 // EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
32 // PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
33 // PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
34 // LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
35 // NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
36 // SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
37 //
38 // Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
39 //
40 // ************************************************************************
41 //@HEADER
42 */
43 
66 
67 #ifndef KOKKOS_ATOMIC_HPP
68 #define KOKKOS_ATOMIC_HPP
69 
70 #include <Kokkos_Macros.hpp>
71 #include <Kokkos_HostSpace.hpp>
72 #include <impl/Kokkos_Traits.hpp>
73 
74 //----------------------------------------------------------------------------
75 #if defined(_WIN32)
76 #define KOKKOS_ATOMICS_USE_WINDOWS
77 #else
78 #if defined( __CUDA_ARCH__ ) && defined( KOKKOS_HAVE_CUDA )
79 
80 // Compiling NVIDIA device code, must use Cuda atomics:
81 
82 #define KOKKOS_ATOMICS_USE_CUDA
83 
84 #elif ! defined( KOKKOS_ATOMICS_USE_GCC ) && \
85  ! defined( KOKKOS_ATOMICS_USE_INTEL ) && \
86  ! defined( KOKKOS_ATOMICS_USE_OMP31 )
87 
88 // Compiling for non-Cuda atomic implementation has not been pre-selected.
89 // Choose the best implementation for the detected compiler.
90 // Preference: GCC, INTEL, OMP31
91 
92 #if defined( KOKKOS_COMPILER_GNU ) || \
93  defined( KOKKOS_COMPILER_CLANG ) || \
94  ( defined ( KOKKOS_COMPILER_NVCC ) && defined ( __GNUC__ ) )
95 
96 #define KOKKOS_ATOMICS_USE_GCC
97 
98 #elif defined( KOKKOS_COMPILER_INTEL ) || \
99  defined( KOKKOS_COMPILER_CRAYC )
100 
101 #define KOKKOS_ATOMICS_USE_INTEL
102 
103 #elif defined( _OPENMP ) && ( 201107 <= _OPENMP )
104 
105 #define KOKKOS_ATOMICS_USE_OMP31
106 
107 #else
108 
109 #error "KOKKOS_ATOMICS_USE : Unsupported compiler"
110 
111 #endif
112 
113 #endif /* Not pre-selected atomic implementation */
114 #endif
115 
116 //----------------------------------------------------------------------------
117 
118 // Forward decalaration of functions supporting arbitrary sized atomics
119 // This is necessary since Kokkos_Atomic.hpp is internally included very early
120 // through Kokkos_HostSpace.hpp as well as the allocation tracker.
121 #ifdef KOKKOS_HAVE_CUDA
122 namespace Kokkos {
123 namespace Impl {
129 __device__ inline
130 bool lock_address_cuda_space(void* ptr);
131 
138 __device__ inline
139 void unlock_address_cuda_space(void* ptr);
140 }
141 }
142 #endif
143 
144 
145 namespace Kokkos {
146 template <typename T>
147 KOKKOS_INLINE_FUNCTION
148 void atomic_add(volatile T * const dest, const T src);
149 
150 // Atomic increment
151 template<typename T>
152 KOKKOS_INLINE_FUNCTION
153 void atomic_increment(volatile T* a);
154 
155 template<typename T>
156 KOKKOS_INLINE_FUNCTION
157 void atomic_decrement(volatile T* a);
158 }
159 
160 #if ! defined(_WIN32)
161 #include<impl/Kokkos_Atomic_Assembly_X86.hpp>
162 #endif
163 
164 namespace Kokkos {
165 
166 
167 inline
168 const char * atomic_query_version()
169 {
170 #if defined( KOKKOS_ATOMICS_USE_CUDA )
171  return "KOKKOS_ATOMICS_USE_CUDA" ;
172 #elif defined( KOKKOS_ATOMICS_USE_GCC )
173  return "KOKKOS_ATOMICS_USE_GCC" ;
174 #elif defined( KOKKOS_ATOMICS_USE_INTEL )
175  return "KOKKOS_ATOMICS_USE_INTEL" ;
176 #elif defined( KOKKOS_ATOMICS_USE_OMP31 )
177  return "KOKKOS_ATOMICS_USE_OMP31" ;
178 #elif defined( KOKKOS_ATOMICS_USE_WINDOWS )
179  return "KOKKOS_ATOMICS_USE_WINDOWS";
180 #endif
181 }
182 
183 } // namespace Kokkos
184 
185 #ifdef _WIN32
186 #include "impl/Kokkos_Atomic_Windows.hpp"
187 #else
188 //#include "impl/Kokkos_Atomic_Assembly_X86.hpp"
189 
190 //----------------------------------------------------------------------------
191 // Atomic exchange
192 //
193 // template< typename T >
194 // T atomic_exchange( volatile T* const dest , const T val )
195 // { T tmp = *dest ; *dest = val ; return tmp ; }
196 
197 #include "impl/Kokkos_Atomic_Exchange.hpp"
198 
199 //----------------------------------------------------------------------------
200 // Atomic compare-and-exchange
201 //
202 // template<class T>
203 // bool atomic_compare_exchange_strong(volatile T* const dest, const T compare, const T val)
204 // { bool equal = compare == *dest ; if ( equal ) { *dest = val ; } return equal ; }
205 
206 #include "impl/Kokkos_Atomic_Compare_Exchange_Strong.hpp"
207 
208 //----------------------------------------------------------------------------
209 // Atomic fetch and add
210 //
211 // template<class T>
212 // T atomic_fetch_add(volatile T* const dest, const T val)
213 // { T tmp = *dest ; *dest += val ; return tmp ; }
214 
215 #include "impl/Kokkos_Atomic_Fetch_Add.hpp"
216 
217 //----------------------------------------------------------------------------
218 // Atomic fetch and sub
219 //
220 // template<class T>
221 // T atomic_fetch_sub(volatile T* const dest, const T val)
222 // { T tmp = *dest ; *dest -= val ; return tmp ; }
223 
224 #include "impl/Kokkos_Atomic_Fetch_Sub.hpp"
225 
226 //----------------------------------------------------------------------------
227 // Atomic fetch and or
228 //
229 // template<class T>
230 // T atomic_fetch_or(volatile T* const dest, const T val)
231 // { T tmp = *dest ; *dest = tmp | val ; return tmp ; }
232 
233 #include "impl/Kokkos_Atomic_Fetch_Or.hpp"
234 
235 //----------------------------------------------------------------------------
236 // Atomic fetch and and
237 //
238 // template<class T>
239 // T atomic_fetch_and(volatile T* const dest, const T val)
240 // { T tmp = *dest ; *dest = tmp & val ; return tmp ; }
241 
242 #include "impl/Kokkos_Atomic_Fetch_And.hpp"
243 #endif /*Not _WIN32*/
244 
245 //----------------------------------------------------------------------------
246 // Memory fence
247 //
248 // All loads and stores from this thread will be globally consistent before continuing
249 //
250 // void memory_fence() {...};
251 #include "impl/Kokkos_Memory_Fence.hpp"
252 
253 //----------------------------------------------------------------------------
254 // Provide volatile_load and safe_load
255 //
256 // T volatile_load(T const volatile * const ptr);
257 //
258 // T const& safe_load(T const * const ptr);
259 // XEON PHI
260 // T safe_load(T const * const ptr
261 
262 #include "impl/Kokkos_Volatile_Load.hpp"
263 
264 #ifndef _WIN32
265 #include "impl/Kokkos_Atomic_Generic.hpp"
266 #endif
267 //----------------------------------------------------------------------------
268 // This atomic-style macro should be an inlined function, not a macro
269 
270 #if defined( KOKKOS_COMPILER_GNU ) && !defined(__PGIC__)
271 
272  #define KOKKOS_NONTEMPORAL_PREFETCH_LOAD(addr) __builtin_prefetch(addr,0,0)
273  #define KOKKOS_NONTEMPORAL_PREFETCH_STORE(addr) __builtin_prefetch(addr,1,0)
274 
275 #else
276 
277  #define KOKKOS_NONTEMPORAL_PREFETCH_LOAD(addr) ((void)0)
278  #define KOKKOS_NONTEMPORAL_PREFETCH_STORE(addr) ((void)0)
279 
280 #endif
281 
282 //----------------------------------------------------------------------------
283 
284 #endif /* KOKKOS_ATOMIC_HPP */
285