Namespaces
namespace	CUDA_Helpers
namespace	Helpers

Classes
class	CellBaseInfo
	Contains the fundamental information that allows interacting with `CellInfoArr` so that it is easier to (pre-)transfer it. More...
class	CellInfoArr
	Contains the per-event cell information: energy, timing, gain, quality and provenance. More...
struct	CellNoiseArr
struct	CellNoiseProperties
class	ClusterBaseInfo
	Contains the fundamental information that allows interacting with `ClusterInfoArr` so that it is easier to (pre-)transfer it. More...
class	ClusterInfoArr
	`cellsPrefixSum`. More...
class	ClusterTag
	The class that actually expresses the cluster assignment. More...
class	ConstantDataHolder
	Holds CPU and GPU versions of the geometry and cell noise information, which are assumed to be constant throughout the run. More...
class	ConstantEnumConversion
class	EtaPhiMapEntry
	Holds an (eta, phi) to cell map for a given sampling. More...
struct	EtaPhiToCellMap
class	EventDataHolder
	Holds the mutable per-event information (clusters and cells) and provides utilities to convert between this representation and the Athena data structures (i. More...
class	GainConversion
	Provides utility functions to handle the gain conversion. More...
class	GenericTagBase
	A base class to implement the constructor, destructor and conversion operators. More...
struct	GeometryArr
class	MomentsOptionsArray
	Holds an array of bools to represent the different moments that may be calculated and transferred to/from the GPU. More...
class	multi_class_holder
	A convenient way to handle a compile-time list of types, useful for several metaprogramming techniques... More...
struct	NeighArr
class	NeighOffset
	A class that expresses the offset from the beginning of the neighbours list for the several neighbour options. More...
struct	NeighPairsArr
class	OtherCellInfo
	Packs the calo sampling, the intra-calorimeter sampling, the subcalo, the region and whether the cell should have its neighbours limited according to the PS and HEICW and FCal options. More...
class	QualityProvenance
	Just two uint16_t bit-packed onto a uint32_t. More...
class	Tag_1_12_1_32_18
	A tag with a bit flag, then 12 bits (for a counter), then another bit flag, then 32 bits (for a float) and finally 18 bits (for a cell index) More...
class	Tag_1_13_32_18
	A tag with a bit flag, then 13 bits (for a counter), then 32 bits (for a float) and finally 18 bits (for a cell index) More...
class	Tag_1_1_12_16_18_16
	A tag with two bit flags, then 12 bits (for a counter), 16 bits (for a half-precision float), 18 bits (for a cell index) and 16 bits (for a cluster index) More...
class	Tag_1_1_12_32_18
	A tag with two bit flags, then 12 bits (for a counter), 32 bits (for a float) and 18 bits (for a cell index) More...
class	Tag_1_1_7_31_8_16
	A tag with two bit flags, then 7 bits (for a counter), 31 bits (for a float with one bit shaved off), 8 bits (for an extra factor) and 16 bits (for a cluster index) More...
class	Tag_1_30_18_15
	A tag with a bit flag, then 30 bit (intended for a truncated float), 18 bit (for a cell index) and 15 bit (for a cluster index). More...
class	Tag_1_7_1_31_8_16
	A tag with a bit flag, then 7 bits (for a counter), then another bit flag, 31 bits (for a float with one bit shaved off), 8 bits (for an extra factor) and 16 bits (for a cluster index) More...
class	TagBase
	The base class for marking cells as belonging to clusters. More...

Typedefs
using	tag_type = TagBase::carrier

Enumerations
enum class	ClusterInformationState : uint8_t { None , Tags , TagsWithBasicInfo , Full , WithBasicInfo , WithMoments , WithExtraMoments }

Functions
template<class F, class... Types, class... Args>
void	apply_to_multi_class (F &&f, multi_class_holder< Types... >, Args &... args)
float	float_unhack (const unsigned int bits)
double	protect_from_zero (const double x)
float	protect_from_zero (const float x)
template<class T, unsigned int ... us, class ... PtrLikes>
__host__ __device__ T *	get_laundered_pointer (unsigned int idx, PtrLikes &&... p)
constexpr unsigned int	get_extra_alignment (const unsigned int base_align, const unsigned int required)
template<class T>
constexpr bool __host__ __device__	check_sufficient_size (unsigned int offset, unsigned int index)
template<class T, unsigned int offset, class PtrLike>
__host__ __device__ T *	get_laundered_pointer (unsigned int idx, PtrLike &&ptr)
template<class T, unsigned int offset, unsigned int ... us, class PtrLike, class ... PtrLikes>
__host__ __device__ T *	get_laundered_pointer (unsigned int idx, PtrLike &&ptr, PtrLikes &&... ps)
template<class T, class ... PtrLikes>
__host__ __device__ T *	get_laundered_pointer_striped (unsigned int jdx, unsigned int idx, PtrLikes &&... ps)
template<class T, class ... PtrLikes>
__host__ __device__ T *	get_laundered_pointer_stacked (unsigned int idx, PtrLikes &&... ps)
template<class T, class PtrLike>
__host__ __device__ T *	get_laundered_pointer_stacked (unsigned int idx, PtrLike &&ptr)
template<class T, class PtrLike, class ... PtrLikes>
__host__ __device__ T *	get_laundered_pointer_stacked (unsigned int idx, PtrLike &&ptr, PtrLikes &&... ps)

Variables
constexpr int	NCaloCells = 187652
constexpr int	LArCellStart = 0
constexpr int	TileCellStart = 182468
constexpr int	LArCellAfterEnd = TileCellStart
constexpr int	TileCellAfterEnd = NCaloCells
constexpr int	NLArCells = LArCellAfterEnd - LArCellStart
constexpr int	NTileCells = TileCellAfterEnd - TileCellStart
constexpr int	NMaxClusters = 0x10000U
constexpr int	NMaxPairs = 0x400000U
constexpr int	NExactPairs = 2560816
constexpr int	NumGainStates = 4
constexpr int	NumSamplings = 28
constexpr int	NumNeighOptions = 12
constexpr int	NMaxNeighbours = 34
constexpr int	NMaxAll2DNeighbours = 12

Detailed Description

Typedef Documentation

◆ tag_type

using CaloRecGPU::tag_type = TagBase::carrier

Definition at line 325 of file TagDefinitions.h.

Enumeration Type Documentation

◆ ClusterInformationState

enum class CaloRecGPU::ClusterInformationState : uint8_t

strong

Enumerator
None
Tags
TagsWithBasicInfo
Full
WithBasicInfo
WithMoments
WithExtraMoments

Definition at line 381 of file EventInfoDefinitions.h.

  {
    None,              //No cell assignment to clusters
    Tags,              //Cell assignment to clusters using the generic tags, unspecified order, some may be invalid
    TagsWithBasicInfo, //As before, but also with the basic information having been calculated
    Full,              //Proper cell assignment, all clusters are valid and sorted, but no info
    WithBasicInfo,     //As before, but also with the basic information having been calculated
    WithMoments,       //Moments have been calculated too
    WithExtraMoments,  //And also the additional moments beyond cluster moments calculation
  };

Function Documentation

◆ apply_to_multi_class()

template<class F, class... Types, class... Args>

void CaloRecGPU::apply_to_multi_class	(	F &&	f,
		multi_class_holder< Types... >	,
		Args &...	args )

Definition at line 24 of file CaloRecUtilities.h.

                                                                             {
  auto&& func = std::forward<F>(f);  // avoid forwarding F multiple times
  [&]<std::size_t... I>(std::index_sequence<I...>) {
    (func(Types{}, I, args...), ...); // avoid forwarding args multiple times
  }(std::make_index_sequence<sizeof...(Types)>{});
}

◆ check_sufficient_size()

template<class T>

bool __host__ __device__ CaloRecGPU::check_sufficient_size	(	unsigned int	offset,
		unsigned int	index )

constexpr

Definition at line 178 of file TemporaryHelpers.h.

  {
    return (NMaxClusters >= get_extra_alignment(offset, alignof(T)) + index * sizeof(T));
  }

◆ float_unhack()

float CaloRecGPU::float_unhack ( const unsigned int bits )

inline

Definition at line 31 of file CaloRecUtilities.h.

                                                    {
    float res;
    std::memcpy(&res, &bits, sizeof(float));
    //In C++20, we should bit-cast. For now, for our platform, works.
    return res;
  }

◆ get_extra_alignment()

unsigned int CaloRecGPU::get_extra_alignment	(	const unsigned int	base_align,
		const unsigned int	required )

inlineconstexpr

Definition at line 170 of file TemporaryHelpers.h.

  {
    const unsigned int delta = base_align % required;
 
    return required * (delta != 0) - delta;
  }

◆ get_laundered_pointer() [1/3]

template<class T, unsigned int offset, class PtrLike>

__host__ __device__ T * CaloRecGPU::get_laundered_pointer	(	unsigned int	idx,
		PtrLike &&	ptr )

Definition at line 184 of file TemporaryHelpers.h.

  {
    using PtrType = std::decay_t<decltype(ptr[0])>;
    using BasePtrType = std::conditional_t<std::is_const_v<PtrType>, const char *, char *>;
 
    constexpr unsigned int base_offset = offset % alignof(T);
 
    constexpr unsigned int extra_alignment = get_extra_alignment(base_offset, alignof(T));
 
    BasePtrType base_ptr = reinterpret_cast<BasePtrType>(&ptr[0]);
 
    return std::launder(reinterpret_cast<T *>(base_ptr + extra_alignment + idx * sizeof(T)));
  }

◆ get_laundered_pointer() [2/3]

template<class T, unsigned int offset, unsigned int ... us, class PtrLike, class ... PtrLikes>

__host__ __device__ T * CaloRecGPU::get_laundered_pointer	(	unsigned int	idx,
		PtrLike &&	ptr,
		PtrLikes &&...	ps )

Definition at line 199 of file TemporaryHelpers.h.

  {
    using PtrType = std::decay_t<decltype(ptr[0])>;
 
    constexpr unsigned int max_size = NMaxClusters * sizeof(PtrType);
 
    constexpr unsigned int base_offset = offset % alignof(T);
    constexpr unsigned int extra_alignment = get_extra_alignment(base_offset, alignof(T));
 
    constexpr unsigned int real_size = (max_size - extra_alignment) / sizeof(T);
 
    return (real_size > idx ? get_laundered_pointer<T, offset>(idx, std::forward<PtrLike>(ptr)) : get_laundered_pointer<T, us...>(idx - real_size, std::forward<PtrLikes>(ps)...));
  }

◆ get_laundered_pointer() [3/3]

template<class T, unsigned int ... us, class ... PtrLikes>

__host__ __device__ T * CaloRecGPU::get_laundered_pointer	(	unsigned int	idx,
		PtrLikes &&...	p )

◆ get_laundered_pointer_stacked() [1/3]

template<class T, class PtrLike>

__host__ __device__ T * CaloRecGPU::get_laundered_pointer_stacked	(	unsigned int	idx,
		PtrLike &&	ptr )

Definition at line 236 of file TemporaryHelpers.h.

  {
    static_assert(alignof(T) <= alignof(double), "We don't support aligning in this case...");
 
    using PtrType = std::decay_t<decltype(ptr[0][0])>;
    using BasePtrType = std::conditional_t<std::is_const_v<PtrType>, const char *, char *>;
 
    constexpr unsigned int num_per_array = (NMaxClusters * sizeof(PtrType)) / sizeof(T);
 
    const unsigned int first_idx = idx / num_per_array;
 
    const unsigned int second_idx = idx % num_per_array;
 
    return get_laundered_pointer<T, 0>(second_idx, ptr[first_idx]);
  }

◆ get_laundered_pointer_stacked() [2/3]

template<class T, class PtrLike, class ... PtrLikes>

__host__ __device__ T * CaloRecGPU::get_laundered_pointer_stacked	(	unsigned int	idx,
		PtrLike &&	ptr,
		PtrLikes &&...	ps )

Definition at line 253 of file TemporaryHelpers.h.

  {
    static_assert(alignof(T) <= alignof(double));
 
    using PtrType = std::decay_t<decltype(ptr[0][0])>;
 
    constexpr unsigned int num_per_array = (NMaxClusters * sizeof(PtrType)) / sizeof(T);
    constexpr unsigned int total_num = num_per_array * NumSamplings;
 
    return (total_num > idx ? get_laundered_pointer_stacked<T>(idx, std::forward<PtrLike>(ptr)) : get_laundered_pointer_stacked<T>(idx - total_num, std::forward<PtrLikes>(ps)...));
  }

◆ get_laundered_pointer_stacked() [3/3]

template<class T, class ... PtrLikes>

__host__ __device__ T * CaloRecGPU::get_laundered_pointer_stacked	(	unsigned int	idx,
		PtrLikes &&...	ps )

◆ get_laundered_pointer_striped()

template<class T, class ... PtrLikes>

__host__ __device__ T * CaloRecGPU::get_laundered_pointer_striped	(	unsigned int	jdx,
		unsigned int	idx,
		PtrLikes &&...	ps )

Definition at line 219 of file TemporaryHelpers.h.

  {
    static_assert(alignof(T) <= alignof(double), "We don't support aligning in this case...");
 
    return get_laundered_pointer<T, (alignof(decltype(ps[0][0])) * 0)...>(idx, ps[jdx]...);
    //We pass a param pack of 0s (with the right size)
    //as we know we won't need alignment.
  }

◆ protect_from_zero() [1/2]

double CaloRecGPU::protect_from_zero ( const double x )

inline

Definition at line 38 of file CaloRecUtilities.h.

                                                 {
    return x == 0 ? 1e-15 : x;
  }

◆ protect_from_zero() [2/2]

float CaloRecGPU::protect_from_zero ( const float x )

inline

Definition at line 42 of file CaloRecUtilities.h.

                                               {
    return x == 0 ? 1e-7 : x;
  }

Variable Documentation

◆ LArCellAfterEnd

int CaloRecGPU::LArCellAfterEnd = TileCellStart

inlineconstexpr

Definition at line 17 of file BaseDefinitions.h.

◆ LArCellStart

int CaloRecGPU::LArCellStart = 0

inlineconstexpr

Definition at line 15 of file BaseDefinitions.h.

◆ NCaloCells

int CaloRecGPU::NCaloCells = 187652

inlineconstexpr

Definition at line 12 of file BaseDefinitions.h.

◆ NExactPairs

int CaloRecGPU::NExactPairs = 2560816

inlineconstexpr

Definition at line 36 of file BaseDefinitions.h.

◆ NLArCells

int CaloRecGPU::NLArCells = LArCellAfterEnd - LArCellStart

inlineconstexpr

Definition at line 19 of file BaseDefinitions.h.

◆ NMaxAll2DNeighbours

int CaloRecGPU::NMaxAll2DNeighbours = 12

inlineconstexpr

Definition at line 63 of file BaseDefinitions.h.

◆ NMaxClusters

int CaloRecGPU::NMaxClusters = 0x10000U

inlineconstexpr

Definition at line 27 of file BaseDefinitions.h.

◆ NMaxNeighbours

int CaloRecGPU::NMaxNeighbours = 34

inlineconstexpr

Definition at line 62 of file BaseDefinitions.h.

◆ NMaxPairs

int CaloRecGPU::NMaxPairs = 0x400000U

inlineconstexpr

Definition at line 35 of file BaseDefinitions.h.

◆ NTileCells

int CaloRecGPU::NTileCells = TileCellAfterEnd - TileCellStart

inlineconstexpr

Definition at line 20 of file BaseDefinitions.h.

◆ NumGainStates

int CaloRecGPU::NumGainStates = 4

inlineconstexpr

Definition at line 39 of file BaseDefinitions.h.

◆ NumNeighOptions

int CaloRecGPU::NumNeighOptions = 12

inlineconstexpr

Definition at line 46 of file BaseDefinitions.h.

◆ NumSamplings

int CaloRecGPU::NumSamplings = 28

inlineconstexpr

Definition at line 43 of file BaseDefinitions.h.

◆ TileCellAfterEnd

int CaloRecGPU::TileCellAfterEnd = NCaloCells

inlineconstexpr

Definition at line 18 of file BaseDefinitions.h.

◆ TileCellStart

int CaloRecGPU::TileCellStart = 182468

inlineconstexpr

Definition at line 16 of file BaseDefinitions.h.

Namespaces

Classes

Typedefs

Enumerations

Functions

Variables

Detailed Description

Typedef Documentation

◆ tag_type

Enumeration Type Documentation

◆ ClusterInformationState

Function Documentation

◆ apply_to_multi_class()

◆ check_sufficient_size()

◆ float_unhack()

◆ get_extra_alignment()

◆ get_laundered_pointer() [1/3]

◆ get_laundered_pointer() [2/3]

◆ get_laundered_pointer() [3/3]

◆ get_laundered_pointer_stacked() [1/3]

◆ get_laundered_pointer_stacked() [2/3]

◆ get_laundered_pointer_stacked() [3/3]

◆ get_laundered_pointer_striped()

◆ protect_from_zero() [1/2]

◆ protect_from_zero() [2/2]

Variable Documentation

◆ LArCellAfterEnd

◆ LArCellStart

◆ NCaloCells

◆ NExactPairs

◆ NLArCells

◆ NMaxAll2DNeighbours

◆ NMaxClusters

◆ NMaxNeighbours

◆ NMaxPairs

◆ NTileCells

◆ NumGainStates

◆ NumNeighOptions

◆ NumSamplings

◆ TileCellAfterEnd

◆ TileCellStart