ATLAS Offline Software
Loading...
Searching...
No Matches
Namespaces | Macros | Functions
TemporaryHelpers.h File Reference
#include "CaloRecGPU/BaseDefinitions.h"
#include <cassert>
#include <new>
#include <utility>
#include <type_traits>
Include dependency graph for TemporaryHelpers.h:
This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Namespaces

namespace  CaloRecGPU
 Copyright (C) 2002-2025 CERN for the benefit of the ATLAS collaboration.

Macros

#define CALORECGPU_TEMP_CONCAT_HELPER_INNER(A, ...)
#define CALORECGPU_TEMP_CONCAT_HELPER(A, B)
#define CALORECGPU_TEMPARR_1(TEMPNAME, BASEVAR, TYPE)
 Defines a temporary variable of type TYPE with name TEMPNAME that will be stored in the 1D array accessible through BASEVAR.
#define CALORECGPU_TEMPARR_2(TEMPNAME, BASEVAR1, BASEVAR2, TYPE)
 Defines a temporary variable of type TYPE with name TEMPNAME that will be split across two 1D arrays accessible through BASEVAR1 and BASEVAR2.
#define CALORECGPU_TEMP2DARR_1(TEMPNAME, BASEVAR, TYPE)
 Defines a temporary variable of type TYPE with name TEMPNAME that will be stored in the 2D array (i.
#define CALORECGPU_TEMP2DARR_2(TEMPNAME, BASEVAR1, BASEVAR2, TYPE)
 Defines a temporary variable of type TYPE with name TEMPNAME that will will be split across two 2D arrays (i.
#define CALORECGPU_TEMPBIGARR_1(TEMPNAME, BASEVAR, TYPE)
 Defines a temporary variable of type TYPE with name TEMPNAME that will be stored as a 1D array split among the entries of the 2D array (i.
#define CALORECGPU_TEMPBIGARR_2(TEMPNAME, BASEVAR1, BASEVAR2, TYPE)
 Defines a temporary variable of type TYPE with name TEMPNAME that will be stored as a 1D array split among the entries of two 2D arrays (i.
#define CALORECGPU_TEMPBIGARR_3(TEMPNAME, BASEVAR1, BASEVAR2, BASEVAR3, TYPE)
 Defines a temporary variable of type TYPE with name TEMPNAME that will be stored as a 1D array split among the entries of three 2D arrays (i.
#define CALORECGPU_TEMPCELLARR_1(TEMPNAME, BASEVAR1, BASEVAR2, BASEVAR3, TYPE)
 Defines a temporary variable of type TYPE with name TEMPNAME that can have up to three times NMaxClusters, split across three variables.
#define CALORECGPU_TEMPCELLARR_2(TEMPNAME, BASEVAR1, BASEVAR2, BASEVAR3, BASEVAR4, BASEVAR5, BASEVAR6, TYPE)
 The definitions for a temporary variable of type TYPE with name TEMPNAME that can have up to three times NMaxClusters, split across six variables.
#define CALORECGPU_TEMPVAR(TEMPNAME, BASEVAR, INDEX, TYPE)
 Defines a temporary variable of type TYPE with name TEMPNAME that will be stored in the INDEX entry of an 1D array.
#define CALORECGPU_TEMPWRAPPER(TEMPNAME, WRAPPED)
 Wraps a function in another name, for better semantics while reusing storage...

Functions

template<class T, unsigned int ... us, class ... PtrLikes>
__host__ __device__ T * CaloRecGPU::get_laundered_pointer (unsigned int idx, PtrLikes &&... p)
constexpr unsigned int CaloRecGPU::get_extra_alignment (const unsigned int base_align, const unsigned int required)
template<class T>
constexpr bool __host__ __device__ CaloRecGPU::check_sufficient_size (unsigned int offset, unsigned int index)
template<class T, unsigned int offset, class PtrLike>
__host__ __device__ T * CaloRecGPU::get_laundered_pointer (unsigned int idx, PtrLike &&ptr)
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)
template<class T, class ... PtrLikes>
__host__ __device__ T * CaloRecGPU::get_laundered_pointer_striped (unsigned int jdx, unsigned int idx, PtrLikes &&... ps)
template<class T, class ... PtrLikes>
__host__ __device__ T * CaloRecGPU::get_laundered_pointer_stacked (unsigned int idx, PtrLikes &&... ps)
template<class T, class PtrLike>
__host__ __device__ T * CaloRecGPU::get_laundered_pointer_stacked (unsigned int idx, PtrLike &&ptr)
template<class T, class PtrLike, class ... PtrLikes>
__host__ __device__ T * CaloRecGPU::get_laundered_pointer_stacked (unsigned int idx, PtrLike &&ptr, PtrLikes &&... ps)

Macro Definition Documentation

◆ CALORECGPU_TEMP2DARR_1

#define CALORECGPU_TEMP2DARR_1 ( TEMPNAME,
BASEVAR,
TYPE )
Value:
_Pragma("nv_diag_suppress 177") \
template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer_striped<const TYPE>(jdx, idx, arr->moments. BASEVAR); \
} \
template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer_striped<TYPE>(jdx, idx, arr->moments. BASEVAR); \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, jdx, idx); \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, jdx, idx); \
} _Pragma("nv_diag_default 177") struct to_end_with_semicolon
CALORECGPU_TEMP_CONCAT_HELPER
#define CALORECGPU_TEMP_CONCAT_HELPER(A, B)
Definition TemporaryHelpers.h:11
TYPE
#define TYPE(CODE, TYP, IOTYP)
CaloRecGPU::get_laundered_pointer_striped
__host__ __device__ T * get_laundered_pointer_striped(unsigned int jdx, unsigned int idx, PtrLikes &&... ps)
Definition TemporaryHelpers.h:219
LArNewCalib_DelayDump_OFC_Cali.idx
idx
Definition LArNewCalib_DelayDump_OFC_Cali.py:69

