ATLAS Offline Software
Namespaces | Macros
GPUClusterInfoAndMomentsCalculatorImplHelper.h File Reference
#include "CaloRecGPU/Helpers.h"
#include "CaloRecGPU/CUDAFriendlyClasses.h"
#include "GPUClusterInfoAndMomentsCalculatorImpl.h"
#include "FPHelpers.h"
#include "CaloGeoHelpers/CaloSampling.h"
#include <cmath>
#include <type_traits>
Include dependency graph for GPUClusterInfoAndMomentsCalculatorImplHelper.h:

Go to the source code of this file.

Namespaces

 CMCTemporaries
 
 ToLoad
 
 ToCalculate
 

Macros

#define CALORECGPU_EXPAND(...)   __VA_ARGS__
 
#define CALORECGPU_CONCAT_HELPER_INNER(A, ...)   A ## __VA_ARGS__
 
#define CALORECGPU_CONCAT_HELPER(A, B)   CALORECGPU_CONCAT_HELPER_INNER(A, B)
 
#define CMC_TEMPARR_1_DECLARE(TEMPNAME, TYPE)
 
#define CMC_TEMPARR_1_SPLIT_DECLARE(TEMPNAME, TYPE)
 
#define CMC_TEMPARR_2_DECLARE(TEMPNAME, TYPE)
 
#define CMC_TEMPARR_2_SPLIT_DECLARE(TEMPNAME, TYPE)
 
#define CMC_TEMPARR_1(TEMPNAME, BASEVAR, TYPE)
 
#define CMC_TEMPARR_1_SPLIT(TEMPNAME, BASEVAR1, BASEVAR2, TYPE)
 
#define CMC_TEMPARR_2(TEMPNAME, BASEVAR, TYPE)
 
#define CMC_TEMPARR_2_SPLIT(TEMPNAME, BASEVAR1, BASEVAR2, TYPE)
 
#define CALORECGPU_CMC_LOAD(NAME, NEEDED, PREVNEEDED, VARS, INIT)
 
#define CALORECGPU_CMC_LOAD_SIMPLE_CELL_INFO(NAME, VARNAME, PROPNAME)
 
#define CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO(NAME, VARNAME, PROPNAME)
 
#define CALORECGPU_CMC_LOAD_SIMPLE_CLUSTER_INFO(NAME, VARNAME, PROPNAME)
 
#define CALORECGPU_CMC_LOAD_SIMPLE_MOMENT_INFO(NAME, VARNAME, PROPNAME)
 
#define CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_MOMENT_INFO(NAME, VARNAME, PROPNAME)
 
#define CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO(NAME, VARNAME, PROPNAME)
 
#define CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_TEMPORARY_INFO(NAME, VARNAME, PROPNAME)
 
#define CALORECGPU_CMC_MOMENT_CALC(NAME, BEFORELOAD, BEFOREEXEC, CELLLOAD, CLUSTERLOAD, CELLEXEC, AFTERLOAD, AFTEREXEC)
 

Macro Definition Documentation

◆ CALORECGPU_CMC_LOAD

#define CALORECGPU_CMC_LOAD (   NAME,
  NEEDED,
  PREVNEEDED,
  VARS,
  INIT 
)
Value:
struct NAME \
{ \
using AssumedList = TypeList<CALORECGPU_EXPAND NEEDED >; \
using AssumedPreviousList = TypeList<CALORECGPU_EXPAND PREVNEEDED >; \
CALORECGPU_EXPAND VARS \
template <class Final> __device__ NAME(const Final & f, Parameters p, const int idx) { CALORECGPU_EXPAND INIT } \
}

Definition at line 834 of file GPUClusterInfoAndMomentsCalculatorImplHelper.h.

◆ CALORECGPU_CMC_LOAD_SIMPLE_CELL_INFO

