FWCore/Utilities/interface/HostDeviceConstant.h - master - githubmirror/cmssw

FWCore/Utilities/interface/HostDeviceConstant.h

25 строк1 KB

Andrea Bocci

Fix HOST_DEVICE_CONSTANT macro for AMD GPUs

15 дек 2024, 17:34

15 дек 2024, 17:347cfb11e

#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.

#if defined(__CUDA_ARCH__) or defined(__HIP_DEVICE_COMPILE__)
#define HOST_DEVICE_CONSTANT __device__ constexpr
#else
#define HOST_DEVICE_CONSTANT constexpr
#endif

#endif  // FWCore_Utilities_HostDeviceConstant_h