ATLAS Offline Software
Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h
Go to the documentation of this file.
1 //
2 // Copyright (C) 2002-2024 CERN for the benefit of the ATLAS collaboration
3 //
4 // Dear emacs, this is -*- c++ -*-
5 //
6 
7 #ifndef CALORECGPU_HELPERS_H
8 #define CALORECGPU_HELPERS_H
9 
10 #include <utility>
11 #include <type_traits>
12 #include <cstring>
13 //For memcpy, of all things...
14 #include <string>
15 #include <cstdio>
16 #include <iostream>
17 #include <thread>
18 #include <mutex>
19 #include <shared_mutex>
20 #include <memory>
21 #include <vector>
22 #include <climits>
23 #include <new>
24 #include <cmath>
25 
26 #if __cpp_lib_math_constants
27  #include <numbers>
28 #endif
29 //This is the best way to have pi,
30 //but we provide a more manual alternative.
31 //Of course, there's also M_PI,
32 //but we wanted to ensure the type matched
33 //to prevent any GPU-based casting shenanigans.
34 
35 namespace CaloRecGPU
36 {
37 
38 #ifndef CUDA_AVAILABLE
39 
40  #ifdef __CUDA_ARCH__
41  #define CUDA_AVAILABLE 1
42  #elif __CUDA__
43  #define CUDA_AVAILABLE 1
44  #elif __CUDACC__
45  #define CUDA_AVAILABLE 1
46  #else
47  #define CUDA_AVAILABLE 0
48  #endif
49 
50 #endif
51 
52 #if CUDA_AVAILABLE
53 
54 #define CUDA_HOS_DEV __host__ __device__
55 
56 
62  CUDA_HOS_DEV inline void CUDA_gpu_assert(cudaError_t code, const char * file, int line, bool abort = true)
63  {
64  if (code != cudaSuccess)
65  {
66  printf("CUDA error: %s (%s %d)\n", cudaGetErrorString(code), file, line);
67  if (abort)
68  {
69 #ifdef __CUDA_ARCH__
70  asm("trap;");
71 #else
72  exit(code);
73 #endif
74  }
75  }
76  }
77 
82 #define CUDA_ERRCHECK(...) CUDA_ERRCHECK_HELPER(__VA_ARGS__, true)
83 
84 #define CUDA_ERRCHECK_HELPER(ans, ...) do { ::CaloRecGPU::CUDA_gpu_assert((ans), __FILE__, __LINE__, CUDA_ERRCHECK_GET_FIRST(__VA_ARGS__, true) ); } while(0)
85 #define CUDA_ERRCHECK_GET_FIRST(x, ...) x
86 
87 
88 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ > 350
89  #if CUDART_VERSION >= 12000
90  #define CUDA_CAN_USE_TAIL_LAUNCH 1
91  #else
92  #define CUDA_CAN_USE_TAIL_LAUNCH 0
93  #endif
94 #elif defined(__CUDA_ARCH__)
95  #error "CUDA compute capability at least 3.5 is needed so we can have dynamic parallelism!"
96 #endif
97 
98 
99 #else
100 
101 #define CUDA_HOS_DEV
102 #define CUDA_ERRCHECK(...)
103 
104 #endif
105 
106  namespace CUDA_Helpers
107  {
108 
110  {
111  void * ptr = nullptr;
112 
113  template <class T = const void>
114  constexpr operator T * () const
115  {
116  return (T *) ptr;
117  }
118 
119  constexpr operator bool() const
120  {
121  return ptr != nullptr;
122  }
123 
124  template <class T>
126  {
127  }
128 
129  CUDAStreamPtrHolder() = default;
130  };
131  //Can't do much more than this
132  //since cudaStream_t is a typedef...
133  //Though not typesafe, it is still
134  //semantically more safe than a naked void *...
135 
139  void * allocate(const size_t num);
140 
144  void deallocate(void * address);
145 
149  void * allocate_pinned(const size_t num);
150 
155 
156 
160  void GPU_to_CPU(void * dest, const void * const source, const size_t num);
161 
165  void CPU_to_GPU(void * dest, const void * const source, const size_t num);
166 
170  void GPU_to_GPU(void * dest, const void * const source, const size_t num);
171 
172 
176  void GPU_to_CPU_async(void * dest, const void * const source, const size_t num, CUDAStreamPtrHolder stream = {});
177 
181  void CPU_to_GPU_async(void * dest, const void * const source, const size_t num, CUDAStreamPtrHolder stream = {});
182 
186  void GPU_to_GPU_async(void * dest, const void * const source, const size_t num, CUDAStreamPtrHolder stream = {});
187 
192 
196  void optimize_block_and_grid_size(void * func, int & block_size, int & grid_size, const int dynamic_memory = 0, const int block_size_limit = 0);
197 
201  void optimize_block_and_grid_size_for_cooperative_launch(void * func, int & block_size, int & grid_size, const int dynamic_memory = 0, const int block_size_limit = 0);
202 
204 
206 
207  std::string GPU_name();
208  }
209 
210  namespace Helpers
211  {
212 
214  inline constexpr int int_ceil_div(const int num, const int denom)
215  {
216  return num / denom + (num % denom != 0);
217  }
218 
220  inline constexpr int int_floor_div(const int num, const int denom)
221  {
222  return num / denom;
223  }
224 
226  template <class Base = float, class Exp = int>
227  inline constexpr Base compile_time_pow2(const Exp exp)
228  {
229  Base ret = 1;
230  if (exp < 0)
231  {
232  for (Exp i = 0; i < -exp; ++i)
233  {
234  ret /= Base(2);
235  }
236  }
237  else
238  {
239  for (Exp i = 0; i < exp; ++i)
240  {
241  ret *= Base(2);
242  }
243  }
244  return ret;
245  }
246  //Though we could possibly bit-hack stuff due to IEEE-754 reliance elsewhere,
247  //it's not valid and type-safe C++...
248  //Since it's compile-time, this being a trifle slower is meaningless.
249 
250 
253  template <class T>
254  inline constexpr unsigned char Pearson_hash(const T number)
255  {
256  constexpr unsigned char initial_value = 42;
257  //The answer.
258 
259  constexpr unsigned char c_mult = 7;
260  constexpr unsigned char c_add = 1;
261  //For our "look up table": table[i] = c_mult * i + c_add
262  //For an appropriate choice of constants (such as this),
263  //this will be bijective (modulo 255), as required.
264 
265  unsigned char ret = initial_value;
266 
267  for (unsigned int i = 0; i < sizeof(T); i += sizeof(unsigned char))
268  {
269  const unsigned char to_hash = number >> (i * CHAR_BIT);
270  const unsigned char operand = ret ^ to_hash;
271  ret = c_mult * operand + c_add;
272  }
273 
274  return ret;
275  }
276 
277 
280  template <class T>
281  inline constexpr unsigned short Pearson_hash_16_bit(const T number)
282  {
283  constexpr unsigned short initial_value = 42754;
284  //The answer and the standard.
285 
286  constexpr unsigned short c_mult = 7;
287  constexpr unsigned short c_add = 1;
288  //For our "look up table": table[i] = c_mult * i + c_add
289  //For an appropriate choice of constants (such as this),
290  //this will be bijective (modulo 255), as required.
291 
292  unsigned short ret = initial_value;
293 
294  for (unsigned int i = 0; i < sizeof(T); i += sizeof(unsigned short))
295  {
296  const unsigned short to_hash = number >> (i * CHAR_BIT);
297  const unsigned short operand = ret ^ to_hash;
298  ret = c_mult * operand + c_add;
299  }
300 
301  return ret;
302  }
303 
304 
306  namespace Constants
307  {
308 #ifdef __cpp_lib_math_constants
309  template <class T>
310  inline constexpr T pi = std::numbers::pi_v<T>;
311 
312  template <class T>
313  inline constexpr T sqrt2 = std::numbers::sqrt2_v<T>;
314 #else
315  template <class T>
316  inline constexpr T pi = T(3.1415926535897932384626433832795028841971693993751058209749445923078164062862089986280348253421170679821480865132823066470938446095505822317253594081284811174502841027019385211055596446229489549303819644288109756659334461284756482337867831652712019091456485669234603486104543266482133936072602491412737245870066063155881748815209209628292540917153643678925903600113305305488204665213841469519415116094330572703657595919530921861173819326117931051185480744623799627495673518857527248912279381830119491298336733624L);
317 
318  template <class T>
319  inline constexpr T sqrt2 = T(1.4142135623730950488016887242096980785696718753769480731766797379907324784621070388503875343276415727350138462309122970249248360558507372126441214970999358314132226659275055927557999505011527820605714701095599716059702745345968620147285174186408891986095523292304843087143214508397626036279952514079896872533965463318088296406206152583523950547457502877599617298355752203375318570113543746034084988471603868999706990048150305440277903164542478230684929369186215805784631115966687130130156185689872372352885092649L);
320 #endif
321 
322  template <class T>
323  inline constexpr T inv_sqrt2 = T(0.70710678118654752440084436210484903928483593768847403658833986899536623923105351942519376716382078636750692311545614851246241802792536860632206074854996791570661133296375279637789997525057639103028573505477998580298513726729843100736425870932044459930477616461524215435716072541988130181399762570399484362669827316590441482031030762917619752737287514387998086491778761016876592850567718730170424942358019344998534950240751527201389515822712391153424646845931079028923155579833435650650780928449361861764425463243L);
324  //Why is this not in the C++ constants?!
325 
326  }
327 
328  CUDA_HOS_DEV static inline
329  float erf_inv_wrapper (const float x)
330  {
331  using namespace std;
332 #ifdef __CUDA_ARCH__
333  return erfinvf(x);
334 #else
335  //Copied directly from ROOT...
336 
337  int kMaxit = 50;
338  float kEps = 1e-14;
339  float kConst = 0.8862269254527579; // sqrt(pi)/2.0
340 
341  if (abs(x) <= kEps)
342  {
343  return kConst * x;
344  }
345 
346  // Newton iterations
347  float erfi, derfi, y0, y1, dy0, dy1;
348  if (fabsf(x) < 1.0f)
349  {
350  erfi = kConst * fabsf(x);
351  y0 = erff(0.9f * erfi);
352  derfi = 0.1 * erfi;
353  for (int iter = 0; iter < kMaxit; iter++)
354  {
355  y1 = 1. - erfc(erfi);
356  dy1 = fabsf(x) - y1;
357  if (fabsf(dy1) < kEps)
358  {
359  if (x < 0)
360  {
361  return -erfi;
362  }
363  else
364  {
365  return erfi;
366  }
367  }
368  dy0 = y1 - y0;
369  derfi *= dy1 / dy0;
370  y0 = y1;
371  erfi += derfi;
372  if (fabsf(derfi / erfi) < kEps)
373  {
374  if (x < 0)
375  {
376  return -erfi;
377  }
378  else
379  {
380  return erfi;
381  }
382  }
383  }
384  }
385  return 0; //did not converge
386 #endif
387  }
388 
389  CUDA_HOS_DEV static inline
390  float regularize_angle(const float b, const float a = 0.f)
391  //a. k. a. proxim in Athena code.
392  {
393  using namespace std;
394  const float diff = b - a;
395  const float divi = (fabsf(diff) - Helpers::Constants::pi<float>) / (2 * Helpers::Constants::pi<float>);
396  return b - ceilf(divi) * ((b > a + Helpers::Constants::pi<float>) - (b < a - Helpers::Constants::pi<float>)) * 2 * Helpers::Constants::pi<float>;
397  }
398 
399  CUDA_HOS_DEV static inline
400  double regularize_angle(const double b, const double a = 0.)
401  //a. k. a. proxim in Athena code.
402  {
403  using namespace std;
404  const float diff = b - a;
405  const float divi = (fabs(diff) - Helpers::Constants::pi<double>) / (2 * Helpers::Constants::pi<double>);
406  return b - ceil(divi) * ((b > a + Helpers::Constants::pi<double>) - (b < a - Helpers::Constants::pi<double>)) * 2 * Helpers::Constants::pi<double>;
407  }
408 
409  template <class T>
410  CUDA_HOS_DEV static inline
411  T angular_difference(const T x, const T y)
412  {
413  return regularize_angle(x - y, T(0));
414  //Might be problematic if x and y have a significant difference
415  //in terms of factors of pi, in which case one should add
416  //a regularize_angle(x) and regularize_angle(y) in there.
417  //For our use case, I think this will be fine.
418  //(The Athena ones are even worse,
419  // being a branchy thing that only
420  // takes care of one factor of 2 pi...)
421  }
422 
423  CUDA_HOS_DEV static inline
424  float eta_from_coordinates(const float x, const float y, const float z)
425  {
426  using namespace std;
427  const float rho2 = x * x + y * y;
428  if (rho2 > 0.)
429  {
430  const float m = sqrtf(rho2 + z * z);
431  return 0.5 * logf((m + z) / (m - z));
432  }
433  else
434  {
435  constexpr float s_etaMax = 22756.0;
436  return z + ((z > 0) - (z < 0)) * s_etaMax;
437  }
438  }
439 
440  CUDA_HOS_DEV static inline
441  double eta_from_coordinates(const double x, const double y, const double z)
442  {
443  using namespace std;
444  const double rho2 = x * x + y * y;
445  if (rho2 > 0.)
446  {
447  const double m = sqrt(rho2 + z * z);
448  return 0.5 * log((m + z) / (m - z));
449  }
450  else
451  {
452  constexpr double s_etaMax = 22756.0;
453  return z + ((z > 0) - (z < 0)) * s_etaMax;
454  }
455  }
456 
458  namespace MemoryContext
459  {
460  struct CPU
461  {
462  constexpr static char const * name = "CPU";
463  };
464  struct CUDAGPU
465  {
466  constexpr static char const * name = "CUDA GPU";
467  };
469  {
470  constexpr static char const * name = "CUDA Pinned CPU";
471  };
472  }
473 
475  template <class T, class indexer>
477  {
478  private:
479  template <class C, class dummy = void> struct unary_helper;
480 
481  template <class dummy> struct unary_helper<MemoryContext::CPU, dummy>
482  {
483  static inline T * allocate(const indexer size)
484  {
485  return new T[size];
486  }
487 
488  static inline void deallocate(T *& arr)
489  {
490  delete[] arr;
491  }
492 
493  };
494 
495  template <class dummy> struct unary_helper<MemoryContext::CUDAGPU, dummy>
496  {
497  static inline T * allocate(const indexer size)
498  {
499  return static_cast<T *>(CUDA_Helpers::allocate(sizeof(T) * size));
500  }
501 
502  static inline void deallocate(T *& arr)
503  {
505  }
506  };
507 
508 
509  template <class dummy> struct unary_helper<MemoryContext::CUDAPinnedCPU, dummy>
510  {
511  static inline T * allocate(const indexer size)
512  {
513  return static_cast<T *>(CUDA_Helpers::allocate_pinned(sizeof(T) * size));
514  }
515 
516  static inline void deallocate(T *& arr)
517  {
519  }
520  };
521 
522  template <class C1, class C2, class dummy = void> struct copy_helper;
523 
524  template <class dummy> struct copy_helper<MemoryContext::CPU, MemoryContext::CPU, dummy>
525  {
526  static inline void copy (T * dest, const T * const source, const indexer sz)
527  {
528  std::memcpy(dest, source, sizeof(T) * sz);
529  }
530  };
531 
532  template <class dummy> struct copy_helper<MemoryContext::CPU, MemoryContext::CUDAGPU, dummy>
533  {
534  static inline void copy (T * dest, const T * const source, const indexer sz)
535  {
536  CUDA_Helpers::GPU_to_CPU(dest, source, sizeof(T) * sz);
537  }
538  };
539 
540  template <class dummy> struct copy_helper<MemoryContext::CUDAGPU, MemoryContext::CUDAGPU, dummy>
541  {
542  static inline void copy (T * dest, const T * const source, const indexer sz)
543  {
544  CUDA_Helpers::GPU_to_GPU(dest, source, sizeof(T) * sz);
545  }
546  };
547 
548  template <class dummy> struct copy_helper<MemoryContext::CUDAGPU, MemoryContext::CPU, dummy>
549  {
550  static inline void copy (T * dest, const T * const source, const indexer sz)
551  {
552  CUDA_Helpers::CPU_to_GPU(dest, source, sizeof(T) * sz);
553  }
554  };
555 
556  template <class dummy> struct copy_helper<MemoryContext::CUDAPinnedCPU, MemoryContext::CPU, dummy>
557  {
558  static inline void copy (T * dest, const T * const source, const indexer sz)
559  {
560  std::memcpy(dest, source, sizeof(T) * sz);
561  }
562  };
563 
564  template <class dummy> struct copy_helper<MemoryContext::CPU, MemoryContext::CUDAPinnedCPU, dummy>
565  {
566  static inline void copy (T * dest, const T * const source, const indexer sz)
567  {
568  std::memcpy(dest, source, sizeof(T) * sz);
569  }
570  };
571 
572  template <class dummy> struct copy_helper<MemoryContext::CUDAPinnedCPU, MemoryContext::CUDAPinnedCPU, dummy>
573  {
574  static inline void copy (T * dest, const T * const source, const indexer sz)
575  {
576  std::memcpy(dest, source, sizeof(T) * sz);
577  }
578  };
579 
580  template <class dummy> struct copy_helper<MemoryContext::CUDAPinnedCPU, MemoryContext::CUDAGPU, dummy>
581  {
582  static inline void copy (T * dest, const T * const source, const indexer sz)
583  {
584  CUDA_Helpers::GPU_to_CPU(dest, source, sizeof(T) * sz);
585  }
586  };
587 
588  template <class dummy> struct copy_helper<MemoryContext::CUDAGPU, MemoryContext::CUDAPinnedCPU, dummy>
589  {
590  static inline void copy (T * dest, const T * const source, const indexer sz)
591  {
592  CUDA_Helpers::CPU_to_GPU(dest, source, sizeof(T) * sz);
593  }
594  };
595 
596 
597  template <class C1, class C2, class dummy = void> struct move_helper;
598 
599  template <class C1, class C2, class dummy> struct move_helper
600  {
601  inline static void move(T *& dest, T *& source, const indexer sz)
602  {
606  }
607  };
608 
609  template <class C, class dummy> struct move_helper<C, C, dummy>
610  {
611  inline static void move(T *& dest, T *& source, const indexer)
612  {
613  dest = source;
614  source = nullptr;
615  }
616  };
617 
618  public:
620  template <class Context> static inline T * allocate(const indexer size)
621  {
622  T * ret = nullptr;
623  if (size > 0)
624  {
626  }
627 #if CALORECGPU_HELPERS_DEBUG
628  std::cerr << "ALLOCATED " << size << " x " << sizeof(T) << " in " << Context::name << ": " << ret << std::endl;
629 #endif
630  return ret;
631  }
632 
634  template <class Context> static inline void deallocate(T *& arr)
635  {
636  if (arr == nullptr)
637  //This check is to ensure the code behaves on non-CUDA enabled platforms
638  //where some destructors might still be called with nullptr.
639  {
640  return;
641  }
643 #if CALORECGPU_HELPERS_DEBUG
644  std::cerr << "DEALLOCATED in " << Context::name << ": " << arr << std::endl;
645 #endif
646  arr = nullptr;
647  }
648 
649 
651  template <class DestContext, class SourceContext>
652  static inline void copy(T * dest, const T * const source, const indexer sz)
653  {
654  if (sz > 0 && source != nullptr)
655  {
657  }
658 #if CALORECGPU_HELPERS_DEBUG
659  std::cerr << "COPIED " << sz << " from " << SourceContext::name << " to " << DestContext::name << ": " << source << " to " << dest << std::endl;
660 #endif
661  }
662 
663 
669  template <class DestContext, class SourceContext>
670  static inline void move(T *& dest, T *& source, const indexer sz)
671  {
672 #if CALORECGPU_HELPERS_DEBUG
673  std::cerr << "MOVED " << sz << " from " << SourceContext::name << " to " << DestContext::name << ": " << source << " to " << dest;
674 #endif
675  if (sz > 0 && source != nullptr)
676  {
678  }
679  else
680  {
681  dest = nullptr;
682  deallocate<SourceContext>(source);
683  }
684 #if CALORECGPU_HELPERS_DEBUG
685  std::cerr << " | " << source << " to " << dest << std::endl;
686 #endif
687  }
688 
689  };
690 
697  template <class T, class indexer, class Context, bool hold_arrays = true>
699 
700  template <class T, class indexer, class Context>
701  class SimpleContainer<T, indexer, Context, true>
702  {
703  static_assert(std::is_trivially_copyable<T>::value, "SimpleContainer only works with a trivially copyable type.");
704  T * m_array;
705  indexer m_size;
706 
707  template <class a, class b, class c, bool d> friend class SimpleContainer;
708 
710 
711  public:
712 
713  CUDA_HOS_DEV inline indexer size() const
714  {
715  return m_size;
716  }
717 
718  CUDA_HOS_DEV inline T & operator[] (const indexer i)
719  {
720  return m_array[i];
721  }
722 
723  CUDA_HOS_DEV inline const T & operator[] (const indexer i) const
724  {
725  return m_array[i];
726  }
727 
728  inline void clear()
729  {
730  Manager::deallocate(m_array);
731  m_size = 0;
732  }
733 
734  inline void resize(const indexer new_size)
735  {
736  if (new_size == 0)
737  {
738  clear();
739  }
740  else if (new_size != m_size)
741  {
742  T * temp = m_array;
743  m_array = Manager::template allocate<Context>(new_size);
744  Manager::template copy<Context, Context>(m_array, temp, (m_size < new_size ? m_size : new_size));
745  Manager::template deallocate<Context>(temp);
746  m_size = new_size;
747  }
748  }
749 
750  SimpleContainer() : m_array(nullptr), m_size(0)
751  {
752  }
753 
754  SimpleContainer(const indexer sz) : m_size(sz)
755  {
756  m_array = Manager::template allocate<Context>(sz);
757  }
758 
762  SimpleContainer(T * other_array, const indexer sz) : m_size(sz)
763  {
764  m_array = Manager::template allocate<Context>(sz);
765  Manager::template copy<Context, Context>(m_array, other_array, sz);
766  }
767 
768  SimpleContainer(const SimpleContainer & other) : m_size(other.m_size)
769  {
770  m_array = Manager::template allocate<Context>(m_size);
771  Manager::template copy<Context, Context>(m_array, other.m_array, m_size);
772  }
773 
775  {
776  m_array = nullptr;
777  Manager::template move<Context, Context>(m_array, other.m_array, m_size);
778  other.m_size = 0;
779  }
780 
781  template <class other_indexer, class other_context, bool other_hold>
783  m_size(other.m_size)
784  {
785 
786  m_array = Manager::template allocate<Context>(m_size);
787  Manager::template copy<Context, other_context>(m_array, other.m_array, m_size);
788  }
789 
790  template <class other_indexer, class other_context>
792  m_size(other.m_size)
793  {
794  m_array = nullptr;
795  Manager::template move<Context, other_context>(m_array, other.m_array, m_size);
796  other.m_size = 0;
797  }
798 
799  SimpleContainer & operator= (const SimpleContainer & other)
800  {
801  if (this == &other)
802  {
803  return (*this);
804  }
805  else
806  {
807  resize(other.size());
808  Manager::template copy<Context, Context>(m_array, other.m_array, m_size);
809  return (*this);
810  }
811  }
812 
814  {
815  if (this == &other)
816  {
817  return (*this);
818  }
819  else
820  {
821  clear();
822  Manager::template move<Context, Context>(m_array, other.m_array, other.size());
823  m_size = other.m_size;
824  other.m_size = 0;
825  return (*this);
826  }
827  }
828 
829 
830  template <class other_indexer, class other_context, bool other_hold>
832  {
833  resize(other.m_size);
834  Manager::template copy<Context, other_context>(m_array, other.m_array, m_size);
835  return (*this);
836  }
837 
838  template <class other_indexer, class other_context>
840  {
841  clear();
842  Manager::template move<Context, other_context>(m_array, other.m_array, other.m_size);
843  m_size = other.m_size;
844  other.m_size = 0;
845  return (*this);
846  }
847 
849  {
850  Manager::template deallocate<Context>(m_array);
851  m_size = 0;
852  }
853 
854  CUDA_HOS_DEV operator const T * () const
855  {
856  return m_array;
857  }
858 
859  CUDA_HOS_DEV operator T * ()
860  {
861  return m_array;
862  }
863 
864  template <class stream, class str = std::basic_string<typename stream::char_type> >
865  void textual_output(stream & s, const str & separator = " ") const
866  {
868  {
869  s << m_size << separator;
870  for (indexer i = 0; i < m_size - 1; ++i)
871  {
872  s << m_array[i] << separator;
873  }
874  s << m_array[m_size - 1];
875  }
876  else
877  {
879  other.textual_output(s, separator);
880  }
881  }
882 
883  template <class stream>
885  {
887  {
888  indexer new_size;
889  s >> new_size >> std::ws;
890  if (s.fail())
891  {
892  //Throw errors, perhaps? Don't know if we can/should use exceptions...
893  std::cerr << "FAILED READING " << this << "!" << std::endl;
894  new_size = 0;
895  }
896  resize(new_size);
897  for (indexer i = 0; i < m_size - 1; ++i)
898  {
899  s >> m_array[i];
900  s >> std::ws;
901  }
902  s >> m_array[m_size - 1];
903  }
904  else
905  {
907  other.textual_input(s);
908  (*this) = other;
909  }
910  }
911 
912  template <class stream>
913  void binary_output(stream & s) const
914  {
916  {
917  s.write((char *) &m_size, sizeof(indexer));
918  for (indexer i = 0; i < m_size; ++i)
919  {
920  s.write((char *) (m_array + i), sizeof(T));
921  }
922  }
923  else
924  {
926  other.binary_output(s);
927  }
928  }
929 
930  template <class stream>
932  {
934  {
935  indexer new_size;
936  s.read((char *) &new_size, sizeof(indexer));
937  if (s.fail())
938  {
939  //Throw errors, perhaps? Don't know if we can/should use exceptions...
940  std::cerr << "FAILED READING " << this << "!" << std::endl;
941  new_size = 0;
942  }
943  resize(new_size);
944  for (indexer i = 0; i < m_size; ++i)
945  {
946  s.read((char *) (m_array + i), sizeof(T));
947  }
948  }
949  else
950  {
952  other.binary_input(s);
953  (*this) = other;
954  }
955  }
956 
957  };
958 
959  template <class T, class indexer, class Context>
960  class SimpleContainer<T, indexer, Context, false>
961  {
962  static_assert(std::is_trivially_copyable<T>::value, "SimpleContainer only works with a trivially copyable type.");
963  T * m_array;
964  indexer m_size;
965 
967 
968  template <class a, class b, class c, bool d> friend class SimpleContainer;
969 
970  public:
971 
972  CUDA_HOS_DEV inline indexer size() const
973  {
974  return m_size;
975  }
976 
977  CUDA_HOS_DEV inline T & operator[] (const indexer i)
978  {
979  return m_array[i];
980  }
981 
982  CUDA_HOS_DEV inline const T & operator[] (const indexer i) const
983  {
984  return m_array[i];
985  }
986 
987  // cppcheck-suppress uninitMemberVar
988  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
989  CUDA_HOS_DEV SimpleContainer() : m_array(nullptr), m_size(0)
990  {
991  }
992 
996  // cppcheck-suppress uninitMemberVar
997  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
998  CUDA_HOS_DEV SimpleContainer(T * other_array, const indexer sz) : m_array(other_array), m_size(sz)
999  {
1000  }
1001 
1002  template <class other_indexer, bool other_hold>
1003  // cppcheck-suppress uninitMemberVar
1004  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1006  m_size(other.m_size),
1007  m_array(other.m_array)
1008  {
1009  }
1010 
1011  // cppcheck-suppress operatorEqVarError
1012  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1014  {
1015  if (this == &other)
1016  {
1017  return (*this);
1018  }
1019  else
1020  {
1021  m_array = other.m_array;
1022  m_size = other.m_size;
1023  return (*this);
1024  }
1025  }
1026 
1027  template <class other_indexer, bool other_hold>
1028  // cppcheck-suppress operatorEqVarError
1029  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1031  {
1032  m_size = other.m_size;
1033  m_array = other.m_array;
1034  return (*this);
1035  }
1036 
1037  CUDA_HOS_DEV operator const T * () const
1038  {
1039  return m_array;
1040  }
1041 
1042  CUDA_HOS_DEV operator T * ()
1043  {
1044  return m_array;
1045  }
1046  };
1047 
1049  template <class T, class indexer = unsigned int>
1051 
1053  template <class T, class indexer = unsigned int>
1055 
1057  template <class T, class indexer = unsigned int>
1059 
1066  template <class T, class Context, bool hold_object = true>
1068 
1069  template <class T, class Context>
1070  class SimpleHolder<T, Context, true>
1071  {
1072  static_assert(std::is_trivially_copyable<T>::value, "SimpleHolder only works with a trivially copyable type.");
1073 
1074  using indexer = unsigned int;
1075 
1077 
1079 
1080  template <class a, class b, bool c> friend class SimpleHolder;
1081 
1082  public:
1083 
1084  CUDA_HOS_DEV const T & operator *() const
1085  {
1086  return *m_object;
1087  }
1088 
1090  {
1091  return *m_object;
1092  }
1093 
1094  CUDA_HOS_DEV const T * operator ->() const
1095  {
1096  return m_object;
1097  }
1098 
1099  CUDA_HOS_DEV T * operator ->()
1100  {
1101  return m_object;
1102  }
1103 
1104  CUDA_HOS_DEV inline bool valid() const
1105  {
1106  return m_object != nullptr;
1107  }
1108 
1109  inline void clear()
1110  {
1111  Manager::template deallocate<Context>(m_object);
1112  }
1113 
1114  inline void allocate()
1115  {
1116  if (m_object == nullptr)
1117  {
1118  m_object = Manager::template allocate<Context>(1);
1119  }
1120  }
1121 
1122  SimpleHolder(): m_object(nullptr)
1123  {
1124  }
1125 
1126  SimpleHolder(const bool really_allocate)
1127  {
1128  if (really_allocate)
1129  {
1130  m_object = Manager::template allocate<Context>(1);
1131  }
1132  else
1133  {
1134  m_object = nullptr;
1135  }
1136  }
1137 
1142  explicit SimpleHolder(X * other_p)
1143  {
1144  m_object = Manager::template allocate<Context>(1);
1145  Manager::template copy<Context, Context>(m_object, other_p, 1);
1146  }
1147 
1152  SimpleHolder(const X & other_v) : SimpleHolder(&other_v)
1153  {
1154 
1155 
1156  }
1157 
1159  {
1160  if (other.valid())
1161  {
1162  m_object = Manager::template allocate<Context>(1);
1163  Manager::template copy<Context, Context>(m_object, other.m_object, other.valid());
1164  }
1165  else
1166  {
1167  m_object = nullptr;
1168  }
1169  }
1170 
1171  template < class X, class other_context, bool other_hold,
1174  {
1175  if (other.valid())
1176  {
1177  m_object = Manager::template allocate<Context>(1);
1178  Manager::template copy<Context, other_context>(m_object, other.m_object, other.valid());
1179  }
1180  else
1181  {
1182  m_object = nullptr;
1183  }
1184  }
1185 
1187  {
1188  m_object = nullptr;
1189  Manager::template move<Context, Context>(m_object, other.m_object, other.valid());
1190  }
1191 
1192  template < class X, class other_context,
1195  {
1196  m_object = nullptr;
1197  Manager::template move<Context, other_context>(m_object, other.m_object, other.valid());
1198  }
1199 
1200  // cppcheck-suppress operatorEqVarError
1201  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1202  SimpleHolder & operator= (const SimpleHolder & other)
1203  {
1204  if (!valid() && other.valid())
1205  {
1206  allocate();
1207  }
1208  if (&other != this)
1209  {
1210  Manager::template copy<Context, Context>(m_object, other.m_object, other.valid());
1211  }
1212  return (*this);
1213  }
1214 
1215  template < class X, class other_context, bool other_hold,
1217  // cppcheck-suppress operatorEqVarError
1218  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1220  {
1221  if (!valid() && other.valid())
1222  {
1223  allocate();
1224  }
1225  Manager::template copy<Context, other_context>(m_object, other.m_object, other.valid());
1226  return (*this);
1227  }
1228 
1229  // cppcheck-suppress operatorEqVarError
1230  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1232  {
1233  if (&other != this)
1234  {
1235  clear();
1236  Manager::template move<Context, Context>(m_object, other.m_object, other.valid());
1237  }
1238  return (*this);
1239  }
1240 
1241  template < class X, class other_context,
1243  // cppcheck-suppress operatorEqVarError
1244  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1246  {
1247  clear();
1248  Manager::template move<Context, other_context>(m_object, other.m_object, other.valid());
1249  return (*this);
1250  }
1251 
1253  {
1254  Manager::template deallocate<Context>(m_object);
1255  }
1256 
1258  CUDA_HOS_DEV operator const X * () const
1259  {
1260  return m_object;
1261  }
1262 
1264  CUDA_HOS_DEV operator X * ()
1265  {
1266  return m_object;
1267  }
1268 
1269  template <class stream, class str = std::basic_string<typename stream::char_type> >
1270  void textual_output(stream & s, const str & separator = " ") const
1271  {
1273  {
1274  if (m_object == nullptr)
1275  {
1276  s << 0;
1277  }
1278  else
1279  {
1280  s << 1 << separator << (*m_object);
1281  }
1282  }
1283  else
1284  {
1286  other.textual_output(s, separator);
1287  }
1288  }
1289 
1290  template <class stream>
1292  {
1294  {
1295  bool is_valid;
1296  s >> is_valid >> std::ws;
1297  if (s.fail())
1298  {
1299  //Throw errors, perhaps? Don't know if we can/should use exceptions...
1300  std::cerr << "FAILED READING " << this << "!" << std::endl;
1301  is_valid = false;
1302  }
1303  if (is_valid)
1304  {
1305  allocate();
1306  s >> (*m_object);
1307  }
1308  else
1309  {
1310  clear();
1311  }
1312  }
1313  else
1314  {
1316  other.textual_input(s);
1317  (*this) = other;
1318  }
1319  }
1320 
1321  template <class stream>
1322  void binary_output(stream & s) const
1323  {
1324  if (m_object == nullptr)
1325  {
1326  return;
1327  }
1329  {
1330  s.write((char *) m_object, sizeof(T));
1331  }
1332  else
1333  {
1335  other.binary_output(s);
1336  }
1337  }
1338 
1339  template <class stream>
1341  {
1343  {
1344  allocate();
1345  s.read((char *) m_object, sizeof(T));
1346  }
1347  else
1348  {
1350  other.binary_input(s);
1351  (*this) = other;
1352  }
1353  }
1354 
1355  };
1356 
1357  template <class T, class Context>
1358  class SimpleHolder<T, Context, false>
1359  {
1360  static_assert(std::is_trivially_copyable<T>::value, "SimpleHolder only works with a trivially copyable type.");
1361 
1362  using indexer = unsigned int;
1363 
1365 
1367 
1368  template <class a, class b, bool c> friend class SimpleHolder;
1369 
1370  public:
1371 
1372  CUDA_HOS_DEV const T & operator *() const
1373  {
1374  return *m_object;
1375  }
1376 
1378  {
1379  return *m_object;
1380  }
1381 
1382  CUDA_HOS_DEV const T * operator ->() const
1383  {
1384  return m_object;
1385  }
1386 
1387  CUDA_HOS_DEV T * operator ->()
1388  {
1389  return m_object;
1390  }
1391 
1392  CUDA_HOS_DEV inline bool valid() const
1393  {
1394  return m_object != nullptr;
1395  }
1396 
1397  // cppcheck-suppress uninitMemberVar
1398  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1399  CUDA_HOS_DEV SimpleHolder() : m_object(nullptr)
1400  {
1401  }
1402 
1407  // cppcheck-suppress uninitMemberVar
1408  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1410  {
1411  m_object = other_p;
1412  }
1413 
1414  template < class X, bool other_hold,
1416  // cppcheck-suppress uninitMemberVar
1417  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1419  {
1420  m_object = other.m_object;
1421  }
1422 
1423  template < class X, bool other_hold,
1425  // cppcheck-suppress operatorEqVarError
1426  //Try to suppress the uninitialized member thing that is probably being thrown off by the CUDA_HOS_DEV macro...
1428  {
1429  m_object = other.m_object;
1430  return (*this);
1431  }
1432 
1434  CUDA_HOS_DEV operator const X * () const
1435  {
1436  return m_object;
1437  }
1438 
1440  CUDA_HOS_DEV operator X * ()
1441  {
1442  return m_object;
1443  }
1444  };
1445 
1447  template <class T>
1449 
1451  template <class T>
1453 
1455  template <class T>
1457 
1459  template <class T>
1461 
1471  template <class T>
1473  {
1474  private:
1475  std::vector< std::unique_ptr<T> > m_held;
1476  std::vector< typename std::thread::id > m_thread_equivs;
1477  //For a sufficiently small number of threads
1478  //(not much more than 100 or so?)
1479  //it's faster to have linear search+insert
1480  //than any other addressing mode
1481  //(e. g. unordered_map)
1482  //We could still consider a more sophisticated solution...
1483 
1484  //Simple alternative: some sort of stack for non-assigned objects,
1485  //pushing and popping instead of linear searching.
1486  //(But with constant memory -> no (de)allocations.)
1487 
1488  mutable std::shared_mutex m_mutex;
1489 
1491  {
1492  std::unique_lock<std::shared_mutex> lock(m_mutex);
1493  m_held.emplace_back(std::make_unique<T>());
1494  m_thread_equivs.emplace_back(std::this_thread::get_id());
1495  return *(m_held.back());
1496  }
1497 
1498  public:
1499  T & get_one()
1500  {
1501  {
1502  std::shared_lock<std::shared_mutex> lock(m_mutex);
1503  std::thread::id this_id = std::this_thread::get_id();
1504  const std::thread::id invalid_id{};
1505  for (size_t i = 0; i < m_thread_equivs.size(); ++i)
1506  {
1507  if (m_thread_equivs[i] == invalid_id)
1508  {
1509  m_thread_equivs[i] = this_id;
1510  return *(m_held[i]);
1511  }
1512  }
1513  }
1514  return add_one_and_return();
1515  }
1516 
1518  T & get_for_thread() const
1519  {
1520  std::shared_lock<std::shared_mutex> lock(m_mutex);
1521  std::thread::id this_id = std::this_thread::get_id();
1522  for (size_t i = 0; i < m_thread_equivs.size(); ++i)
1523  {
1524  if (m_thread_equivs[i] == this_id)
1525  {
1526  return *(m_held[i]);
1527  }
1528  }
1529  //Here would be a good place for an unreachable.
1530  //C++23?
1531  return *(m_held.back());
1532  }
1533 
1535  {
1536  std::unique_lock<std::shared_mutex> lock(m_mutex);
1537  std::thread::id this_id = std::this_thread::get_id();
1538  const std::thread::id invalid_id{};
1539  for (size_t i = 0; i < m_thread_equivs.size(); ++i)
1540  {
1541  if (m_thread_equivs[i] == this_id)
1542  {
1543  m_thread_equivs[i] = invalid_id;
1544  }
1545  }
1546  }
1547 
1548  void resize(const size_t new_size)
1549  {
1550  std::unique_lock<std::shared_mutex> lock(m_mutex);
1551  if (new_size < m_held.size())
1552  {
1553  m_held.resize(new_size);
1554  m_thread_equivs.resize(new_size);
1555  }
1556  else if (new_size > m_held.size())
1557  {
1558  const size_t to_add = new_size - m_held.size();
1559  const std::thread::id invalid_id{};
1560  for (size_t i = 0; i < to_add; ++i)
1561  {
1562  m_held.emplace_back(std::make_unique<T>());
1563  m_thread_equivs.emplace_back(invalid_id);
1564  }
1565  }
1566  }
1567 
1568  template <class F, class ... Args>
1569  void operate_on_all(F && f, Args && ... args)
1570  {
1571  std::unique_lock<std::shared_mutex> lock(m_mutex);
1572  for (std::unique_ptr<T> & obj : m_held)
1573  {
1574  f(*obj, std::forward<Args>(args)...);
1575  }
1576  }
1577 
1578  size_t held_size() const
1579  {
1580  std::shared_lock<std::shared_mutex> lock(m_mutex);
1581  return m_held.size();
1582  }
1583 
1584  size_t available_size() const
1585  {
1586  std::shared_lock<std::shared_mutex> lock(m_mutex);
1587  size_t count = 0;
1588  const std::thread::id invalid_id{};
1589  for (const auto & id : m_thread_equivs)
1590  {
1591  if (id == invalid_id)
1592  {
1593  ++count;
1594  }
1595  }
1596  return count;
1597  }
1598 
1599  size_t filled_size() const
1600  {
1601  std::shared_lock<std::shared_mutex> lock(m_mutex);
1602  size_t count = 0;
1603  const std::thread::id invalid_id{};
1604  for (const auto & id : m_thread_equivs)
1605  {
1606  if (id == invalid_id)
1607  {
1608  ++count;
1609  }
1610  }
1611  return m_held.size() - count;
1612  }
1613  };
1614 
1617  template <class T>
1619  {
1620  private:
1622  T * m_held;
1623  public:
1625  m_sth(s), m_held(nullptr)
1626  {
1627  }
1628  T & get_one()
1629  {
1630  if (m_held == nullptr)
1631  {
1632  m_held = &(m_sth.get_one());
1633  }
1634  return *m_held;
1635  }
1637  {
1638  if (m_held != nullptr)
1639  {
1640  m_sth.release_one();
1641  m_held = nullptr;
1642  }
1643  }
1645  {
1646  if (m_held != nullptr)
1647  {
1648  m_sth.release_one();
1649  }
1650  }
1653  {
1654  get_one();
1655  ptr = m_held;
1656  }
1657  };
1658 
1663  template <class T>
1665  {
1666  private:
1667 
1668  alignas(T) char m_buf[sizeof(T)];
1669  T * m_object = nullptr;
1670 
1671  public:
1672 
1673  maybe_allocate(const bool allocate, const T & t)
1674  {
1675  if (allocate)
1676  {
1677  m_object = new (m_buf) T(t);
1678  }
1679  }
1680 
1681  maybe_allocate(const bool allocate, T && t)
1682  {
1683  if (allocate)
1684  {
1685  m_object = new (m_buf) T(t);
1686  }
1687  }
1688 
1689  template <class ... Args>
1690  maybe_allocate(const bool allocate, Args && ... args)
1691  {
1692  if (allocate)
1693  {
1694  m_object = new (m_buf) T(std::forward<Args>(args)...);
1695  }
1696  }
1697 
1699  {
1700  }
1701 
1702 
1704  {
1705  }
1706 
1708  {
1709  if (&other != this)
1710  {
1711  if (m_object != nullptr)
1712  {
1713  (*m_object) = other.get();
1714  }
1715  else
1716  {
1717  m_object = new (m_buf) T(other.get());
1718  }
1719  }
1720  return (*this);
1721  }
1722 
1723 
1725  {
1726  if (&other != this)
1727  {
1728  if (m_object != nullptr)
1729  {
1730  (*m_object) = other.get();
1731  }
1732  else
1733  {
1734  m_object = new (m_buf) T(other.get());
1735  }
1736  }
1737  return (*this);
1738  }
1739 
1741  {
1742  if (m_object != nullptr)
1743  {
1744  m_object->~T();
1745  }
1746  }
1747 
1748  bool valid() const
1749  {
1750  return m_object != nullptr;
1751  }
1752 
1753  T && get() &&
1754  {
1755  return *m_object;
1756  }
1757 
1758  T & get() &
1759  {
1760  return *m_object;
1761  }
1762 
1763  const T & get() const &
1764  {
1765  return *m_object;
1766  }
1767 
1768  const T * operator ->() const
1769  {
1770  return m_object;
1771  }
1772 
1774  {
1775  return m_object;
1776  }
1777 
1778  operator T & ()
1779  {
1780  return *m_object;
1781  }
1782 
1783  operator T && () &&
1784  {
1785  return *m_object;
1786  }
1787 
1788  operator const T & () const
1789  {
1790  return *m_object;
1791  }
1792  };
1793  }
1794 
1795 }
1796 
1797 #endif // CALORECGPU_HELPERS_H
CaloRecGPU::Helpers::maybe_allocate::maybe_allocate
maybe_allocate(maybe_allocate &&other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1703
AllowedVariables::e
e
Definition: AsgElectronSelectorTool.cxx:37
CaloRecGPU::CUDA_Helpers::CPU_to_GPU_async
void CPU_to_GPU_async(void *dest, const void *const source, const size_t num, CUDAStreamPtrHolder stream={})
Copies num bytes from source in CPU memory to dest in GPU memory, asynchronously.
temp
Definition: JetEventDict.h:21
CaloRecGPU::Helpers::maybe_allocate::maybe_allocate
maybe_allocate(const bool allocate, T &&t)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1681
CaloRecGPU::Helpers::maybe_allocate::operator=
maybe_allocate & operator=(const maybe_allocate &other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1707
CaloRecGPU::Helpers::MemoryManagement
! Handles allocation of a type T, using indexer as the integer type to indicate sizes.
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:477
CaloRecGPU::CUDA_Helpers::optimize_block_and_grid_size
void optimize_block_and_grid_size(void *func, int &block_size, int &grid_size, const int dynamic_memory=0, const int block_size_limit=0)
Optimizes block and grid size according to cudaOccupancyMaxPotentialBlockSize.
CaloRecGPU::CUDA_Helpers::supports_cooperative_launches
bool supports_cooperative_launches()
CaloRecGPU::Helpers::MemoryManagement::copy_helper< MemoryContext::CUDAGPU, MemoryContext::CPU, dummy >::copy
static void copy(T *dest, const T *const source, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:550
CaloRecGPU::Helpers::MemoryManagement::copy_helper< MemoryContext::CPU, MemoryContext::CUDAPinnedCPU, dummy >::copy
static void copy(T *dest, const T *const source, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:566
checkFileSG.line
line
Definition: checkFileSG.py:75
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::SimpleHolder
SimpleHolder(SimpleHolder &&other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1186
CaloRecGPU::CUDA_Helpers::GPU_to_CPU_async
void GPU_to_CPU_async(void *dest, const void *const source, const size_t num, CUDAStreamPtrHolder stream={})
Copies num bytes from source in GPU memory to dest in CPU memory, asynchronously.
python.SystemOfUnits.s
int s
Definition: SystemOfUnits.py:131
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::SimpleContainer
SimpleContainer(const SimpleContainer< T, other_indexer, other_context, other_hold > &other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:782
xAOD::short
short
Definition: Vertex_v1.cxx:165
fitman.sz
sz
Definition: fitman.py:527
CaloRecGPU::Helpers::MemoryManagement::move_helper
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:597
python.SystemOfUnits.m
int m
Definition: SystemOfUnits.py:91
CaloRecGPU::Helpers::MemoryManagement::unary_helper< MemoryContext::CUDAPinnedCPU, dummy >::allocate
static T * allocate(const indexer size)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:511
CaloRecGPU::Helpers::Constants::pi
constexpr T pi
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:316
CaloRecGPU::Helpers::MemoryManagement::unary_helper
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:479
CaloRecGPU::CUDA_Helpers::optimize_block_and_grid_size_for_cooperative_launch
void optimize_block_and_grid_size_for_cooperative_launch(void *func, int &block_size, int &grid_size, const int dynamic_memory=0, const int block_size_limit=0)
Optimizes block and grid size for a cooperative launch.
CaloCellPos2Ntuple.int
int
Definition: CaloCellPos2Ntuple.py:24
CaloRecGPU::Helpers::MemoryManagement::deallocate
static void deallocate(T *&arr)
!
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:634
CaloRecGPU::CUDA_Helpers::CPU_to_GPU
void CPU_to_GPU(void *dest, const void *const source, const size_t num)
Copies num bytes from source in CPU memory to dest in GPU memory.
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::valid
CUDA_HOS_DEV bool valid() const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1104
CaloRecGPU::Helpers::maybe_allocate::maybe_allocate
maybe_allocate(const maybe_allocate &other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1698
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::binary_output
void binary_output(stream &s) const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1322
CaloRecGPU::Helpers::separate_thread_holder::resize
void resize(const size_t new_size)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1548
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::textual_output
void textual_output(stream &s, const str &separator=" ") const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:865
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::SimpleHolder
SimpleHolder(const bool really_allocate)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1126
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::~SimpleContainer
~SimpleContainer()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:848
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::clear
void clear()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1109
CaloRecGPU::Helpers::separate_thread_holder::add_one_and_return
T & add_one_and_return()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1490
xAOD::char
char
Definition: TrigDecision_v1.cxx:38
CaloRecGPU::Helpers::MemoryManagement::allocate
static T * allocate(const indexer size)
!
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:620
DMTest::C
C_v1 C
Definition: C.h:26
CaloRecGPU::Helpers::SimpleHolder
Holds one objects of type \T in memory context Context.
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1067
CaloRecGPU::Helpers::Pearson_hash_16_bit
constexpr unsigned short Pearson_hash_16_bit(const T number)
Calculates a 16-bit Pearson hash from @ number.
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:281
CaloRecGPU::Helpers::MemoryManagement::unary_helper< MemoryContext::CUDAGPU, dummy >::allocate
static T * allocate(const indexer size)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:497
CaloRecGPU::Helpers::int_ceil_div
constexpr int int_ceil_div(const int num, const int denom)
Returns the ceiling of num/denom, with proper rounding.
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:214
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::allocate
void allocate()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1114
CaloRecGPU::Helpers::separate_thread_accessor::separate_thread_accessor
separate_thread_accessor(separate_thread_holder< T > &s)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1624
taskman.template
dictionary template
Definition: taskman.py:317
CaloRecGPU::Helpers::SimpleHolder< T, Context, false >::valid
CUDA_HOS_DEV bool valid() const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1392
CaloRecGPU::CUDA_Helpers::GPU_synchronize
void GPU_synchronize(CUDAStreamPtrHolder stream={})
Synchronizes the stream.
CaloRecGPU::CUDA_Helpers::deallocate_pinned
void deallocate_pinned(void *address)
Deallocates address in CPU pinned memory.
CaloRecGPU::Helpers::SimpleHolder< T, Context, false >::indexer
unsigned int indexer
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1362
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::SimpleHolder
SimpleHolder(const X &other_v)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1152
CaloRecGPU::Helpers::separate_thread_holder::m_held
std::vector< std::unique_ptr< T > > m_held
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1475
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::SimpleContainer
SimpleContainer(SimpleContainer &&other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:774
mc.diff
diff
Definition: mc.SFGenPy8_MuMu_DD.py:14
CaloRecGPU::Helpers::maybe_allocate::operator->
const T * operator->() const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1768
CaloRecGPU::Helpers::separate_thread_holder::operate_on_all
void operate_on_all(F &&f, Args &&... args)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1569
CaloRecGPU::Helpers::MemoryContext::CUDAGPU
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:465
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, false >::SimpleContainer
CUDA_HOS_DEV SimpleContainer(T *other_array, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:998
athena.value
value
Definition: athena.py:124
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::textual_input
void textual_input(stream &s)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:884
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::SimpleContainer
SimpleContainer(SimpleContainer< T, other_indexer, other_context, true > &&other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:791
perfmonmt-printer.dest
dest
Definition: perfmonmt-printer.py:189
CaloRecGPU::Helpers::MemoryManagement::unary_helper< MemoryContext::CPU, dummy >::deallocate
static void deallocate(T *&arr)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:488
CaloRecGPU::Helpers::MemoryManagement::unary_helper< MemoryContext::CUDAPinnedCPU, dummy >::deallocate
static void deallocate(T *&arr)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:516
read_hist_ntuple.t
t
Definition: read_hist_ntuple.py:5
CaloRecGPU::Helpers::MemoryManagement::unary_helper< MemoryContext::CPU, dummy >::allocate
static T * allocate(const indexer size)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:483
dbg::ptr
void * ptr(T *p)
Definition: SGImplSvc.cxx:74
const
bool const RAWDATA *ch2 const
Definition: LArRodBlockPhysicsV0.cxx:560
CaloRecGPU::Helpers::MemoryManagement::unary_helper< MemoryContext::CUDAGPU, dummy >::deallocate
static void deallocate(T *&arr)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:502
Args
Definition: test_lwtnn_fastgraph.cxx:12
fitman.rho2
rho2
Definition: fitman.py:544
CaloRecGPU::Helpers::maybe_allocate::maybe_allocate
maybe_allocate(const bool allocate, Args &&... args)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1690
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::SimpleHolder
SimpleHolder()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1122
drawFromPickle.exp
exp
Definition: drawFromPickle.py:36
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::SimpleContainer
SimpleContainer(const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:754
x
#define x
CaloRecGPU::Helpers::MemoryManagement::copy_helper< MemoryContext::CUDAGPU, MemoryContext::CUDAGPU, dummy >::copy
static void copy(T *dest, const T *const source, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:542
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::SimpleHolder
SimpleHolder(SimpleHolder< X, other_context, true > &&other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1194
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::indexer
unsigned int indexer
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1074
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::binary_input
void binary_input(stream &s)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:931
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::binary_output
void binary_output(stream &s) const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:913
CaloRecGPU::Helpers::int_floor_div
constexpr int int_floor_div(const int num, const int denom)
Returns the floor of num/denom, with proper rounding.
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:220
AthenaPoolTestWrite.stream
string stream
Definition: AthenaPoolTestWrite.py:12
XMLtoHeader.count
count
Definition: XMLtoHeader.py:85
makeTRTBarrelCans.y1
tuple y1
Definition: makeTRTBarrelCans.py:15
Monitored::X
@ X
Definition: HistogramFillerUtils.h:24
CaloRecGPU::Helpers::MemoryContext::CUDAPinnedCPU
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:469
CaloRecGPU::CUDA_Helpers::deallocate
void deallocate(void *address)
Deallocates address in GPU memory.
CaloRecGPU::Helpers::MemoryManagement::move_helper< C, C, dummy >::move
static void move(T *&dest, T *&source, const indexer)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:611
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::textual_input
void textual_input(stream &s)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1291
CaloRecGPU::Helpers::Pearson_hash
constexpr unsigned char Pearson_hash(const T number)
Calculates a Pearson hash from @ number.
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:254
CaloRecGPU::Helpers::MemoryContext::CUDAGPU::name
constexpr static char const * name
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:466
python.setupRTTAlg.size
int size
Definition: setupRTTAlg.py:39
calibdata.valid
list valid
Definition: calibdata.py:45
CaloRecGPU::Helpers::separate_thread_holder::filled_size
size_t filled_size() const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1599
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::SimpleHolder
SimpleHolder(X *other_p)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1142
histSizes.code
code
Definition: histSizes.py:129
InDetDD::operator*
SiLocalPosition operator*(const SiLocalPosition &position, const double factor)
Definition: SiLocalPosition.cxx:98
python.utils.AtlRunQueryDQUtils.p
p
Definition: AtlRunQueryDQUtils.py:210
CaloRecGPU::Helpers::MemoryManagement::copy_helper< MemoryContext::CPU, MemoryContext::CUDAGPU, dummy >::copy
static void copy(T *dest, const T *const source, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:534
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::SimpleContainer
SimpleContainer()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:750
CaloRecGPU::Helpers::separate_thread_holder::m_mutex
std::shared_mutex m_mutex
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1488
CUDA_HOS_DEV
#define CUDA_HOS_DEV
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:101
CaloRecGPU::CUDA_Helpers::GPU_name
std::string GPU_name()
CaloRecGPU::Helpers::MemoryContext::CUDAPinnedCPU::name
constexpr static char const * name
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:470
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::size
CUDA_HOS_DEV indexer size() const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:713
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::textual_output
void textual_output(stream &s, const str &separator=" ") const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1270
lumiFormat.i
int i
Definition: lumiFormat.py:85
z
#define z
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::~SimpleHolder
~SimpleHolder()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1252
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::binary_input
void binary_input(stream &s)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1340
CaloRecGPU::Helpers::separate_thread_holder::m_thread_equivs
std::vector< typename std::thread::id > m_thread_equivs
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1476
Base
CaloRecGPU::CUDA_Helpers::GPU_to_GPU
void GPU_to_GPU(void *dest, const void *const source, const size_t num)
Copies num bytes from source to dest, both in GPU memory.
CaloRecGPU::Helpers::MemoryContext::CPU::name
constexpr static char const * name
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:462
file
TFile * file
Definition: tile_monitor.h:29
CaloRecGPU::CUDA_Helpers::GPU_to_CPU
void GPU_to_CPU(void *dest, const void *const source, const size_t num)
Copies num bytse from source in GPU memory to dest in CPU memory.
CaloRecGPU::Helpers::maybe_allocate::valid
bool valid() const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1748
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::clear
void clear()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:728
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::SimpleHolder
SimpleHolder(const SimpleHolder< X, other_context, other_hold > &other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1173
python.xAODType.dummy
dummy
Definition: xAODType.py:4
CaloRecGPU::CUDA_Helpers::supports_dynamic_parallelism
bool supports_dynamic_parallelism()
hist_file_dump.f
f
Definition: hist_file_dump.py:135
CaloRecGPU::Helpers::SimpleHolder< T, Context, false >::SimpleHolder
CUDA_HOS_DEV SimpleHolder()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1399
CaloRecGPU::Helpers::separate_thread_holder::get_one
T & get_one()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1499
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::m_object
T * m_object
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1076
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, false >::SimpleContainer
CUDA_HOS_DEV SimpleContainer(const SimpleContainer< T, other_indexer, Context, other_hold > &other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1005
CaloRecGPU::Helpers::MemoryManagement::copy_helper< MemoryContext::CUDAPinnedCPU, MemoryContext::CUDAPinnedCPU, dummy >::copy
static void copy(T *dest, const T *const source, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:574
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::m_array
T * m_array
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:703
CaloRecGPU::Helpers::SimpleHolder< T, Context, false >::m_object
T * m_object
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1364
CaloRecGPU::Helpers::SimpleContainer
Holds a run-time amount of objects of type \T, measuring sizes with indexer, in memory context Contex...
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:698
calibdata.exit
exit
Definition: calibdata.py:236
compute_lumi.denom
denom
Definition: compute_lumi.py:76
CaloRecGPU::Helpers::MemoryManagement::copy_helper< MemoryContext::CPU, MemoryContext::CPU, dummy >::copy
static void copy(T *dest, const T *const source, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:526
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::SimpleContainer
SimpleContainer(T *other_array, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:762
CaloRecGPU::Helpers::MemoryContext::CPU
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:461
CaloRecGPU::Helpers::MemoryManagement::copy_helper< MemoryContext::CUDAPinnedCPU, MemoryContext::CUDAGPU, dummy >::copy
static void copy(T *dest, const T *const source, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:582
trigbs_pickEvents.num
num
Definition: trigbs_pickEvents.py:76
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, false >::size
CUDA_HOS_DEV indexer size() const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:972
MuonR4::SegmentFit::ParamDefs::y0
@ y0
CaloRecGPU::CUDA_Helpers::allocate
void * allocate(const size_t num)
Allocates and returns the address of num bytes from GPU memory.
CaloRecGPU::Helpers::maybe_allocate::m_object
T * m_object
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1669
CaloRecGPU::CUDA_Helpers::allocate_pinned
void * allocate_pinned(const size_t num)
Allocates and returns the address of num bytes from CPU pinned memory.
CaloRecGPU::Helpers::Constants::sqrt2
constexpr T sqrt2
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:319
CaloRecGPU::Helpers::MemoryManagement::copy_helper< MemoryContext::CUDAPinnedCPU, MemoryContext::CPU, dummy >::copy
static void copy(T *dest, const T *const source, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:558
id
SG::auxid_t id
Definition: Control/AthContainers/Root/debug.cxx:227
python.selection.number
number
Definition: selection.py:20
name
std::string name
Definition: Control/AthContainers/Root/debug.cxx:228
plotBeamSpotMon.b
b
Definition: plotBeamSpotMon.py:77
CaloRecGPU::Helpers::separate_thread_holder
Manages objects of type T in a thread-safe way, ensuring that there's an object available for each se...
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1473
CaloRecGPU::Helpers::separate_thread_holder::get_for_thread
T & get_for_thread() const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1518
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, false >::m_array
T * m_array
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:962
CaloRecGPU::Helpers::maybe_allocate::get
const T & get() const &
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1763
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, false >::SimpleContainer
CUDA_HOS_DEV SimpleContainer()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:989
CaloRecGPU::Helpers::MemoryManagement::move
static void move(T *&dest, T *&source, const indexer sz)
!
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:670
CaloRecGPU::Helpers::separate_thread_accessor::separate_thread_accessor
separate_thread_accessor(separate_thread_holder< T > &s, T *&ptr)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1651
CaloRecGPU::CUDA_Helpers::CUDAStreamPtrHolder
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:110
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::m_size
indexer m_size
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:705
RTTAlgmain.address
address
Definition: RTTAlgmain.py:55
CaloRecGPU::CUDA_Helpers::GPU_to_GPU_async
void GPU_to_GPU_async(void *dest, const void *const source, const size_t num, CUDAStreamPtrHolder stream={})
Copies num bytes from source to dest, both in GPU memory, asynchronously.
CaloRecGPU::Helpers::SimpleHolder< T, Context, false >::SimpleHolder
CUDA_HOS_DEV SimpleHolder(const SimpleHolder< X, Context, other_hold > &other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1418
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::SimpleContainer
SimpleContainer(const SimpleContainer &other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:768
CaloRecGPU::Helpers::SimpleHolder< T, Context, true >::SimpleHolder
SimpleHolder(const SimpleHolder &other)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1158
CaloRecGPU::Helpers::MemoryManagement::copy
static void copy(T *dest, const T *const source, const indexer sz)
!
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:652
CaloRecGPU::Helpers::maybe_allocate::maybe_allocate
maybe_allocate(const bool allocate, const T &t)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1673
CaloRecGPU::Helpers::maybe_allocate
Possibly holds an object in its internal buffer.
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1665
a
TList * a
Definition: liststreamerinfos.cxx:10
InDetDD::other
@ other
Definition: InDetDD_Defs.h:16
VKalVrtAthena::varHolder_detail::clear
void clear(T &var)
Definition: NtupleVars.h:48
y
#define y
CaloRecGPU::CUDA_Helpers::CUDAStreamPtrHolder::CUDAStreamPtrHolder
CUDAStreamPtrHolder()=default
CaloRecGPU::Helpers::separate_thread_holder::release_one
void release_one()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1534
CaloRecGPU::Helpers::maybe_allocate::m_buf
char m_buf[sizeof(T)]
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1668
python.CaloScaleNoiseConfig.type
type
Definition: CaloScaleNoiseConfig.py:78
F
#define F(x, y, z)
Definition: MD5.cxx:112
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, false >::m_size
indexer m_size
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:964
python.CaloCondTools.log
log
Definition: CaloCondTools.py:20
ExpressionParsing::ast::operand
boost::variant< nil, double, unsigned int, bool, std::string, boost::recursive_wrapper< unaryexpr_ >, boost::recursive_wrapper< expression > > operand
Definition: ParsingInternals.h:43
CaloRecGPU::Helpers::compile_time_pow2
constexpr Base compile_time_pow2(const Exp exp)
Returns 2 to the power of exp.
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:227
CaloRecGPU::Helpers::MemoryManagement::copy_helper
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:522
CaloRecGPU::Helpers::MemoryManagement::copy_helper< MemoryContext::CUDAGPU, MemoryContext::CUDAPinnedCPU, dummy >::copy
static void copy(T *dest, const T *const source, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:590
CaloRecGPU::Helpers::separate_thread_accessor::release_one
void release_one()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1636
copySelective.source
string source
Definition: copySelective.py:32
str
Definition: BTagTrackIpAccessor.cxx:11
calibdata.copy
bool copy
Definition: calibdata.py:27
CaloRecGPU::Helpers::MemoryManagement::move_helper::move
static void move(T *&dest, T *&source, const indexer sz)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:601
CaloRecGPU::CUDA_Helpers::CUDAStreamPtrHolder::CUDAStreamPtrHolder
CUDAStreamPtrHolder(T *p)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:125
CaloRecGPU::CUDA_Helpers::CUDAStreamPtrHolder::ptr
void * ptr
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:111
CaloRecGPU::Helpers::Constants::inv_sqrt2
constexpr T inv_sqrt2
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:323
CaloRecGPU::Helpers::separate_thread_accessor::m_held
T * m_held
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1622
python.PyAthena.obj
obj
Definition: PyAthena.py:132
CaloRecGPU
Definition: BaseDefinitions.h:11
CaloRecGPU::Helpers::separate_thread_holder::available_size
size_t available_size() const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1584
xAOD::bool
setBGCode setTAP setLVL2ErrorBits bool
Definition: TrigDecision_v1.cxx:60
CaloRecGPU::Helpers::separate_thread_accessor::m_sth
separate_thread_holder< T > & m_sth
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1621
python.CaloScaleNoiseConfig.args
args
Definition: CaloScaleNoiseConfig.py:80
CaloRecGPU::Helpers::separate_thread_accessor::get_one
T & get_one()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1628
CaloRecGPU::Helpers::separate_thread_accessor::~separate_thread_accessor
~separate_thread_accessor()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1644
TSU::T
unsigned long long T
Definition: L1TopoDataTypes.h:35
CaloRecGPU::Helpers::SimpleHolder< T, Context, false >::SimpleHolder
CUDA_HOS_DEV SimpleHolder(X *other_p)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1409
CaloRecGPU::Helpers::maybe_allocate::get
T && get() &&
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1753
CaloRecGPU::Helpers::separate_thread_accessor
!
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1619
CaloRecGPU::Helpers::maybe_allocate::~maybe_allocate
~maybe_allocate()
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1740
CaloRecGPU::Helpers::SimpleContainer< T, indexer, Context, true >::resize
void resize(const indexer new_size)
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:734
CaloRecGPU::Helpers::separate_thread_holder::held_size
size_t held_size() const
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1578
CaloRecGPU::Helpers::maybe_allocate::get
T & get() &
Definition: Calorimeter/CaloRecGPU/CaloRecGPU/Helpers.h:1758