#define CALORECGPU_CMC_LOAD_SIMPLE_CELL_INFO (   NAME,
  VARNAME,
  PROPNAME 
)
Value:
CALORECGPU_CMC_LOAD(NAME, \
(), \
(), \
(std::decay_t<decltype(std::declval<CaloRecGPU::CellInfoArr>().PROPNAME[0])> VARNAME;), \
(VARNAME = p.cell_info_arr->PROPNAME[idx];) \
);

Definition at line 845 of file GPUClusterInfoAndMomentsCalculatorImplHelper.h.

◆ CALORECGPU_CMC_LOAD_SIMPLE_CLUSTER_INFO

#define CALORECGPU_CMC_LOAD_SIMPLE_CLUSTER_INFO (   NAME,
  VARNAME,
  PROPNAME 
)
Value:
CALORECGPU_CMC_LOAD(NAME, \
(), \
(), \
(std::decay_t<decltype(std::declval<CaloRecGPU::ClusterInfoArr>().PROPNAME[0])> VARNAME;), \
(VARNAME = p.clusters_arr->PROPNAME[idx];) \
);

Definition at line 861 of file GPUClusterInfoAndMomentsCalculatorImplHelper.h.

◆ CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO

#define CALORECGPU_CMC_LOAD_SIMPLE_GEOMETRY_INFO (   NAME,
  VARNAME,
  PROPNAME 
)
Value:
CALORECGPU_CMC_LOAD(NAME, \
(), \
(), \
(std::decay_t<decltype(std::declval<CaloRecGPU::GeometryArr>().PROPNAME[0])> VARNAME;), \
(VARNAME = p.geometry->PROPNAME[idx];) \
);

Definition at line 853 of file GPUClusterInfoAndMomentsCalculatorImplHelper.h.

◆ CALORECGPU_CMC_LOAD_SIMPLE_MOMENT_INFO

#define CALORECGPU_CMC_LOAD_SIMPLE_MOMENT_INFO (   NAME,
  VARNAME,
  PROPNAME 
)
Value:
CALORECGPU_CMC_LOAD(NAME, \
(), \
(), \
(std::decay_t<decltype(std::declval<CaloRecGPU::ClusterMomentsArr>().PROPNAME[0])> VARNAME;), \
(VARNAME = p.moments_arr->PROPNAME[idx];) \
);

Definition at line 869 of file GPUClusterInfoAndMomentsCalculatorImplHelper.h.

◆ CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_MOMENT_INFO

#define CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_MOMENT_INFO (   NAME,
  VARNAME,
  PROPNAME 
)
Value:
CALORECGPU_CMC_LOAD(NAME, \
(), \
(), \
(std::decay_t<decltype(std::declval<CaloRecGPU::ClusterMomentsArr>().PROPNAME[0][0])> VARNAME;), \
(VARNAME = p.moments_arr->PROPNAME[f.sampling][idx];) \
);

Definition at line 879 of file GPUClusterInfoAndMomentsCalculatorImplHelper.h.

◆ CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_TEMPORARY_INFO

#define CALORECGPU_CMC_LOAD_SIMPLE_PER_SAMPLING_TEMPORARY_INFO (   NAME,
  VARNAME,
  PROPNAME 
)
Value:
CALORECGPU_CMC_LOAD(NAME, \
(), \
(), \
(std::decay_t<decltype(CMCTemporaries::PROPNAME(std::declval<CaloRecGPU::ClusterMomentsArr*>(),0,0))> VARNAME;), \
(VARNAME = CMCTemporaries::PROPNAME(p.moments_arr, f.sampling, idx);) \
);

Definition at line 897 of file GPUClusterInfoAndMomentsCalculatorImplHelper.h.

◆ CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO

#define CALORECGPU_CMC_LOAD_SIMPLE_TEMPORARY_INFO (   NAME,
  VARNAME,
  PROPNAME 
)
Value:
CALORECGPU_CMC_LOAD(NAME, \
(), \
(), \
(std::decay_t<decltype(CMCTemporaries::PROPNAME(std::declval<CaloRecGPU::ClusterMomentsArr *>(),0))> VARNAME;), \
(VARNAME = CMCTemporaries::PROPNAME(p.moments_arr, idx);) \
);