Defines a temporary variable of type TYPE with name TEMPNAME that will be stored in the 2D array (i.

e. per sampling) accessible through BASEVAR.

Remarks
This would be a 32-bit per sampling temporary variable per cluster.

Definition at line 319 of file TemporaryHelpers.h.

319#define CALORECGPU_TEMP2DARR_1(TEMPNAME, BASEVAR, TYPE) _Pragma("nv_diag_suppress 177") \
320 template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
321 { \
322 return CaloRecGPU::get_laundered_pointer_striped<const TYPE>(jdx, idx, arr->moments. BASEVAR); \
323 } \
324 template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
325 { \
326 return CaloRecGPU::get_laundered_pointer_striped<TYPE>(jdx, idx, arr->moments. BASEVAR); \
327 } \
328 template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
329 { \
330 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, jdx, idx); \
331 } \
332 template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
333 { \
334 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, jdx, idx); \
335 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon

◆ CALORECGPU_TEMP2DARR_2

#define CALORECGPU_TEMP2DARR_2 ( TEMPNAME,
BASEVAR1,
BASEVAR2,
TYPE )
Value:
_Pragma("nv_diag_suppress 177") \
template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer_striped<const TYPE>(jdx, idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
} \
template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer_striped<TYPE>(jdx, idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, jdx, idx); \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, jdx, idx); \
} _Pragma("nv_diag_default 177") struct to_end_with_semicolon

Defines a temporary variable of type TYPE with name TEMPNAME that will will be split across two 2D arrays (i.

e. per sampling) accessible through BASEVAR1 and BASEVAR2.

Remarks
The most common case of a 32-bit temporary variable per cluster.

Definition at line 341 of file TemporaryHelpers.h.

341#define CALORECGPU_TEMP2DARR_2(TEMPNAME, BASEVAR1, BASEVAR2, TYPE) _Pragma("nv_diag_suppress 177") \
342 template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
343 { \
344 return CaloRecGPU::get_laundered_pointer_striped<const TYPE>(jdx, idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
345 } \
346 template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
347 { \
348 return CaloRecGPU::get_laundered_pointer_striped<TYPE>(jdx, idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
349 } \
350 template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
351 { \
352 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, jdx, idx); \
353 } \
354 template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int jdx, const unsigned int idx) \
355 { \
356 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, jdx, idx); \
357 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon

◆ CALORECGPU_TEMP_CONCAT_HELPER

#define CALORECGPU_TEMP_CONCAT_HELPER ( A,
B )
Value:
CALORECGPU_TEMP_CONCAT_HELPER_INNER(A, B)
CALORECGPU_TEMP_CONCAT_HELPER_INNER
#define CALORECGPU_TEMP_CONCAT_HELPER_INNER(A,...)
Definition TemporaryHelpers.h:10
A
hold the test vectors and ease the comparison

Definition at line 11 of file TemporaryHelpers.h.

◆ CALORECGPU_TEMP_CONCAT_HELPER_INNER

#define CALORECGPU_TEMP_CONCAT_HELPER_INNER ( A,
... )
Value:
A ## __VA_ARGS__

Definition at line 10 of file TemporaryHelpers.h.

◆ CALORECGPU_TEMPARR_1

#define CALORECGPU_TEMPARR_1 ( TEMPNAME,
BASEVAR,
TYPE )
Value:
_Pragma("nv_diag_suppress 177") \
template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer<const TYPE, offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR)>(idx, arr->moments. BASEVAR); \
} \
template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer<TYPE, offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR)>(idx, arr->moments. BASEVAR); \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} _Pragma("nv_diag_default 177") struct to_end_with_semicolon
CaloRecGPU::get_laundered_pointer
__host__ __device__ T * get_laundered_pointer(unsigned int idx, PtrLikes &&... p)

Defines a temporary variable of type TYPE with name TEMPNAME that will be stored in the 1D array accessible through BASEVAR.

Remarks
The most common case of a 32-bit temporary variable per cluster.

Definition at line 269 of file TemporaryHelpers.h.

269#define CALORECGPU_TEMPARR_1(TEMPNAME, BASEVAR, TYPE) _Pragma("nv_diag_suppress 177") \
270 template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
271 { \
272 return CaloRecGPU::get_laundered_pointer<const TYPE, offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR)>(idx, arr->moments. BASEVAR); \
273 } \
274 template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
275 { \
276 return CaloRecGPU::get_laundered_pointer<TYPE, offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR)>(idx, arr->moments. BASEVAR); \
277 } \
278 template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
279 { \
280 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
281 } \
282 template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
283 { \
284 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
285 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon

◆ CALORECGPU_TEMPARR_2

#define CALORECGPU_TEMPARR_2 ( TEMPNAME,
BASEVAR1,
BASEVAR2,
TYPE )
Value:
_Pragma("nv_diag_suppress 177") \
template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer<const TYPE, \
offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR1), \
offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR2)> \
(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
} \
template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer<TYPE, \
offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR1), \
offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR2)> \
(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} _Pragma("nv_diag_default 177") struct to_end_with_semicolon

Defines a temporary variable of type TYPE with name TEMPNAME that will be split across two 1D arrays accessible through BASEVAR1 and BASEVAR2.

Remarks
This would be a 64-bit temporary variable per cluster.

Definition at line 291 of file TemporaryHelpers.h.

291#define CALORECGPU_TEMPARR_2(TEMPNAME, BASEVAR1, BASEVAR2, TYPE) _Pragma("nv_diag_suppress 177") \
292 template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
293 { \
294 return CaloRecGPU::get_laundered_pointer<const TYPE, \
295 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR1), \
296 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR2)> \
297 (idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
298 } \
299 template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
300 { \
301 return CaloRecGPU::get_laundered_pointer<TYPE, \
302 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR1), \
303 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR2)> \
304 (idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
305 } \
306 template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
307 { \
308 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
309 } \
310 template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
311 { \
312 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
313 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon

◆ CALORECGPU_TEMPBIGARR_1

#define CALORECGPU_TEMPBIGARR_1 ( TEMPNAME,
BASEVAR,
TYPE )
Value:
_Pragma("nv_diag_suppress 177") \
template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer_stacked<const TYPE>(idx, arr->moments. BASEVAR); \
} \
template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer_stacked<TYPE>(idx, arr->moments. BASEVAR); \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} _Pragma("nv_diag_default 177") struct to_end_with_semicolon
CaloRecGPU::get_laundered_pointer_stacked
__host__ __device__ T * get_laundered_pointer_stacked(unsigned int idx, PtrLikes &&... ps)

