Backend.h

/// \file cuda/Backend.h
/// \author Johannes de Fine Licht (johannes.definelicht@cern.ch)

#ifndef VECGEOM_BACKEND_CUDABACKEND_H_
#define VECGEOM_BACKEND_CUDABACKEND_H_

#include "base/Global.h"

#include "backend/scalar/Backend.h"
#include "backend/cuda/Interface.h"

namespace vecgeom {
#ifdef VECCORE_CUDA
inline
#endif
    namespace cuda {

struct kCuda {
  typedef int int_v;
  typedef Precision precision_v;
  typedef bool bool_v;
  typedef Inside_t inside_v;
  static constexpr precision_v kOne  = 1.0;
  static constexpr precision_v kZero = 0.0;
  const static bool_v kTrue          = true;
  const static bool_v kFalse         = false;
  // alternative typedefs ( might supercede above typedefs )
  typedef int Int_t;
  typedef Precision Double_t;
  typedef bool Bool_t;
  typedef int Index_t;
};

typedef kCuda::int_v CudaInt;
typedef kCuda::precision_v CudaPrecision;
typedef kCuda::bool_v CudaBool;

#if defined(VECGEOM_ENABLE_CUDA) && !defined(VECGEOM_BACKEND_TYPE)
constexpr size_t kVectorSize = 1;
#define VECGEOM_BACKEND_TYPE vecgeom::kScalar
#define VECGEOM_BACKEND_PRECISION_FROM_PTR(P) (*(P))
#define VECGEOM_BACKEND_PRECISION_TYPE Precision
#define VECGEOM_BACKEND_PRECISION_TYPE_SIZE 1
//#define VECGEOM_BACKEND_PRECISION_NOT_SCALAR
#define VECGEOM_BACKEND_BOOL vecgeom::ScalarBool
#define VECGEOM_BACKEND_INSIDE vecgeom::kScalar::inside_v
#endif

static const unsigned kThreadsPerBlock = 256;

// Auxiliary GPU functions
#ifdef VECCORE_CUDA

VECCORE_ATT_DEVICE
VECGEOM_FORCE_INLINE
int ThreadIndex()
{
  return blockDim.x * blockIdx.x + threadIdx.x;
}

VECCORE_ATT_DEVICE
VECGEOM_FORCE_INLINE
int ThreadOffset()
{
  return blockDim.x * gridDim.x;
}

#endif

/**