Definition at line 887 of file GPUClusterInfoAndMomentsCalculatorImplHelper.h.

◆ CALORECGPU_CMC_MOMENT_CALC

#define CALORECGPU_CMC_MOMENT_CALC (   NAME,
  BEFORELOAD,
  BEFOREEXEC,
  CELLLOAD,
  CLUSTERLOAD,
  CELLEXEC,
  AFTERLOAD,
  AFTEREXEC 
)
Value:
struct NAME \
{ \
using BeforeLoading = TypeList<CALORECGPU_EXPAND BEFORELOAD>; \
template <class T> __device__ static void before(Parameters p, \
const T & data, \
const int cluster) \
{ CALORECGPU_EXPAND BEFOREEXEC } \
using CellLoading = TypeList<CALORECGPU_EXPAND CELLLOAD>; \
using ClusterLoading = TypeList<CALORECGPU_EXPAND CLUSTERLOAD>; \
template <class T> __device__ static void per_cell(Parameters p, \
const T & data, \
const int cell, \
const int cluster) \
{ CALORECGPU_EXPAND CELLEXEC } \
using AfterLoading = TypeList<CALORECGPU_EXPAND AFTERLOAD>; \
template <class T> __device__ static void after(Parameters p, \
const T & data, \
const int cluster) \
{ CALORECGPU_EXPAND AFTEREXEC } \
}

Definition at line 1357 of file GPUClusterInfoAndMomentsCalculatorImplHelper.h.

◆ CALORECGPU_CONCAT_HELPER

#define CALORECGPU_CONCAT_HELPER (   A,
  B 
)    CALORECGPU_CONCAT_HELPER_INNER(A, B)

◆ CALORECGPU_CONCAT_HELPER_INNER

#define CALORECGPU_CONCAT_HELPER_INNER (   A,
  ... 
)    A ## __VA_ARGS__

◆ CALORECGPU_EXPAND

#define CALORECGPU_EXPAND (   ...)    __VA_ARGS__

◆ CMC_TEMPARR_1

#define CMC_TEMPARR_1 (   TEMPNAME,
  BASEVAR,
  TYPE 
)
Value:
template <class PtrLike> __host__ __device__ const TYPE * TEMPNAME (const PtrLike & arr) \
{ \
constexpr auto misalignment = offsetof(CaloRecGPU::ClusterMomentsArr, BASEVAR) % alignof(TYPE); \
const char * aligned = ((const char *) arr->BASEVAR) + (alignof(TYPE) - misalignment) * (misalignment > 0); \
return (const TYPE *) (aligned); \
} \
template <class PtrLike> __host__ __device__ TYPE * TEMPNAME (PtrLike & arr) \
{ \
constexpr auto misalignment = offsetof(CaloRecGPU::ClusterMomentsArr, BASEVAR) % alignof(TYPE); \
char * aligned = ((char *) arr->BASEVAR) + (alignof(TYPE) - misalignment) * (misalignment > 0); \
return (TYPE *) (aligned); \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const int idx) \
{ \
return TEMPNAME (arr) [idx]; \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const int idx) \
{ \
return TEMPNAME (arr) [idx]; \
} struct to_end_with_semicolon

◆ CMC_TEMPARR_1_DECLARE

#define CMC_TEMPARR_1_DECLARE (   TEMPNAME,
  TYPE 
)
Value:
template <class PtrLike> __host__ __device__ const TYPE * TEMPNAME (const PtrLike & arr); \
template <class PtrLike> __host__ __device__ TYPE * TEMPNAME (PtrLike & arr); \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const int idx); \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const int idx); \
struct to_end_with_semicolon

◆ CMC_TEMPARR_1_SPLIT