Defines a temporary variable of type TYPE with name TEMPNAME that will be stored as a 1D array split among the entries of the 2D array (i.

e. per sampling) accessible through BASEVAR.

Remarks
This would be enough for a 32-bit temporary variable per possible neighbour pair (NExactPairs) if NMaxClusters == NCaloCells, otherwise it's not enough.

Definition at line 365 of file TemporaryHelpers.h.

365#define CALORECGPU_TEMPBIGARR_1(TEMPNAME, BASEVAR, TYPE) _Pragma("nv_diag_suppress 177") \
366 template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
367 { \
368 return CaloRecGPU::get_laundered_pointer_stacked<const TYPE>(idx, arr->moments. BASEVAR); \
369 } \
370 template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
371 { \
372 return CaloRecGPU::get_laundered_pointer_stacked<TYPE>(idx, arr->moments. BASEVAR); \
373 } \
374 template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
375 { \
376 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
377 } \
378 template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
379 { \
380 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
381 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon

◆ CALORECGPU_TEMPBIGARR_2

#define CALORECGPU_TEMPBIGARR_2 ( TEMPNAME,
BASEVAR1,
BASEVAR2,
TYPE )
Value:
_Pragma("nv_diag_suppress 177") \
template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer_stacked<const TYPE>(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
} \
template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer_stacked<TYPE>(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} _Pragma("nv_diag_default 177") struct to_end_with_semicolon

Defines a temporary variable of type TYPE with name TEMPNAME that will be stored as a 1D array split among the entries of two 2D arrays (i.

e. per sampling) accessible through BASEVAR1 and BASEVAR2.

Remarks
This is always enough for a 32-bit temporary variable per possible neighbour pair (NExactPairs), or even a 64-bit per neighbour pair if NMaxClusters == NCaloCells.

Definition at line 388 of file TemporaryHelpers.h.

388#define CALORECGPU_TEMPBIGARR_2(TEMPNAME, BASEVAR1, BASEVAR2, TYPE) _Pragma("nv_diag_suppress 177") \
389 template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
390 { \
391 return CaloRecGPU::get_laundered_pointer_stacked<const TYPE>(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
392 } \
393 template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
394 { \
395 return CaloRecGPU::get_laundered_pointer_stacked<TYPE>(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2); \
396 } \
397 template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
398 { \
399 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
400 } \
401 template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
402 { \
403 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
404 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon

◆ CALORECGPU_TEMPBIGARR_3

#define CALORECGPU_TEMPBIGARR_3 ( TEMPNAME,
BASEVAR1,
BASEVAR2,
BASEVAR3,
TYPE )
Value:
_Pragma("nv_diag_suppress 177") \
template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer_stacked<const TYPE>(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2, arr->moments. BASEVAR3); \
} \
template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer_stacked<TYPE>(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2, arr->moments. BASEVAR3); \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} _Pragma("nv_diag_default 177") struct to_end_with_semicolon

Defines a temporary variable of type TYPE with name TEMPNAME that will be stored as a 1D array split among the entries of three 2D arrays (i.

e. per sampling) accessible through BASEVAR1, BASEVAR2 and BASEVAR3.

Remarks
This is a good match for fitting two NExactPairs arrays (in different directions) while still being able to fit two in the six per-sampling variables in the moments.

Definition at line 411 of file TemporaryHelpers.h.

