CandsGPU::BestCands< MaxCandsPerSeed, BlockSize > Struct Template Reference

#include <BestCands.h>

Public Member Functions
__device__ bool	better (int icand_fst, int fst, int icand_snd, int snd)
__device__ void	copy_node (int icand_fst, int fst, int icand_snd, int snd)
__device__ int	count_valid_cands (int itrack)
__device__ void	get_cand_info (const int tid, const int cid, int &my_trkIdx, int &my_hitIdx, int &my_nhits, float &my_chi2)
__device__ int	get_nhits (const int tid, const int cid)
__device__ void	heap_sort (int icand, int heap_size)
__device__ void	heapify (int itrack, int idx, int heap_size)
__device__ int	left (int idx)
__device__ void	merge_cands_for_seed (int iseed, int icand)
__device__ void	print_heap (const int tid)
__device__ void	reset (int itrack)
__device__ int	right (int idx)
__device__ void	swap_nodes (int icand_fst, int fst, int icand_snd, int snd)
__device__ void	update (int itrack, int cand_trIdx, int cand_hitIdx, int cand_nhits, float cand_chi2)

Public Attributes
float	chi2 [MaxCandsPerSeed][BlockSize]
int	hitIdx [MaxCandsPerSeed][BlockSize]
int	nhits [MaxCandsPerSeed][BlockSize]
int	trkIdx [MaxCandsPerSeed][BlockSize]

Detailed Description

template<int MaxCandsPerSeed, int BlockSize>
struct CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >

Definition at line 24 of file BestCands.h.

Member Function Documentation

better()

template<int M, int B>

__device__ bool CandsGPU::BestCands< M, B >::better	(	int	icand_fst,
		int	fst,
		int	icand_snd,
		int	snd
	)

Definition at line 101 of file BestCands.h.

References nano_mu_local_reco_cff::chi2, and nhits.

                                                                                        {
    return (nhits[fst][icand_fst] > nhits[snd][icand_snd]) ||
           ((nhits[fst][icand_fst] == nhits[snd][icand_snd]) && (chi2[fst][icand_fst] < chi2[snd][icand_snd]));
  }

copy_node()

template<int M, int B>

__device__ void CandsGPU::BestCands< M, B >::copy_node	(	int	icand_fst,
		int	fst,
		int	icand_snd,
		int	snd
	)

Definition at line 115 of file BestCands.h.

References nano_mu_local_reco_cff::chi2, and nhits.

                                                                                           {
    trkIdx[snd][icand_snd] = trkIdx[fst][icand_fst];
    hitIdx[snd][icand_snd] = hitIdx[fst][icand_fst];
    nhits[snd][icand_snd] = nhits[fst][icand_fst];
    chi2[snd][icand_snd] = chi2[fst][icand_fst];
  }

count_valid_cands()

template<int M, int B>

__device__ int CandsGPU::BestCands< M, B >::count_valid_cands

(

int

itrack

)

Definition at line 222 of file BestCands.h.

References submitPVResolutionJobs::count, mps_fire::i, and CandsGPU::trkIdx_sentinel.

                                                              {
    int count = 0;
    for (int i = 0; i < M; ++i) {
      if (trkIdx[i][itrack] != trkIdx_sentinel)
        ++count;
    }
    return count;
  }

get_cand_info()

template<int MaxCandsPerSeed, int BlockSize>

__device__ void CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >::get_cand_info	(	const int	tid,
		const int	cid,
		int &	my_trkIdx,
		int &	my_hitIdx,
		int &	my_nhits,
		float &	my_chi2
	)

Definition at line 211 of file BestCands.h.

References ecaldqm::BlockSize, nano_mu_local_reco_cff::chi2, and nhits.

                                                                                                   {
    if (cid < MaxCandsPerSeed && tid < BlockSize) {
      my_trkIdx = trkIdx[cid][tid];
      my_hitIdx = hitIdx[cid][tid];
      my_nhits = nhits[cid][tid];
      my_chi2 = chi2[cid][tid];
    }
  }

get_nhits()

template<int MaxCandsPerSeed, int BlockSize>