#define CMC_TEMPARR_1_SPLIT (   TEMPNAME,
  BASEVAR1,
  BASEVAR2,
  TYPE 
)
Value:
CMC_TEMPARR_1(CALORECGPU_CONCAT_HELPER(TEMPNAME, _1), BASEVAR1, TYPE); \
CMC_TEMPARR_1(CALORECGPU_CONCAT_HELPER(TEMPNAME, _2), BASEVAR2, TYPE); \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const int idx) \
{ \
if (idx >= split_size) \
{ \
return CALORECGPU_CONCAT_HELPER(TEMPNAME, _2)(arr, idx - split_size); \
} \
else \
{ \
return CALORECGPU_CONCAT_HELPER(TEMPNAME, _1)(arr, idx); \
} \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const int idx) \
{ \
if (idx >= split_size) \
{ \
return CALORECGPU_CONCAT_HELPER(TEMPNAME, _2)(arr, idx - split_size); \
} \
else \
{ \
return CALORECGPU_CONCAT_HELPER(TEMPNAME, _1)(arr, idx); \
} \
} struct to_end_with_semicolon

◆ CMC_TEMPARR_1_SPLIT_DECLARE

#define CMC_TEMPARR_1_SPLIT_DECLARE (   TEMPNAME,
  TYPE 
)
Value:
CMC_TEMPARR_1_DECLARE(CALORECGPU_CONCAT_HELPER(TEMPNAME, _1), TYPE); \
CMC_TEMPARR_1_DECLARE(CALORECGPU_CONCAT_HELPER(TEMPNAME, _2), TYPE); \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const int idx); \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const int idx); \
struct to_end_with_semicolon

◆ CMC_TEMPARR_2

#define CMC_TEMPARR_2 (   TEMPNAME,
  BASEVAR,
  TYPE 
)
Value:
template <class PtrLike> __host__ __device__ const nested_type<TYPE> * TEMPNAME (const PtrLike & arr) \
{ \
constexpr auto misalignment = offsetof(CaloRecGPU::ClusterMomentsArr, BASEVAR) % alignof(nested_type<TYPE>); \
const char * aligned = ((const char *) arr->BASEVAR) + (alignof(nested_type<TYPE>) - misalignment) * (misalignment > 0); \
return (const nested_type<TYPE> *) (aligned); \
} \
template <class PtrLike> __host__ __device__ nested_type<TYPE> * TEMPNAME (PtrLike & arr) \
{ \
constexpr auto misalignment = offsetof(CaloRecGPU::ClusterMomentsArr, BASEVAR) % alignof(nested_type<TYPE>); \
char * aligned = ((char *) arr->BASEVAR) + (alignof(nested_type<TYPE>) - misalignment) * (misalignment > 0); \
return (nested_type<TYPE> *) (aligned); \
} \
template <class PtrLike> __host__ __device__ const nested_type<TYPE> & TEMPNAME (const PtrLike & arr, const int idx) \
{ \
return TEMPNAME (arr) [idx]; \
} \
template <class PtrLike> __host__ __device__ nested_type<TYPE> & TEMPNAME (PtrLike & arr, const int idx) \
{ \
return TEMPNAME (arr) [idx]; \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const int idx, const int jdx) \
{ \
return TEMPNAME (arr) [idx] [jdx]; \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const int idx, const int jdx) \
{ \
return TEMPNAME (arr) [idx] [jdx]; \
} struct to_end_with_semicolon

◆ CMC_TEMPARR_2_DECLARE

#define CMC_TEMPARR_2_DECLARE (   TEMPNAME,
  TYPE 
)
Value:
template <class PtrLike> __host__ __device__ const nested_type<TYPE> * TEMPNAME (const PtrLike & arr); \
template <class PtrLike> __host__ __device__ nested_type<TYPE> * TEMPNAME (PtrLike & arr); \
template <class PtrLike> __host__ __device__ const nested_type<TYPE> & TEMPNAME (const PtrLike & arr, const int idx); \
template <class PtrLike> __host__ __device__ nested_type<TYPE> & TEMPNAME (PtrLike & arr, const int idx); \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const int idx, const int jdx); \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const int idx, const int jdx); \
struct to_end_with_semicolon

