VecGeom/base/Cuda.h

0001 #ifndef VECGEOM_CUDA_H
0002 #define VECGEOM_CUDA_H
0003
0004 #include <memory>
0005
0006 #include "VecGeom/base/Config.h"
0007
0008 // Keep compact version of the macros.
0009 // clang-format off
0010
0011 #if (defined(__CUDACC__) || defined(__NVCC__))
0012   // Compiling with nvcc
0013   #define VECGEOM_IMPL_NAMESPACE cuda
0014   #define VECGEOM_NAMESPACE ::vecgeom
0015   #define VECGEOM_ALIGNED __align__((64))
0016
0017   namespace wide {
0018    // 2-lane double
0019   struct alignas(16) d2 {
0020     double2 v;
0021     __host__ __device__ d2() : v{0.0, 0.0} {}
0022     __host__ __device__ d2(double a, double b) : v{a, b} {}
0023     __host__ __device__ inline double x() const { return v.x; }
0024     __host__ __device__ inline double y() const { return v.y; }
0025     __host__ __device__ inline void set(double a, double b){ v = make_double2(a, b); }
0026   };
0027
0028   // 4-lane float
0029   struct alignas(16) f4 {
0030     float4 v;
0031     __host__ __device__ f4() : v{0.f, 0.f, 0.f, 0.f} {}
0032     __host__ __device__ f4(float a,float b,float c,float d) : v{a, b, c, d} {}
0033     __host__ __device__ inline float x() const { return v.x; }
0034     __host__ __device__ inline float y() const { return v.y; }
0035     __host__ __device__ inline float z() const { return v.z; }
0036     __host__ __device__ inline float w() const { return v.w; }
0037     __host__ __device__ inline void set(float a, float b, float c, float d) { v = make_float4(a, b, c, d); }
0038   };
0039   } // namespace wide
0040
0041   #define VECGEOM_HOST_FORWARD_DECLARE(X) namespace cxx { X }
0042   #define VECGEOM_DEVICE_FORWARD_DECLARE(X) class __QuietSemi
0043   #define VECGEOM_DEVICE_DECLARE_CONV(classOrStruct,X) class __QuietSemi
0044   #define VECGEOM_DEVICE_DECLARE_NS_CONV(NS,classOrStruct,X,Def) class __QuietSemi
0045   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE(classOrStruct,X,ArgType) class __QuietSemi
0046   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1v(classOrStruct,X,ArgType1,Def1) class __QuietSemi
0047   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1v_1t(classOrStruct,X,ArgType1,ArgType2) class __QuietSemi
0048   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1specv_1t(classOrStruct,X,ArgType1,Def1,ArgType2,Def2,ArgType3) class __QuietSemi
0049   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_2t(classOrStruct,X,ArgType1,ArgType2) class __QuietSemi
0050   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_2v(classOrStruct,X,ArgType1,Def1,ArgType2,Def2) class __QuietSemi
0051   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_2v_1t(classOrStruct,X,ArgType1,Def1,ArgType2,Def2,ArgType3) class __QuietSemi
0052   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1t_2v(classOrStruct,X,ArgType1,ArgType2,Def2,ArgType3,Def3) class __QuietSemi
0053   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_3v(classOrStruct,X,ArgType1,Def1,ArgType2,Def2,ArgType3,Def3) class __QuietSemi
0054   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_4v(classOrStruct,X,ArgType1,Def1,ArgType2,Def2,ArgType3,Def3,ArgType4,Def4) class __QuietSemi
0055   #undef VECGEOM_VC
0056   #undef VECGEOM_CILK
0057   #undef VECGEOM_ROOT
0058   #undef VECGEOM_GEANT4
0059 #else
0060   // Not compiling with NVCC
0061   #define VECGEOM_IMPL_NAMESPACE cxx
0062   #define VECGEOM_NAMESPACE ::vecgeom
0063   namespace wide {
0064   // 2-lane double
0065   struct alignas(16) d2 {
0066     double x_, y_;
0067     d2() : x_(0.0), y_(0.0) {}
0068     d2(double a, double b) : x_(a), y_(b) {}
0069     inline double x() const { return x_; }
0070     inline double y() const { return y_; }
0071     inline void set(double a, double b){ x_=a; y_=b; }
0072   };
0073
0074   // 4-lane float
0075   struct alignas(16) f4 {
0076     float x_, y_, z_, w_;
0077     f4() : x_(0.f), y_(0.f), z_(0.f), w_(0.f) {}
0078     f4(float a, float b, float c, float d) : x_(a), y_(b), z_(c), w_(d) {}
0079     inline float x() const { return x_; }
0080     inline float y() const { return y_; }
0081     inline float z() const { return z_; }
0082     inline float w() const { return w_; }
0083     inline void set(float a, float b, float c, float d){ x_=a; y_=b; z_=c; w_=d; }
0084   };
0085
0086   } // namespace wide
0087   #ifdef VECGEOM_ENABLE_CUDA
0088     // CUDA is enabled, but currently compiling regular C++ code.
0089     // This enables methods that interface between C++ and CUDA environments
0090     #define VECGEOM_CUDA_INTERFACE
0091   #endif
0092   namespace vecgeom {
0093      template <typename DataType> struct kCudaType;
0094      template <typename DataType> using CudaType_t = typename kCudaType<DataType>::type_t;
0095      template <> struct kCudaType<float> { using type_t = float; };
0096      template <> struct kCudaType<double> { using type_t = double; };
0097      template <> struct kCudaType<int> { using type_t = int; };
0098      template <> struct kCudaType<size_t> { using type_t = size_t; };
0099      template <typename DataType> struct kCudaType<DataType*> {
0100        using type_t = CudaType_t<DataType>*;
0101      };
0102      template <typename DataType> struct kCudaType<const DataType> {
0103        using type_t = const CudaType_t<DataType>;
0104      };
0105   }
0106   #define VECGEOM_HOST_FORWARD_DECLARE(X)  class __QuietSemi
0107   #define VECGEOM_DEVICE_FORWARD_DECLARE(X)  namespace cuda { X }  class __QuietSemi
0108
0109   #define VECGEOM_DEVICE_DECLARE_CONV(classOrStruct,X)                 \
0110      namespace cuda { classOrStruct X; }                               \
0111      inline namespace cxx  { classOrStruct X; }                        \
0112      template <> struct kCudaType<cxx::X> { using type_t = cuda::X; }
0113
0114   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE(classOrStruct,X,ArgType) \
0115      namespace cuda { template <ArgType Arg> classOrStruct X; }         \
0116      inline namespace cxx  { template <ArgType Arg> classOrStruct X; }  \
0117      template <ArgType Arg> struct kCudaType<cxx::X<Arg> >              \
0118      { using type_t = cuda::X<CudaType_t<Arg> >; }
0119
0120   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_2t(classOrStruct,X,ArgType1,ArgType2)   \
0121      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }        \
0122      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; } \
0123      template <ArgType1 Arg1,ArgType2 Arg2> struct kCudaType<cxx::X<Arg1,Arg2> >       \
0124      { using type_t = cuda::X<CudaType_t<Arg1>, CudaType_t<Arg2> >; }
0125
0126   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1v_1t(classOrStruct,X,ArgType1,ArgType2) \
0127      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }         \
0128      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }  \
0129      template <ArgType1 Arg1,ArgType2 Arg2> struct kCudaType<cxx::X<Arg1,Arg2> >        \
0130      { using type_t = cuda::X<Arg1, CudaType_t<Arg2> >; }
0131
0132 #ifdef VECGEOM_CUDA_VOLUME_SPECIALIZATION
0133
0134   #define VECGEOM_DEVICE_DECLARE_NS_CONV(NS,classOrStruct,X,Def)               \
0135      namespace cuda { namespace NS { classOrStruct X; } }                      \
0136      inline namespace cxx { namespace NS { classOrStruct X; } }                \
0137      template <> struct kCudaType<cxx::NS::X> { using type_t = cuda::NS::X; }
0138
0139   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1v(classOrStruct,X,ArgType1,Def1) \
0140      namespace cuda { template <ArgType1 Arg1> classOrStruct X; }                \
0141      inline namespace cxx  { template <ArgType1 Arg1> classOrStruct X; }         \
0142      template <ArgType1 Arg1> struct kCudaType<cxx::X<Arg1> >                    \
0143      { using type_t = cuda::X<Arg>; }
0144   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1specv_1t(classOrStruct,X,ArgType1,ArgType2) \
0145      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }             \
0146      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }      \
0147      template <ArgType1 Arg1,ArgType2 Arg2> struct kCudaType<cxx::X<Arg1,Arg2> >            \
0148      { using type_t = cuda::X<Arg1, CudaType_t<Arg2> >; }
0149   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_2v(classOrStruct,X,ArgType1,Def1,ArgType2,Def2) \
0150      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }                \
0151      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }         \
0152      template <ArgType1 Arg1,ArgType2 Arg2> struct kCudaType<cxx::X<Arg1,Arg2> >               \
0153      { using type_t = cuda::X<Arg1,Arg2 >; }
0154   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_2v_1t(classOrStruct,X,ArgType1,Def1,ArgType2,Def2,ArgType3) \
0155      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }              \
0156      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }       \
0157      template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> struct kCudaType<cxx::X<Arg1,Arg2,Arg3> >        \
0158      { using type_t = cuda::X<Arg1, Arg2, CudaType_t<Arg3> >; }
0159   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1t_2v(classOrStruct,X,ArgType1,ArgType2,Def2,ArgType3,Def3) \
0160      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }              \
0161      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }       \
0162      template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> struct kCudaType<cxx::X<Arg1,Arg2,Arg3> >        \
0163      { using type_t = cuda::X<CudaType_t<Arg1>, Arg2, Arg3 >; }
0164   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_3v(classOrStruct,X,ArgType1,Def1,ArgType2,Def2,ArgType3,Def3) \
0165      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }                \
0166      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }         \
0167      template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> struct kCudaType<cxx::X<Arg1,Arg2,Arg3> >          \
0168      { using type_t = cuda::X<Arg1,Arg2,Arg3 >; }
0169
0170   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_4v(classOrStruct,X,ArgType1,Def1,ArgType2,Def2,ArgType3,Def3,ArgType4,Def4) \
0171     namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3,ArgType4 Arg4> classOrStruct X; }                 \
0172     inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3,ArgType4 Arg4> classOrStruct X; }          \
0173     template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3,ArgType4 Arg4> struct kCudaType<cxx::X<Arg1,Arg2,Arg3,Arg4> >      \
0174     { using type_t = cuda::X<Arg1,Arg2,Arg3,Arg4 >; }
0175
0176 #else // VECGEOM_CUDA_VOLUME_SPECIALIZATION
0177
0178   #define VECGEOM_DEVICE_DECLARE_NS_CONV(NS,classOrStruct,X,Def)                 \
0179      namespace cuda { namespace NS { classOrStruct Def; } }                      \
0180      inline namespace cxx { namespace NS { classOrStruct X; } }                  \
0181      template <> struct kCudaType<cxx::NS::X> { using type_t = cuda::NS::Def; }
0182
0183   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1v(classOrStruct,X,ArgType1,Def1) \
0184      namespace cuda { template <ArgType1 Arg1> classOrStruct X; }                \
0185      inline namespace cxx  { template <ArgType1 Arg1> classOrStruct X; }         \
0186      template <ArgType1 Arg1> struct kCudaType<cxx::X<Arg1> >                    \
0187      { using type_t = cuda::X<Def1>; }
0188   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1specv_1t(classOrStruct,X,ArgType1,Def1,ArgType2) \
0189      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }             \
0190      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }      \
0191      template <ArgType1 Arg1,ArgType2 Arg2> struct kCudaType<cxx::X<Arg1,Arg2> >            \
0192      { using type_t = cuda::X<Def1, CudaType_t<Arg2> >; }
0193   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_2v(classOrStruct,X,ArgType1,Def1,ArgType2,Def2) \
0194      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }                \
0195      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2> classOrStruct X; }         \
0196      template <ArgType1 Arg1,ArgType2 Arg2> struct kCudaType<cxx::X<Arg1,Arg2> >               \
0197      { using type_t = cuda::X<Def1, Def2 >; }
0198   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_2v_1t(classOrStruct,X,ArgType1,Def1,ArgType2,Def2,ArgType3) \
0199      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }              \
0200      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }       \
0201      template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> struct kCudaType<cxx::X<Arg1,Arg2,Arg3> >        \
0202      { using type_t = cuda::X<Def1, Def2, CudaType_t<Arg3> >; }
0203   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_1t_2v(classOrStruct,X,ArgType1,ArgType2,Def2,ArgType3,Def3) \
0204      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }              \
0205      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }       \
0206      template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> struct kCudaType<cxx::X<Arg1,Arg2,Arg3> >        \
0207      { using type_t = cuda::X<CudaType_t<Arg1>, Def2, Def3 >; }
0208   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_3v(classOrStruct,X,ArgType1,Def1,ArgType2,Def2,ArgType3,Def3) \
0209      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }                \
0210      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> classOrStruct X; }         \
0211      template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3> struct kCudaType<cxx::X<Arg1,Arg2,Arg3> >          \
0212      { using type_t = cuda::X<Def1,Def2,Def3 >; }
0213   #define VECGEOM_DEVICE_DECLARE_CONV_TEMPLATE_4v(classOrStruct,X,ArgType1,Def1,ArgType2,Def2,ArgType3,Def3,ArgType4,Def4) \
0214      namespace cuda { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3,ArgType4 Arg4> classOrStruct X; }                \
0215      inline namespace cxx  { template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3,ArgType4 Arg4> classOrStruct X; }         \
0216      template <ArgType1 Arg1,ArgType2 Arg2,ArgType3 Arg3,ArgType4 Arg4> struct kCudaType<cxx::X<Arg1,Arg2,Arg3,Arg4> >     \
0217      { using type_t = cuda::X<Def1,Def2,Def3,Def4 >; }
0218
0219 #endif // VECGEOM_CUDA_VOLUME_SPECIALIZATION
0220
0221 /* Instead of multiple macro, when we have auto expansion of Template pack we could use:
0222 template <typename... Arguments>
0223 struct kCudaType<cxx::BoxImplementation<Arguments...>  >
0224    { using type_t = typename cuda::BoxImplementation<CudaType_t<Arguments...> >; };
0225 */
0226 #endif
0227
0228 // clang-format on
0229
0230 namespace vecgeom {
0231 inline namespace VECGEOM_IMPL_NAMESPACE {
0232
0233 #ifndef VECCORE_CUDA
0234 using std::unique_ptr;
0235 #else
0236 template <typename T>
0237 class unique_ptr {
0238   T *fValue;
0239
0240 public:
0241   VECCORE_ATT_HOST_DEVICE
0242   unique_ptr(T *in) : fValue(in) {}
0243
0244   VECCORE_ATT_HOST_DEVICE
0245   ~unique_ptr() { delete fValue; }
0246
0247   VECCORE_ATT_HOST_DEVICE
0248   T *operator->() { return fValue; }
0249 };
0250
0251 template <typename T>
0252 class unique_ptr<T[]> {
0253   T *fValue;
0254
0255 public:
0256   VECCORE_ATT_HOST_DEVICE
0257   unique_ptr(T *in) : fValue(in) {}
0258
0259   VECCORE_ATT_HOST_DEVICE
0260   ~unique_ptr() { delete[] fValue; }
0261
0262   VECCORE_ATT_HOST_DEVICE
0263   T &operator[](size_t idx) { return fValue[idx]; }
0264 };
0265 #endif
0266 } // namespace VECGEOM_IMPL_NAMESPACE
0267 } // namespace vecgeom
0268
0269 #endif