__device__ int CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >::get_nhits ( const int  tid, const int  cid  )

inline

Definition at line 49 of file BestCands.h.

References CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >::nhits.

{ return nhits[cid][tid]; }

heap_sort()

template<int M, int B>

__device__ void CandsGPU::BestCands< M, B >::heap_sort	(	int	icand,
		int	heap_size
	)

Definition at line 201 of file BestCands.h.

References mps_fire::i.

                                                                     {
    int num_unsorted_elts = heap_size;
    // Assume that we have a heap with the worst one at the root.
    for (int i = heap_size - 1; i > 0; --i) {
      swap_nodes(icand, 0, icand, i);  // worst at the end
      heapify(icand, 0, --num_unsorted_elts);
    }
  }

heapify()

template<int M, int B>

__device__ void CandsGPU::BestCands< M, B >::heapify	(	int	itrack,
		int	idx,
		int	heap_size
	)

Definition at line 123 of file BestCands.h.

References heavyIonCSV_trainingSettings::idx.

                                                                            {
    // We want to move idx down so the smallest value is at the root
    int smallest = -1;
    while (idx != smallest) {
      if (idx < 0 || idx >= heap_size / 2)
        return;

      smallest = idx;
      if (heap_size > left(idx) && better(icand, smallest, icand, left(idx)))
        smallest = left(idx);
      if (heap_size > right(idx) && better(icand, smallest, icand, right(idx)))
        smallest = right(idx);

      if (smallest != idx) {
        swap_nodes(icand, smallest, icand, idx);
        idx = smallest;
        smallest = -1;
      }
    }
  }

left()

template<int M, int B>

__device__ int CandsGPU::BestCands< M, B >::left

(

int

idx

)

Definition at line 91 of file BestCands.h.

References heavyIonCSV_trainingSettings::idx.

Referenced by svgfig.Curve.Samples::__len__().

                                              {
    return (++idx << 1) - 1;
  }

merge_cands_for_seed()

template<int M, int B>

__device__ void CandsGPU::BestCands< M, B >::merge_cands_for_seed	(	int	iseed,
		int	icand
	)

Definition at line 145 of file BestCands.h.

References cms::cudacompat::__syncthreads(), B, mps_fire::i, iseed, and dqmiolumiharvest::j.

                                                                            {
    int itrack = iseed * M + icand;
// TODO: Need a better way to reduce candidates.
//       So far, binary tree reduction is a bit slower than the naive approach
#if 1
    if (icand) {
      heap_sort(itrack, M);
    }
    __syncthreads();  // cand 0 waits;
    if (icand)
      return;  // reduction by the first cand of each seed

    for (int i = itrack + 1; i < itrack + M; ++i) {  // over cands
      for (int j = 0; j < M; ++j) {                  // inside heap
        if (better(i, j, itrack, 0)) {
          copy_node(i, j, itrack, 0);
          heapify(itrack, 0, M);
        } else {
          break;
        }
      }
    }
    heap_sort(itrack, M);
    __syncthreads();  // TODO: Volta: sync only on MaxCandsPerSeeds threads
#else

    for (int step = 2; step <= Config::maxCandsPerSeed; step <<= 1) {
      if (icand % step == step / 2) {
        heap_sort(itrack, M);
      }
      __syncthreads();

      if (icand % step == 0) {
        int i = itrack + step / 2;
        if ((i < iseed * M + M) && (i < B) && (icand + step / 2 < M)) {
          for (int j = 0; j < M; ++j) {  // inside heap
            if (better(i, j, itrack, 0)) {
              copy_node(i, j, itrack, 0);
              heapify(itrack, 0, M);
            } else {
              break;
            }
          }
        }
      }
      //__syncthreads();
    }

    if (icand == 0) {
      heap_sort(itrack, M);
    }
    __syncthreads();
#endif
  }

print_heap()

template<int M, int B>

__device__ void CandsGPU::BestCands< M, B >::print_heap

(

const int

tid

)

Definition at line 79 of file BestCands.h.

