Kokkos Core Kernels Package
Version of the Day
Main Page
Namespaces
Classes
Files
File List
core
src
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
Kokkos
Definition:
Kokkos_Array.hpp:51
Generated by
1.8.11