-
Notifications
You must be signed in to change notification settings - Fork 4.3k
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Define HOST_DEVICE_CONSTANT to implement constexpr aggregate host/dev…
…ice constants
- Loading branch information
Showing
1 changed file
with
25 additions
and
0 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,25 @@ | ||
#ifndef FWCore_Utilities_HostDeviceConstant_h | ||
#define FWCore_Utilities_HostDeviceConstant_h | ||
|
||
// The use of host-side consexpr constants in device code is limited to: | ||
// - scalars (other than `long double`) | ||
// - scalar elements of aggregates used inside `constexpr` functions | ||
// | ||
// See https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#constexpr-variables . | ||
// | ||
// In particular, it's not possible to use constexpr scalars by pointer or reference (e.g. std::min() takes arguments | ||
// by reference), or pass constexpr arrays as pointers, or access elements of constexpr arrays outside of constexpr | ||
// functions. | ||
// | ||
// The workaround is to define a macro that evaluates to "constexpr" on the host, and "__device__ constexpr" on the | ||
// device. Such macro can be used to declare aggregate objects that are available both on the host and on the device. | ||
// Note these objects may be at different memory addresses on the host and device, so their pointers will be different | ||
// -- but the actual values should be the same. | ||
|
||
#ifdef __CUDA_ARCH__ | ||
#define HOST_DEVICE_CONSTANT __device__ constexpr | ||
#else | ||
#define HOST_DEVICE_CONSTANT constexpr | ||
#endif | ||
|
||
#endif // FWCore_Utilities_HostDeviceConstant_h |