411#define CALORECGPU_TEMPBIGARR_3(TEMPNAME, BASEVAR1, BASEVAR2, BASEVAR3, TYPE) _Pragma("nv_diag_suppress 177") \
412 template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
413 { \
414 return CaloRecGPU::get_laundered_pointer_stacked<const TYPE>(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2, arr->moments. BASEVAR3); \
415 } \
416 template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
417 { \
418 return CaloRecGPU::get_laundered_pointer_stacked<TYPE>(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2, arr->moments. BASEVAR3); \
419 } \
420 template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
421 { \
422 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
423 } \
424 template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
425 { \
426 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
427 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon

◆ CALORECGPU_TEMPCELLARR_1

#define CALORECGPU_TEMPCELLARR_1 ( TEMPNAME,
BASEVAR1,
BASEVAR2,
BASEVAR3,
TYPE )
Value:
_Pragma("nv_diag_suppress 177") \
template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer<const TYPE, \
offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR1), \
offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR2), \
offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR3)> \
(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2, arr->moments. BASEVAR3); \
} \
template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
{ \
return CaloRecGPU::get_laundered_pointer<TYPE, \
offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR1), \
offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR2), \
offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR3)> \
(idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2, arr->moments. BASEVAR3); \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
} _Pragma("nv_diag_default 177") struct to_end_with_semicolon

Defines a temporary variable of type TYPE with name TEMPNAME that can have up to three times NMaxClusters, split across three variables.

Remarks
This would be a 32-bit temporary variable per cell regardless of NMaxClusters.

Definition at line 432 of file TemporaryHelpers.h.

432#define CALORECGPU_TEMPCELLARR_1(TEMPNAME, BASEVAR1, BASEVAR2, BASEVAR3, TYPE) _Pragma("nv_diag_suppress 177") \
433 template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
434 { \
435 return CaloRecGPU::get_laundered_pointer<const TYPE, \
436 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR1), \
437 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR2), \
438 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR3)> \
439 (idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2, arr->moments. BASEVAR3); \
440 } \
441 template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
442 { \
443 return CaloRecGPU::get_laundered_pointer<TYPE, \
444 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR1), \
445 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR2), \
446 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR3)> \
447 (idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2, arr->moments. BASEVAR3); \
448 } \
449 template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
450 { \
451 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
452 } \
453 template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
454 { \
455 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
456 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon

◆ CALORECGPU_TEMPCELLARR_2

#define CALORECGPU_TEMPCELLARR_2 ( TEMPNAME,
BASEVAR1,
BASEVAR2,
BASEVAR3,
BASEVAR4,
BASEVAR5,
BASEVAR6,
TYPE )

The definitions for a temporary variable of type TYPE with name TEMPNAME that can have up to three times NMaxClusters, split across six variables.

Remarks
This would be a 64-bit temporary variable per cell regardless of NMaxClusters.

Definition at line 462 of file TemporaryHelpers.h.

462#define CALORECGPU_TEMPCELLARR_2(TEMPNAME, BASEVAR1, BASEVAR2, BASEVAR3, BASEVAR4, BASEVAR5, BASEVAR6, TYPE) _Pragma("nv_diag_suppress 177") \
463 template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr, const unsigned int idx) \
464 { \
465 return CaloRecGPU::get_laundered_pointer<const TYPE, \
466 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR1), \
467 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR2), \
468 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR3), \
469 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR4), \
470 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR5), \
471 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR6)> \
472 (idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2, arr->moments. BASEVAR3, \
473 arr->moments. BASEVAR4, arr->moments. BASEVAR5, arr->moments. BASEVAR6); \
474 } \
475 template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr, const unsigned int idx) \
476 { \
477 return CaloRecGPU::get_laundered_pointer<TYPE, \
478 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR1), \
479 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR2), \
480 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR3), \
481 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR4), \
482 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR5), \
483 offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR6)> \
484 (idx, arr->moments. BASEVAR1, arr->moments. BASEVAR2, arr->moments. BASEVAR3, \
485 arr->moments. BASEVAR4, arr->moments. BASEVAR5, arr->moments. BASEVAR6); \
486 } \
487 template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr, const unsigned int idx) \
488 { \
489 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
490 } \
491 template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr, const unsigned int idx) \
492 { \
493 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr, idx); \
494 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon

◆ CALORECGPU_TEMPVAR