References nano_mu_local_reco_cff::chi2, and nhits.

                                                           {
    for (int cid = 0; cid < M; cid++) {
      printf(">>>>> tid %d rowIdx %d hitIdx %d nhits %d chi2 %f\n",
             tid,
             cid,
             hitIdx[cid][tid],
             nhits[cid][tid],
             chi2[cid][tid]);
    }
  }

reset()

template<int M, int B>

__device__ void CandsGPU::BestCands< M, B >::reset

(

int

itrack

)

Definition at line 55 of file BestCands.h.

References nano_mu_local_reco_cff::chi2, CandsGPU::chi2_sentinel, CandsGPU::hitIdx_sentinel, dqmiolumiharvest::j, nhits, CandsGPU::nhits_sentinel, and CandsGPU::trkIdx_sentinel.

                                                   {
    for (auto j = 0; j < M; ++j) {
      trkIdx[j][itrack] = trkIdx_sentinel;
      hitIdx[j][itrack] = hitIdx_sentinel;
      nhits[j][itrack] = nhits_sentinel;
      chi2[j][itrack] = chi2_sentinel;
    }
  }

right()

template<int M, int B>

__device__ int CandsGPU::BestCands< M, B >::right

(

int

idx

)

Definition at line 96 of file BestCands.h.

References heavyIonCSV_trainingSettings::idx.

                                               {
    return ++idx << 1;
  }

swap_nodes()

template<int M, int B>

__device__ void CandsGPU::BestCands< M, B >::swap_nodes	(	int	icand_fst,
		int	fst,
		int	icand_snd,
		int	snd
	)

Definition at line 107 of file BestCands.h.

References nano_mu_local_reco_cff::chi2, nhits, and CandsGPU::swap_values().

                                                                                            {
    swap_values(trkIdx[fst][icand_fst], trkIdx[snd][icand_snd]);
    swap_values(hitIdx[fst][icand_fst], hitIdx[snd][icand_snd]);
    swap_values(nhits[fst][icand_fst], nhits[snd][icand_snd]);
    swap_values(chi2[fst][icand_fst], chi2[snd][icand_snd]);
  }

update()

template<int M, int B>

__device__ void CandsGPU::BestCands< M, B >::update	(	int	itrack,
		int	cand_trIdx,
		int	cand_hitIdx,
		int	cand_nhits,
		float	cand_chi2
	)

Definition at line 65 of file BestCands.h.

References nano_mu_local_reco_cff::chi2, and nhits.

Referenced by progressbar.ProgressBar::__next__(), MatrixUtil.Matrix::__setitem__(), MatrixUtil.Steps::__setitem__(), progressbar.ProgressBar::finish(), and MatrixUtil.Steps::overwrite().

                                                                                                                      {
    if (cand_nhits < nhits[0][itrack])
      return;
    if (cand_chi2 > chi2[0][itrack])
      return;
    trkIdx[0][itrack] = cand_trIdx;
    hitIdx[0][itrack] = cand_hitIdx;
    nhits[0][itrack] = cand_nhits;
    chi2[0][itrack] = cand_chi2;

    heapify(itrack, 0, M);
  }

Member Data Documentation

chi2

template<int MaxCandsPerSeed, int BlockSize>

float CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >::chi2[MaxCandsPerSeed][BlockSize]

Definition at line 29 of file BestCands.h.

hitIdx

template<int MaxCandsPerSeed, int BlockSize>

int CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >::hitIdx[MaxCandsPerSeed][BlockSize]

Definition at line 27 of file BestCands.h.

Referenced by ntupleDataFormat._RecoHitAdaptor::_hits().

nhits

template<int MaxCandsPerSeed, int BlockSize>

int CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >::nhits[MaxCandsPerSeed][BlockSize]

Definition at line 28 of file BestCands.h.

Referenced by CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >::get_nhits().

trkIdx

template<int MaxCandsPerSeed, int BlockSize>

int CandsGPU::BestCands< MaxCandsPerSeed, BlockSize >::trkIdx[MaxCandsPerSeed][BlockSize]

Definition at line 26 of file BestCands.h.