#ifndef BUNCH_CTX_INTERNAL_H_ #define BUNCH_CTX_INTERNAL_H_ /** @file bunch-ctx-internal.h Internal includes for bunch-ctx*.cu files */ #define USE_SHARED_TUBES 0 // tubes loaded per one shared memory iteration #define SHARED_NTUBES 64 // hits loaded into shared memory per iteration; only times are loaded #define SHARED_NHITS (4 * SHARED_NTUBES) // skewlets loaded into shared memory per iteration #define SHARED_NSKEWLETS 128 //#define SHARED_NSKEWLETS 256 // block size for eval_tracks_*_k kernel #if USE_SHARED_TUBES #define EVAL_TRACKS_BS 256 //TODO: MIN BLOCKS are not optimized for this branch #define EVAL_TRACKS_RASTER_MIN_BLOCKS 1 #define EVAL_TRACKS_SKEWLETS_MIN_BLOCKS 1 #else //#define EVAL_TRACKS_BS 192 #define EVAL_TRACKS_BS 256 #define EVAL_TRACKS_RASTER_MIN_BLOCKS 6 #define EVAL_TRACKS_SKEWLETS_MIN_BLOCKS 8 #endif // block size for half-skewlet pairing #define SKEWLETS_BSX 128 #define SKEWLETS_BSY 1 //#define SKEWLETS_BSX 64 //#define SKEWLETS_BSY 4 #define MAX_PHIS \ {{0, F_PI / 6, F_PI / 2}, {0, 0, F_PI / 180 * 35}, {0, 0, 0}} #define MIN_COS_PHIS \ {{0, (float)cos(F_PI / 6), (float)cos(F_PI / 2)}, {0, 0, (float)cos(F_PI / \ 180 * 35)}, {0, 0, 0}} // functor classes for using with Thrust-like libraries /** functor class for checking whether the hit is pivot */ class hit_is_pivot { layptr tubes; public: __host__ __device__ hit_is_pivot(layptr tubes) { this->tubes = tubes; } __host__ __device__ bool operator()(const Hit &h) const { return tubes[h.tube_id].pivot_range_id != -1; } }; // hit_is_pivot /** functor class for checking whether the hit is skewed */ class hit_is_skewed { layptr tubes; public: __host__ __device__ hit_is_skewed(layptr tubes) { this->tubes = tubes; } __host__ __device__ bool operator()(const Hit &h) const { return tubes[h.tube_id].skewed_range_id != -1; } }; // hit_is_skewed /** this functor checks whether the track is good */ class is_track_good { int min_nhits0, min_nhits1, row_diff; layptr tubes; public: __host__ __device__ is_track_good (int min_nhits0, int min_nhits1, int row_diff, layptr tubes) { this->min_nhits0 = min_nhits0; this->min_nhits1 = min_nhits1; this->row_diff = row_diff; this->tubes = tubes; } __host__ __device__ bool operator()(const Track &track) const { int track_row_diff = tubes[track.inner_tube_id].row - tubes[track.outer_tube_id].row; track_row_diff = track_row_diff < 0 ? -track_row_diff : track_row_diff; int min_nhits = track_row_diff < row_diff ? min_nhits0 : min_nhits1; return track.nhits >= min_nhits; } }; // functor is_track_good // kernel declarations; required as they are now referenced in many compilation // units __global__ void tube_hit_histo_k(BunchCtx *b); __global__ void tube_hit_group_k(BunchCtx *b); __global__ void nz_tubes_k(BunchCtx *b); __global__ void find_triplets_k(BunchCtx *b); __global__ void find_half_skewlets_k(BunchCtx *b); __global__ void combine_half_skewlets_k(BunchCtx *b, int irange); __global__ void poca_test_k(BunchCtx *b, int irange); __global__ void compute_circles_k (layptr cand_tracks, int ncand_tracks); __global__ void eval_tracks_tubes_raster_k(BunchCtx *b); __global__ void eval_tracks_tubes_shared_k(BunchCtx *b); __global__ void eval_tracks_skewlets_all_k(BunchCtx *b); __global__ void eval_tracks_skewlets_bins_k(BunchCtx *b); __global__ void combine_triplets_k (BunchCtx *b, int inner_id, int outer_id, float max_phi); #endif