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
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
|
/*
* Copyright 1993-2023 NVIDIA Corporation. All rights reserved.
*
* NOTICE TO LICENSEE:
*
* This source code and/or documentation ("Licensed Deliverables") are
* subject to NVIDIA intellectual property rights under U.S. and
* international Copyright laws.
*
* These Licensed Deliverables contained herein is PROPRIETARY and
* CONFIDENTIAL to NVIDIA and is being provided under the terms and
* conditions of a form of NVIDIA software license agreement by and
* between NVIDIA and Licensee ("License Agreement") or electronically
* accepted by Licensee. Notwithstanding any terms or conditions to
* the contrary in the License Agreement, reproduction or disclosure
* of the Licensed Deliverables to any third party without the express
* written consent of NVIDIA is prohibited.
*
* NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
* LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE
* SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS
* PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND.
* NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED
* DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY,
* NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
* NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
* LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY
* SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY
* DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,
* WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS
* ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE
* OF THESE LICENSED DELIVERABLES.
*
* U.S. Government End Users. These Licensed Deliverables are a
* "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT
* 1995), consisting of "commercial computer software" and "commercial
* computer software documentation" as such terms are used in 48
* C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government
* only as a commercial end item. Consistent with 48 C.F.R.12.212 and
* 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all
* U.S. Government End Users acquire the Licensed Deliverables with
* only those rights set forth herein.
*
* Any use of the Licensed Deliverables in individual and commercial
* software must include, in the user documentation and internal
* comments to the code, the above Disclaimer and U.S. Government End
* Users Notice.
*/
//NOTE: For NVRTC, these declarations have been moved into the compiler (to reduce compile time)
#define EXCLUDE_FROM_RTC
#if !defined(__SM_60_ATOMIC_FUNCTIONS_H__)
#define __SM_60_ATOMIC_FUNCTIONS_H__
#if defined(__CUDACC_RTC__)
#define __SM_60_ATOMIC_FUNCTIONS_DECL__ __device__
#elif defined(_NVHPC_CUDA)
#define __SM_60_ATOMIC_FUNCTIONS_DECL__ extern __device__ __cudart_builtin__
#else /* __CUDACC_RTC__ */
#define __SM_60_ATOMIC_FUNCTIONS_DECL__ static __inline__ __device__
#endif /* __CUDACC_RTC__ */
#if defined(__cplusplus) && defined(__CUDACC__)
#if defined(_NVHPC_CUDA) || !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600
/*******************************************************************************
* *
* *
* *
*******************************************************************************/
#include "cuda_runtime_api.h"
/* Add !defined(_NVHPC_CUDA) to avoid empty function definition in CUDA
* C++ compiler where the macro __CUDA_ARCH__ is not defined. */
#if !defined(__CUDA_ARCH__) && !defined(_NVHPC_CUDA)
#define __DEF_IF_HOST { }
#else /* !__CUDA_ARCH__ */
#define __DEF_IF_HOST ;
#endif /* __CUDA_ARCH__ */
/*******************************************************************************
* *
* *
* *
*******************************************************************************/
__SM_60_ATOMIC_FUNCTIONS_DECL__ double atomicAdd(double *address, double val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicAdd_block(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicAdd_system(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicAdd_block(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicAdd_system(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicAdd_block(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicAdd_system(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
float atomicAdd_block(float *address, float val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
float atomicAdd_system(float *address, float val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
double atomicAdd_block(double *address, double val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
double atomicAdd_system(double *address, double val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicSub_block(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicSub_system(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicSub_block(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicSub_system(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicExch_block(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicExch_system(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicExch_block(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicExch_system(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicExch_block(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicExch_system(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
float atomicExch_block(float *address, float val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
float atomicExch_system(float *address, float val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicMin_block(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicMin_system(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
long long atomicMin_block(long long *address, long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
long long atomicMin_system(long long *address, long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicMin_block(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicMin_system(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicMin_block(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicMin_system(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicMax_block(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicMax_system(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
long long atomicMax_block(long long *address, long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
long long atomicMax_system(long long *address, long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicMax_block(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicMax_system(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicMax_block(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicMax_system(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicInc_block(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicInc_system(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicDec_block(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicDec_system(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicCAS_block(int *address, int compare, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicCAS_system(int *address, int compare, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicCAS_block(unsigned int *address, unsigned int compare,
unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicCAS_system(unsigned int *address, unsigned int compare,
unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long int atomicCAS_block(unsigned long long int *address,
unsigned long long int compare,
unsigned long long int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long int atomicCAS_system(unsigned long long int *address,
unsigned long long int compare,
unsigned long long int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicAnd_block(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicAnd_system(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
long long atomicAnd_block(long long *address, long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
long long atomicAnd_system(long long *address, long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicAnd_block(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicAnd_system(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicAnd_block(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicAnd_system(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicOr_block(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicOr_system(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
long long atomicOr_block(long long *address, long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
long long atomicOr_system(long long *address, long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicOr_block(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicOr_system(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicOr_block(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicOr_system(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicXor_block(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
int atomicXor_system(int *address, int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
long long atomicXor_block(long long *address, long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
long long atomicXor_system(long long *address, long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicXor_block(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned int atomicXor_system(unsigned int *address, unsigned int val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicXor_block(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
__SM_60_ATOMIC_FUNCTIONS_DECL__
unsigned long long atomicXor_system(unsigned long long *address, unsigned long long val) __DEF_IF_HOST
#endif /* !__CUDA_ARCH__ || __CUDA_ARCH__ >= 600 */
#endif /* __cplusplus && __CUDACC__ */
#undef __SM_60_ATOMIC_FUNCTIONS_DECL__
#undef __DEF_IF_HOST
#if !defined(__CUDACC_RTC__) && defined(__CUDA_ARCH__)
#include "sm_60_atomic_functions.hpp"
#endif /* !__CUDACC_RTC__ && defined(__CUDA_ARCH__) */
#endif /* !__SM_60_ATOMIC_FUNCTIONS_H__ */
#undef EXCLUDE_FROM_RTC
|