#define CALORECGPU_TEMPVAR ( TEMPNAME,
BASEVAR,
INDEX,
TYPE )
Value:
_Pragma("nv_diag_suppress 177") \
template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr) \
{ \
static_assert(CaloRecGPU::check_sufficient_size<TYPE>(offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR), INDEX)); \
return CaloRecGPU::get_laundered_pointer<const TYPE, offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR)>(INDEX, arr->moments. BASEVAR); \
} \
template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr) \
{ \
static_assert(CaloRecGPU::check_sufficient_size<TYPE>(offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR), INDEX)); \
return CaloRecGPU::get_laundered_pointer<TYPE, offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR)>(INDEX, arr->moments. BASEVAR); \
} \
template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr); \
} \
template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr) \
{ \
return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr); \
} _Pragma("nv_diag_default 177") struct to_end_with_semicolon
CaloRecGPU::check_sufficient_size
constexpr bool __host__ __device__ check_sufficient_size(unsigned int offset, unsigned int index)
Definition TemporaryHelpers.h:178

Defines a temporary variable of type TYPE with name TEMPNAME that will be stored in the INDEX entry of an 1D array.

Remarks
The INDEX represents the index of the array of TYPE, not the offset within the original array.
Warning
The address will be appropriately aligned, so, for bigger types, the maximum number of variables that can be stored in a single array may be smaller, depending on the necessary alignment.

Definition at line 504 of file TemporaryHelpers.h.

504#define CALORECGPU_TEMPVAR(TEMPNAME, BASEVAR, INDEX, TYPE) _Pragma("nv_diag_suppress 177") \
505 template <class PtrLike> __host__ __device__ const TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (const PtrLike & arr) \
506 { \
507 static_assert(CaloRecGPU::check_sufficient_size<TYPE>(offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR), INDEX)); \
508 return CaloRecGPU::get_laundered_pointer<const TYPE, offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR)>(INDEX, arr->moments. BASEVAR); \
509 } \
510 template <class PtrLike> __host__ __device__ TYPE * CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr) (PtrLike & arr) \
511 { \
512 static_assert(CaloRecGPU::check_sufficient_size<TYPE>(offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR), INDEX)); \
513 return CaloRecGPU::get_laundered_pointer<TYPE, offsetof(CaloRecGPU::ClusterInfoArr::ClusterMomentsArr, BASEVAR)>(INDEX, arr->moments. BASEVAR); \
514 } \
515 template <class PtrLike> __host__ __device__ const TYPE & TEMPNAME (const PtrLike & arr) \
516 { \
517 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr); \
518 } \
519 template <class PtrLike> __host__ __device__ TYPE & TEMPNAME (PtrLike & arr) \
520 { \
521 return *CALORECGPU_TEMP_CONCAT_HELPER(TEMPNAME, _ptr)(arr); \
522 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon

◆ CALORECGPU_TEMPWRAPPER

#define CALORECGPU_TEMPWRAPPER ( TEMPNAME,
WRAPPED )
Value:
_Pragma("nv_diag_suppress 177") \
template <class PtrLike, class ... Args> __host__ __device__ decltype(auto) TEMPNAME (PtrLike && p, Args && ... args) \
{ \
return std::forward<PtrLike>(p)-> WRAPPED (std::forward<Args>(args)...); \
} _Pragma("nv_diag_default 177") struct to_end_with_semicolon

Wraps a function in another name, for better semantics while reusing storage...

Definition at line 525 of file TemporaryHelpers.h.

525#define CALORECGPU_TEMPWRAPPER(TEMPNAME, WRAPPED) _Pragma("nv_diag_suppress 177") \
526 template <class PtrLike, class ... Args> __host__ __device__ decltype(auto) TEMPNAME (PtrLike && p, Args && ... args) \
527 { \
528 return std::forward<PtrLike>(p)-> WRAPPED (std::forward<Args>(args)...); \
529 } _Pragma("nv_diag_default 177") struct to_end_with_semicolon
Generated on for ATLAS Offline Software by doxygen 1.14.0