alpaka
Abstraction Library for Parallel Kernel Acceleration
Loading...
Searching...
No Matches
include
alpaka
core
common.hpp
Go to the documentation of this file.
1
/* Copyright 2024 Axel Hübl, Benjamin Worpitz, Matthias Werner, Jan Stephan, René Widera, Andrea Bocci, Aurora Perego,
2
* Tim Hanel SPDX-License-Identifier: MPL-2.0
3
*/
4
5
#pragma once
6
7
#include "
alpaka/core/config.hpp
"
8
9
#include <type_traits>
10
11
12
#if ALPAKA_LANG_HIP
13
// HIP defines some keywords like __forceinline__ in header files.
14
# include <hip/hip_runtime.h>
15
#endif
16
17
//! All functions that can be used on an accelerator have to be attributed with ALPAKA_FN_ACC or ALPAKA_FN_HOST_ACC.
18
//!
19
//! \code{.cpp}
20
//! Usage:
21
//! ALPAKA_FN_ACC
22
//! auto add(std::int32_t a, std::int32_t b)
23
//! -> std::int32_t;
24
//! \endcode
25
//! @{
26
#if ALPAKA_LANG_CUDA || ALPAKA_LANG_HIP
27
# define ALPAKA_FN_ACC __device__ __host__
28
# define ALPAKA_FN_HOST_ACC __device__ __host__
29
# define ALPAKA_FN_HOST __host__
30
#else
31
# define ALPAKA_FN_ACC
32
# define ALPAKA_FN_HOST_ACC
33
# define ALPAKA_FN_HOST
34
#endif
35
//! @}
36
37
//! All functions marked with ALPAKA_FN_ACC or ALPAKA_FN_HOST_ACC that are exported to / imported from different
38
//! translation units have to be attributed with ALPAKA_FN_EXTERN. Note that this needs to be applied to both the
39
//! declaration and the definition.
40
//!
41
//! Usage:
42
//! ALPAKA_FN_ACC ALPAKA_FN_EXTERN auto add(std::int32_t a, std::int32_t b) -> std::int32_t;
43
//!
44
//! Warning: If this is used together with the SYCL back-end make sure that your SYCL runtime supports generic
45
//! address spaces. Otherwise it is forbidden to use pointers as parameter or return type for functions marked
46
//! with ALPAKA_FN_EXTERN.
47
#if ALPAKA_LANG_SYCL
48
/*
49
This is required by the SYCL standard, section 5.10.1 "SYCL functions and member functions linkage":
50
51
The default behavior in SYCL applications is that all the definitions and declarations of the functions and member
52
functions are available to the SYCL compiler, in the same translation unit. When this is not the case, all the
53
symbols that need to be exported to a SYCL library or from a C++ library to a SYCL application need to be defined
54
using the macro: SYCL_EXTERNAL.
55
*/
56
# define ALPAKA_FN_EXTERN SYCL_EXTERNAL
57
#else
58
# define ALPAKA_FN_EXTERN
59
#endif
60
61
//! Disable nvcc warning:
62
//! 'calling a __host__ function from __host__ __device__ function.'
63
//! Usage:
64
//! ALPAKA_NO_HOST_ACC_WARNING
65
//! ALPAKA_FN_HOST_ACC function_declaration()
66
//! WARNING: Only use this method if there is no other way.
67
//! Most cases can be solved by #if ALPAKA_ARCH_PTX or #if ALPAKA_LANG_CUDA.
68
#if (ALPAKA_LANG_CUDA && !ALPAKA_COMP_CLANG_CUDA)
69
# if ALPAKA_COMP_MSVC
70
# define ALPAKA_NO_HOST_ACC_WARNING __pragma(hd_warning_disable)
71
# else
72
# define ALPAKA_NO_HOST_ACC_WARNING _Pragma("hd_warning_disable")
73
# endif
74
#else
75
# define ALPAKA_NO_HOST_ACC_WARNING
76
#endif
77
78
//! Macro defining the inline function attribute.
79
//!
80
//! The macro should stay on the left hand side of keywords, e.g. 'static', 'constexpr', 'explicit' or the return type.
81
#if ALPAKA_LANG_CUDA || ALPAKA_LANG_HIP
82
# define ALPAKA_FN_INLINE __forceinline__
83
#elif ALPAKA_COMP_MSVC
84
// TODO: With C++20 [[msvc::forceinline]] can be used.
85
# define ALPAKA_FN_INLINE __forceinline
86
#else
87
// For gcc, clang, and clang-based compilers like Intel icpx
88
# define ALPAKA_FN_INLINE [[gnu::always_inline]] inline
89
#endif
90
91
//! This macro defines a variable lying in global accelerator device memory.
92
//!
93
//! Example:
94
//! ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal<TAcc, int> variable;
95
//!
96
//! Those variables behave like ordinary variables when used in file-scope,
97
//! but inside kernels the get() method must be used to access the variable.
98
//! They are declared inline to resolve to a single instance across multiple
99
//! translation units.
100
//! Like ordinary variables, only one definition is allowed (ODR)
101
//! Failure to do so might lead to linker errors.
102
//!
103
//! In contrast to ordinary variables, you can not define such variables
104
//! as static compilation unit local variables with internal linkage
105
//! because this is forbidden by CUDA.
106
//!
107
//! \attention It is not allowed to initialize the variable together with the declaration.
108
//! To initialize the variable alpaka::memcpy must be used.
109
//! \code{.cpp}
110
//! ALPAKA_STATIC_ACC_MEM_GLOBAL alpaka::DevGlobal<TAcc, int> foo;
111
//!
112
//! struct DeviceMemoryKernel
113
//! {
114
//! ALPAKA_NO_HOST_ACC_WARNING
115
//! template<typename TAcc>
116
//! ALPAKA_FN_ACC void operator()(TAcc const& acc) const
117
//! {
118
//! auto a = foo<TAcc>.get();
119
//! }
120
//! }
121
//!
122
//! void initFoo() {
123
//! auto extent = alpaka::Vec<alpaka::DimInt<1u>, size_t>{1};
124
//! int initialValue = 42;
125
//! alpaka::ViewPlainPtr<DevHost, int, alpaka::DimInt<1u>, size_t> bufHost(&initialValue, devHost, extent);
126
//! alpaka::memcpy(queue, foo<Acc>, bufHost, extent);
127
//! }
128
//! \endcode
129
#if ( \
130
(ALPAKA_LANG_CUDA && ALPAKA_COMP_CLANG_CUDA) || (ALPAKA_LANG_CUDA && ALPAKA_COMP_NVCC && ALPAKA_ARCH_PTX) \
131
|| ALPAKA_LANG_HIP)
132
# if defined(__CUDACC_RDC__) || defined(__CLANG_RDC__)
133
# define ALPAKA_STATIC_ACC_MEM_GLOBAL \
134
template<typename TAcc> \
135
__device__ inline
136
# else
137
# define ALPAKA_STATIC_ACC_MEM_GLOBAL \
138
template<typename TAcc> \
139
__device__ static
140
# endif
141
#else
142
# define ALPAKA_STATIC_ACC_MEM_GLOBAL \
143
template<typename TAcc> \
144
inline
145
#endif
146
147
/** Perfectly forward an instance as argument. */
148
#define ALPAKA_FORWARD(instance) std::forward<decltype(instance)>(instance)
149
150
/** Get the type of instance
151
*
152
* References will be removed which is often required because traits are mostly defined for the type only.
153
*/
154
#define ALPAKA_TYPEOF(...) std::decay_t<decltype(__VA_ARGS__)>
config.hpp
Generated on
for alpaka by
1.16.1
1.16.1