Kokkos Core Kernels Package  Version of the Day
Kokkos_Macros.hpp
1 /*
2 //@HEADER
3 // ************************************************************************
4 //
5 // Kokkos v. 3.0
6 // Copyright (2020) National Technology & Engineering
7 // Solutions of Sandia, LLC (NTESS).
8 //
9 // Under the terms of Contract DE-NA0003525 with NTESS,
10 // the U.S. Government retains certain rights in this software.
11 //
12 // Redistribution and use in source and binary forms, with or without
13 // modification, are permitted provided that the following conditions are
14 // met:
15 //
16 // 1. Redistributions of source code must retain the above copyright
17 // notice, this list of conditions and the following disclaimer.
18 //
19 // 2. Redistributions in binary form must reproduce the above copyright
20 // notice, this list of conditions and the following disclaimer in the
21 // documentation and/or other materials provided with the distribution.
22 //
23 // 3. Neither the name of the Corporation nor the names of the
24 // contributors may be used to endorse or promote products derived from
25 // this software without specific prior written permission.
26 //
27 // THIS SOFTWARE IS PROVIDED BY NTESS "AS IS" AND ANY
28 // EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
29 // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
30 // PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NTESS OR THE
31 // CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
32 // EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
33 // PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
34 // PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
35 // LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
36 // NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
37 // SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
38 //
39 // Questions? Contact Christian R. Trott (crtrott@sandia.gov)
40 //
41 // ************************************************************************
42 //@HEADER
43 */
44 
45 #ifndef KOKKOS_MACROS_HPP
46 #define KOKKOS_MACROS_HPP
47 
48 //----------------------------------------------------------------------------
63 #ifndef KOKKOS_DONT_INCLUDE_CORE_CONFIG_H
64 #include <KokkosCore_config.h>
65 #endif
66 
67 //----------------------------------------------------------------------------
98 //----------------------------------------------------------------------------
99 
100 #if !defined(KOKKOS_ENABLE_THREADS) && !defined(KOKKOS_ENABLE_CUDA) && \
101  !defined(KOKKOS_ENABLE_OPENMP) && !defined(KOKKOS_ENABLE_HPX) && \
102  !defined(KOKKOS_ENABLE_OPENMPTARGET) && !defined(KOKKOS_ENABLE_HIP) && \
103  !defined(KOKKOS_ENABLE_SYCL)
104 #define KOKKOS_INTERNAL_NOT_PARALLEL
105 #endif
106 
107 #define KOKKOS_ENABLE_CXX11_DISPATCH_LAMBDA
108 
109 #include <KokkosCore_Config_SetupBackend.hpp>
110 
111 //----------------------------------------------------------------------------
112 // Mapping compiler built-ins to KOKKOS_COMPILER_*** macros
113 
114 #if defined(__NVCC__)
115 // NVIDIA compiler is being used.
116 // Code is parsed and separated into host and device code.
117 // Host code is compiled again with another compiler.
118 // Device code is compile to 'ptx'.
119 #define KOKKOS_COMPILER_NVCC __NVCC__
120 #endif // #if defined( __NVCC__ )
121 
122 #if !defined(KOKKOS_LAMBDA)
123 #define KOKKOS_LAMBDA [=]
124 #endif
125 
126 #if (defined(KOKKOS_ENABLE_CXX17) || defined(KOKKOS_ENABLE_CXX20)) && \
127  !defined(KOKKOS_CLASS_LAMBDA)
128 #define KOKKOS_CLASS_LAMBDA [ =, *this ]
129 #endif
130 
131 //#if !defined( __CUDA_ARCH__ ) // Not compiling Cuda code to 'ptx'.
132 
133 // Intel compiler for host code.
134 
135 #if defined(__INTEL_COMPILER)
136 #define KOKKOS_COMPILER_INTEL __INTEL_COMPILER
137 #elif defined(__INTEL_LLVM_COMPILER)
138 #define KOKKOS_COMPILER_INTEL __INTEL_LLVM_COMPILER
139 #elif defined(__ICC)
140 // Old define
141 #define KOKKOS_COMPILER_INTEL __ICC
142 #elif defined(__ECC)
143 // Very old define
144 #define KOKKOS_COMPILER_INTEL __ECC
145 #endif
146 
147 // CRAY compiler for host code
148 #if defined(_CRAYC)
149 #define KOKKOS_COMPILER_CRAYC _CRAYC
150 #endif
151 
152 #if defined(__IBMCPP__)
153 // IBM C++
154 #define KOKKOS_COMPILER_IBM __IBMCPP__
155 #elif defined(__IBMC__)
156 #define KOKKOS_COMPILER_IBM __IBMC__
157 #elif defined(__ibmxl_vrm__) // xlclang++
158 #define KOKKOS_COMPILER_IBM __ibmxl_vrm__
159 #endif
160 
161 #if defined(__APPLE_CC__)
162 #define KOKKOS_COMPILER_APPLECC __APPLE_CC__
163 #endif
164 
165 #if defined(__clang__) && !defined(KOKKOS_COMPILER_INTEL) && \
166  !defined(KOKKOS_COMPILER_IBM)
167 #define KOKKOS_COMPILER_CLANG \
168  __clang_major__ * 100 + __clang_minor__ * 10 + __clang_patchlevel__
169 #endif
170 
171 #if !defined(__clang__) && !defined(KOKKOS_COMPILER_INTEL) && defined(__GNUC__)
172 #define KOKKOS_COMPILER_GNU \
173  __GNUC__ * 100 + __GNUC_MINOR__ * 10 + __GNUC_PATCHLEVEL__
174 
175 #if (530 > KOKKOS_COMPILER_GNU)
176 #error "Compiling with GCC version earlier than 5.3.0 is not supported."
177 #endif
178 #endif
179 
180 #if defined(__PGIC__)
181 #define KOKKOS_COMPILER_PGI \
182  __PGIC__ * 100 + __PGIC_MINOR__ * 10 + __PGIC_PATCHLEVEL__
183 
184 #if (1740 > KOKKOS_COMPILER_PGI)
185 #error "Compiling with PGI version earlier than 17.4 is not supported."
186 #endif
187 #endif
188 
189 #if defined(_MSC_VER) && !defined(KOKKOS_COMPILER_INTEL)
190 #define KOKKOS_COMPILER_MSVC _MSC_VER
191 #endif
192 
193 #if defined(_OPENMP)
194 // Compiling with OpenMP.
195 // The value of _OPENMP is an integer value YYYYMM
196 // where YYYY and MM are the year and month designation
197 // of the supported OpenMP API version.
198 #endif // #if defined( _OPENMP )
199 
200 //----------------------------------------------------------------------------
201 // Intel compiler macros
202 
203 #if defined(KOKKOS_COMPILER_INTEL)
204 // FIXME_SYCL
205 #if !defined(KOKKOS_ENABLE_SYCL)
206 #define KOKKOS_ENABLE_PRAGMA_UNROLL 1
207 #define KOKKOS_ENABLE_PRAGMA_LOOPCOUNT 1
208 #define KOKKOS_ENABLE_PRAGMA_VECTOR 1
209 #endif
210 #if (1800 > KOKKOS_COMPILER_INTEL)
211 #define KOKKOS_ENABLE_PRAGMA_SIMD 1
212 #endif
213 
214 // FIXME_SYCL
215 #if !defined(KOKKOS_ENABLE_SYCL)
216 #define KOKKOS_ENABLE_PRAGMA_IVDEP 1
217 #endif
218 
219 #if !defined(KOKKOS_MEMORY_ALIGNMENT)
220 #define KOKKOS_MEMORY_ALIGNMENT 64
221 #endif
222 
223 #define KOKKOS_RESTRICT __restrict__
224 
225 #ifndef KOKKOS_IMPL_ALIGN_PTR
226 #define KOKKOS_IMPL_ALIGN_PTR(size) __attribute__((align_value(size)))
227 #endif
228 
229 #if (1700 > KOKKOS_COMPILER_INTEL)
230 #error "Compiling with Intel version earlier than 17.0 is not supported."
231 #endif
232 
233 #if !defined(KOKKOS_ENABLE_ASM) && !defined(_WIN32)
234 #define KOKKOS_ENABLE_ASM 1
235 #endif
236 
237 #if !defined(KOKKOS_IMPL_FORCEINLINE_FUNCTION)
238 #if !defined(_WIN32)
239 #define KOKKOS_IMPL_FORCEINLINE_FUNCTION inline __attribute__((always_inline))
240 #define KOKKOS_IMPL_FORCEINLINE __attribute__((always_inline))
241 #else
242 #define KOKKOS_IMPL_FORCEINLINE_FUNCTION inline
243 #endif
244 #endif
245 
246 #if defined(KOKKOS_ARCH_AVX512MIC)
247 #define KOKKOS_ENABLE_RFO_PREFETCH 1
248 #if (KOKKOS_COMPILER_INTEL < 1800) && !defined(KOKKOS_KNL_USE_ASM_WORKAROUND)
249 #define KOKKOS_KNL_USE_ASM_WORKAROUND 1
250 #endif
251 #endif
252 
253 #if (1800 > KOKKOS_COMPILER_INTEL)
254 #define KOKKOS_IMPL_INTEL_WORKAROUND_NOEXCEPT_SPECIFICATION_VIRTUAL_FUNCTION
255 #endif
256 
257 #if defined(__MIC__)
258 // Compiling for Xeon Phi
259 #endif
260 #endif
261 
262 //----------------------------------------------------------------------------
263 // Cray compiler macros
264 
265 #if defined(KOKKOS_COMPILER_CRAYC)
266 #endif
267 
268 //----------------------------------------------------------------------------
269 // IBM Compiler macros
270 
271 #if defined(KOKKOS_COMPILER_IBM)
272 #define KOKKOS_ENABLE_PRAGMA_UNROLL 1
273 //#define KOKKOS_ENABLE_PRAGMA_IVDEP 1
274 //#define KOKKOS_ENABLE_PRAGMA_LOOPCOUNT 1
275 //#define KOKKOS_ENABLE_PRAGMA_VECTOR 1
276 //#define KOKKOS_ENABLE_PRAGMA_SIMD 1
277 
278 #if !defined(KOKKOS_ENABLE_ASM)
279 #define KOKKOS_ENABLE_ASM 1
280 #endif
281 #endif
282 
283 //----------------------------------------------------------------------------
284 // CLANG compiler macros
285 
286 #if defined(KOKKOS_COMPILER_CLANG)
287 //#define KOKKOS_ENABLE_PRAGMA_UNROLL 1
288 //#define KOKKOS_ENABLE_PRAGMA_IVDEP 1
289 //#define KOKKOS_ENABLE_PRAGMA_LOOPCOUNT 1
290 //#define KOKKOS_ENABLE_PRAGMA_VECTOR 1
291 //#define KOKKOS_ENABLE_PRAGMA_SIMD 1
292 
293 #if !defined(KOKKOS_IMPL_FORCEINLINE_FUNCTION)
294 #define KOKKOS_IMPL_FORCEINLINE_FUNCTION inline __attribute__((always_inline))
295 #define KOKKOS_IMPL_FORCEINLINE __attribute__((always_inline))
296 #endif
297 
298 #if !defined(KOKKOS_IMPL_ALIGN_PTR)
299 #define KOKKOS_IMPL_ALIGN_PTR(size) __attribute__((aligned(size)))
300 #endif
301 
302 #endif
303 
304 //----------------------------------------------------------------------------
305 // GNU Compiler macros
306 
307 #if defined(KOKKOS_COMPILER_GNU)
308 //#define KOKKOS_ENABLE_PRAGMA_UNROLL 1
309 //#define KOKKOS_ENABLE_PRAGMA_IVDEP 1
310 //#define KOKKOS_ENABLE_PRAGMA_LOOPCOUNT 1
311 //#define KOKKOS_ENABLE_PRAGMA_VECTOR 1
312 //#define KOKKOS_ENABLE_PRAGMA_SIMD 1
313 
314 #if defined(KOKKOS_ARCH_AVX512MIC)
315 #define KOKKOS_ENABLE_RFO_PREFETCH 1
316 #endif
317 
318 #if !defined(KOKKOS_IMPL_FORCEINLINE_FUNCTION)
319 #define KOKKOS_IMPL_FORCEINLINE_FUNCTION inline __attribute__((always_inline))
320 #define KOKKOS_IMPL_FORCEINLINE __attribute__((always_inline))
321 #endif
322 
323 #define KOKKOS_RESTRICT __restrict__
324 
325 #if !defined(KOKKOS_ENABLE_ASM) && !defined(__PGIC__) && \
326  (defined(__amd64) || defined(__amd64__) || defined(__x86_64) || \
327  defined(__x86_64__) || defined(__PPC64__))
328 #define KOKKOS_ENABLE_ASM 1
329 #endif
330 #endif
331 
332 //----------------------------------------------------------------------------
333 
334 #if defined(KOKKOS_COMPILER_PGI)
335 #define KOKKOS_ENABLE_PRAGMA_UNROLL 1
336 #define KOKKOS_ENABLE_PRAGMA_IVDEP 1
337 //#define KOKKOS_ENABLE_PRAGMA_LOOPCOUNT 1
338 #define KOKKOS_ENABLE_PRAGMA_VECTOR 1
339 //#define KOKKOS_ENABLE_PRAGMA_SIMD 1
340 #endif
341 
342 //----------------------------------------------------------------------------
343 
344 #if defined(KOKKOS_COMPILER_NVCC)
345 #if defined(__CUDA_ARCH__)
346 #define KOKKOS_ENABLE_PRAGMA_UNROLL 1
347 #endif
348 #endif
349 
350 //----------------------------------------------------------------------------
351 // Define function marking macros if compiler specific macros are undefined:
352 
353 #if !defined(KOKKOS_IMPL_FORCEINLINE_FUNCTION)
354 #define KOKKOS_IMPL_FORCEINLINE_FUNCTION inline
355 #endif
356 
357 #if !defined(KOKKOS_IMPL_FORCEINLINE)
358 #define KOKKOS_IMPL_FORCEINLINE inline
359 #endif
360 
361 #if !defined(KOKKOS_IMPL_INLINE_FUNCTION)
362 #define KOKKOS_IMPL_INLINE_FUNCTION inline
363 #endif
364 
365 #if !defined(KOKKOS_IMPL_FUNCTION)
366 #define KOKKOS_IMPL_FUNCTION
367 #endif
368 
369 #if !defined(KOKKOS_INLINE_FUNCTION_DELETED)
370 #define KOKKOS_INLINE_FUNCTION_DELETED inline
371 #endif
372 
373 #if !defined(KOKKOS_DEFAULTED_FUNCTION)
374 #define KOKKOS_DEFAULTED_FUNCTION inline
375 #endif
376 
377 #if !defined(KOKKOS_IMPL_HOST_FUNCTION)
378 #define KOKKOS_IMPL_HOST_FUNCTION
379 #endif
380 
381 #if !defined(KOKKOS_IMPL_DEVICE_FUNCTION)
382 #define KOKKOS_IMPL_DEVICE_FUNCTION
383 #endif
384 
385 // Temporary solution for SYCL not supporting printf in kernels.
386 // Might disappear at any point once we have found another solution.
387 #if !defined(KOKKOS_IMPL_DO_NOT_USE_PRINTF)
388 #define KOKKOS_IMPL_DO_NOT_USE_PRINTF(...) printf(__VA_ARGS__)
389 #endif
390 
391 //----------------------------------------------------------------------------
392 // Define final version of functions. This is so that clang tidy can find these
393 // macros more easily
394 #if defined(__clang_analyzer__)
395 #define KOKKOS_FUNCTION \
396  KOKKOS_IMPL_FUNCTION __attribute__((annotate("KOKKOS_FUNCTION")))
397 #define KOKKOS_INLINE_FUNCTION \
398  KOKKOS_IMPL_INLINE_FUNCTION \
399  __attribute__((annotate("KOKKOS_INLINE_FUNCTION")))
400 #define KOKKOS_FORCEINLINE_FUNCTION \
401  KOKKOS_IMPL_FORCEINLINE_FUNCTION \
402  __attribute__((annotate("KOKKOS_FORCEINLINE_FUNCTION")))
403 #else
404 #define KOKKOS_FUNCTION KOKKOS_IMPL_FUNCTION
405 #define KOKKOS_INLINE_FUNCTION KOKKOS_IMPL_INLINE_FUNCTION
406 #define KOKKOS_FORCEINLINE_FUNCTION KOKKOS_IMPL_FORCEINLINE_FUNCTION
407 #endif
408 
409 //----------------------------------------------------------------------------
410 // Define empty macro for restrict if necessary:
411 
412 #if !defined(KOKKOS_RESTRICT)
413 #define KOKKOS_RESTRICT
414 #endif
415 
416 //----------------------------------------------------------------------------
417 // Define Macro for alignment:
418 
419 #if !defined(KOKKOS_MEMORY_ALIGNMENT)
420 #define KOKKOS_MEMORY_ALIGNMENT 64
421 #endif
422 
423 #if !defined(KOKKOS_MEMORY_ALIGNMENT_THRESHOLD)
424 #define KOKKOS_MEMORY_ALIGNMENT_THRESHOLD 1
425 #endif
426 
427 #if !defined(KOKKOS_IMPL_ALIGN_PTR)
428 #define KOKKOS_IMPL_ALIGN_PTR(size) /* */
429 #endif
430 
431 //----------------------------------------------------------------------------
432 // Determine the default execution space for parallel dispatch.
433 // There is zero or one default execution space specified.
434 
435 #if 1 < ((defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_CUDA) ? 1 : 0) + \
436  (defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_HIP) ? 1 : 0) + \
437  (defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_SYCL) ? 1 : 0) + \
438  (defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMPTARGET) ? 1 : 0) + \
439  (defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMP) ? 1 : 0) + \
440  (defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_THREADS) ? 1 : 0) + \
441  (defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_HPX) ? 1 : 0) + \
442  (defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_SERIAL) ? 1 : 0))
443 #error "More than one KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_* specified."
444 #endif
445 
446 // If default is not specified then chose from enabled execution spaces.
447 // Priority: CUDA, HIP, SYCL, OPENMPTARGET, OPENMP, THREADS, HPX, SERIAL
448 #if defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_CUDA)
449 #elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_HIP)
450 #elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_SYCL)
451 #elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMPTARGET)
452 #elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMP)
453 #elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_THREADS)
454 #elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_HPX)
455 #elif defined(KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_SERIAL)
456 #elif defined(KOKKOS_ENABLE_CUDA)
457 #define KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_CUDA
458 #elif defined(KOKKOS_ENABLE_HIP)
459 #define KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_HIP
460 #if defined(__HIP__)
461 // mark that HIP-clang can use __host__ and __device__
462 // as valid overload criteria
463 #define KOKKOS_IMPL_ENABLE_OVERLOAD_HOST_DEVICE
464 #endif
465 #elif defined(KOKKOS_ENABLE_SYCL)
466 #define KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_SYCL
467 #elif defined(KOKKOS_ENABLE_OPENMPTARGET)
468 #define KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMPTARGET
469 #elif defined(KOKKOS_ENABLE_OPENMP)
470 #define KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_OPENMP
471 #elif defined(KOKKOS_ENABLE_THREADS)
472 #define KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_THREADS
473 #elif defined(KOKKOS_ENABLE_HPX)
474 #define KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_HPX
475 #else
476 #define KOKKOS_ENABLE_DEFAULT_DEVICE_TYPE_SERIAL
477 #endif
478 
479 //----------------------------------------------------------------------------
480 // Determine for what space the code is being compiled:
481 
482 #if defined(__CUDACC__) && defined(__CUDA_ARCH__) && defined(KOKKOS_ENABLE_CUDA)
483 #define KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_CUDA
484 #elif defined(__SYCL_DEVICE_ONLY__) && defined(KOKKOS_ENABLE_SYCL)
485 #define KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_SYCL
486 #elif defined(__HIPCC__) && defined(__HIP_DEVICE_COMPILE__) && \
487  defined(KOKKOS_ENABLE_HIP)
488 #define KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_HIP_GPU
489 #else
490 #define KOKKOS_ACTIVE_EXECUTION_MEMORY_SPACE_HOST
491 #endif
492 
493 //----------------------------------------------------------------------------
494 
495 #if (defined(_POSIX_C_SOURCE) && _POSIX_C_SOURCE >= 200112L) || \
496  (defined(_XOPEN_SOURCE) && _XOPEN_SOURCE >= 600)
497 #if defined(KOKKOS_ENABLE_PERFORMANCE_POSIX_MEMALIGN)
498 #define KOKKOS_ENABLE_POSIX_MEMALIGN 1
499 #endif
500 #endif
501 
502 //----------------------------------------------------------------------------
503 // If compiling with CUDA, we must use relocateable device code
504 // to enable the task policy.
505 
506 #if defined(KOKKOS_ENABLE_CUDA)
507 #if defined(KOKKOS_ENABLE_CUDA_RELOCATABLE_DEVICE_CODE)
508 #define KOKKOS_ENABLE_TASKDAG
509 #endif
510 #elif !defined(KOKKOS_ENABLE_HIP) && !defined(KOKKOS_ENABLE_SYCL)
511 #define KOKKOS_ENABLE_TASKDAG
512 #endif
513 
514 #if defined(KOKKOS_ENABLE_CUDA)
515 #define KOKKOS_IMPL_CUDA_VERSION_9_WORKAROUND
516 #if (__CUDA_ARCH__)
517 #define KOKKOS_IMPL_CUDA_SYNCWARP_NEEDS_MASK
518 #endif
519 #endif
520 
521 #define KOKKOS_INVALID_INDEX (~std::size_t(0))
522 
523 #define KOKKOS_IMPL_CTOR_DEFAULT_ARG KOKKOS_INVALID_INDEX
524 
525 #define KOKKOS_CONSTEXPR_14 constexpr
526 #define KOKKOS_DEPRECATED [[deprecated]]
527 #define KOKKOS_DEPRECATED_TRAILING_ATTRIBUTE
528 
529 // DJS 05/28/2019: Bugfix: Issue 2155
530 // Use KOKKOS_ENABLE_CUDA_LDG_INTRINSIC to avoid memory leak in RandomAccess
531 // View
532 #if defined(KOKKOS_ENABLE_CUDA) && !defined(KOKKOS_ENABLE_CUDA_LDG_INTRINSIC)
533 #define KOKKOS_ENABLE_CUDA_LDG_INTRINSIC
534 #endif
535 
536 #if defined(KOKKOS_ENABLE_CXX17) || defined(KOKKOS_ENABLE_CXX20)
537 #define KOKKOS_ATTRIBUTE_NODISCARD [[nodiscard]]
538 #else
539 #define KOKKOS_ATTRIBUTE_NODISCARD
540 #endif
541 
542 #if (defined(KOKKOS_COMPILER_GNU) || defined(KOKKOS_COMPILER_CLANG) || \
543  defined(KOKKOS_COMPILER_INTEL) || defined(KOKKOS_COMPILER_PGI)) && \
544  !defined(KOKKOS_COMPILER_MSVC)
545 #define KOKKOS_IMPL_ENABLE_STACKTRACE
546 #define KOKKOS_IMPL_ENABLE_CXXABI
547 #endif
548 
549 // WORKAROUND for AMD aomp which apparently defines CUDA_ARCH when building for
550 // AMD GPUs with OpenMP Target ???
551 #if defined(__CUDA_ARCH__) && !defined(__CUDACC__) && \
552  !defined(KOKKOS_ENABLE_HIP) && !defined(KOKKOS_ENABLE_CUDA)
553 #undef __CUDA_ARCH__
554 #endif
555 
556 #if defined(KOKKOS_COMPILER_MSVC) && !defined(KOKKOS_COMPILER_CLANG)
557 #define KOKKOS_THREAD_LOCAL __declspec(thread)
558 #else
559 #define KOKKOS_THREAD_LOCAL __thread
560 #endif
561 
562 #if (defined(KOKKOS_IMPL_WINDOWS_CUDA) || defined(KOKKOS_COMPILER_MSVC)) && \
563  !defined(KOKKOS_COMPILER_CLANG)
564 // MSVC (as of 16.5.5 at least) does not do empty base class optimization by
565 // default when there are multiple bases, even though the standard requires it
566 // for standard layout types.
567 #define KOKKOS_IMPL_ENFORCE_EMPTY_BASE_OPTIMIZATION __declspec(empty_bases)
568 #else
569 #define KOKKOS_IMPL_ENFORCE_EMPTY_BASE_OPTIMIZATION
570 #endif
571 
572 #endif // #ifndef KOKKOS_MACROS_HPP