◆ CMC_TEMPARR_2_SPLIT

#define CMC_TEMPARR_2_SPLIT (   TEMPNAME,
  BASEVAR1,
  BASEVAR2,
  TYPE 
)
Value:
CMC_TEMPARR_2(CALORECGPU_CONCAT_HELPER(TEMPNAME, _1), BASEVAR1, TYPE); \
CMC_TEMPARR_2(CALORECGPU_CONCAT_HELPER(TEMPNAME, _2), BASEVAR2, TYPE); \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const int idx, const int jdx) \
{ \
if (jdx >= split_size) \
{ \
return CALORECGPU_CONCAT_HELPER(TEMPNAME, _2)(arr, idx, jdx - split_size); \
} \
else \
{ \
return CALORECGPU_CONCAT_HELPER(TEMPNAME, _1)(arr, idx, jdx); \
} \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const int idx, const int jdx) \
{ \
if (jdx >= split_size) \
{ \
return CALORECGPU_CONCAT_HELPER(TEMPNAME, _2)(arr, idx, jdx - split_size); \
} \
else \
{ \
return CALORECGPU_CONCAT_HELPER(TEMPNAME, _1)(arr, idx, jdx); \
} \
} struct to_end_with_semicolon

◆ CMC_TEMPARR_2_SPLIT_DECLARE

#define CMC_TEMPARR_2_SPLIT_DECLARE (   TEMPNAME,
  TYPE 
)
Value:
CMC_TEMPARR_2_DECLARE(CALORECGPU_CONCAT_HELPER(TEMPNAME, _1), TYPE); \
CMC_TEMPARR_2_DECLARE(CALORECGPU_CONCAT_HELPER(TEMPNAME, _2), TYPE); \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const int idx, const int jdx); \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const int idx, const int jdx); \
struct to_end_with_semicolon
data
char data[hepevt_bytes_allocation_ATLAS]
Definition: HepEvt.cxx:11
INIT
#define INIT(__TYPE)
ReadCellNoiseFromCool.cell
cell
Definition: ReadCellNoiseFromCool.py:53
MuonR4::SegmentFit::Parameters
AmgVector(toInt(ParamDefs::nPars)) Parameters
Definition: MuonHoughDefs.h:48
CALORECGPU_CONCAT_HELPER
#define CALORECGPU_CONCAT_HELPER(A, B)
CMC_TEMPARR_2
#define CMC_TEMPARR_2(TEMPNAME, BASEVAR, TYPE)
CALORECGPU_CMC_LOAD
#define CALORECGPU_CMC_LOAD(NAME, NEEDED, PREVNEEDED, VARS, INIT)
Definition: GPUClusterInfoAndMomentsCalculatorImplHelper.h:833
CMC_TEMPARR_1
#define CMC_TEMPARR_1(TEMPNAME, BASEVAR, TYPE)
python.utils.AtlRunQueryDQUtils.p
p
Definition: AtlRunQueryDQUtils.py:209
hist_file_dump.f
f
Definition: hist_file_dump.py:140
CMC_TEMPARR_2_DECLARE
#define CMC_TEMPARR_2_DECLARE(TEMPNAME, TYPE)
TYPE
#define TYPE(CODE, TYP, IOTYP)
CALORECGPU_EXPAND
#define CALORECGPU_EXPAND(...)
CaloRecGPU::ClusterMomentsArr
Definition: EventInfoDefinitions.h:341
LArNewCalib_DelayDump_OFC_Cali.idx
idx
Definition: LArNewCalib_DelayDump_OFC_Cali.py:69
CMC_TEMPARR_1_DECLARE
#define CMC_TEMPARR_1_DECLARE(TEMPNAME, TYPE)
TSU::T
unsigned long long T
Definition: L1TopoDataTypes.h:35
Generated on Mon Sep 1 2025 21:25:20 for ATLAS Offline Software by   doxygen 1.8.18