#ifndef TRIPLETS_GLOBAL_CTX_H_ #define TRIPLETS_GLOBAL_CTX_H_ /** @file global-ctx.h definition of the global context class */ #include using namespace std; #include "data.h" #include "layout-ptr.h" #include "options.h" #include "sector-layer.h" #include "track.h" #include "triplet.h" #include "util.h" /** the maximum number of neighbor tubes */ #define MAX_NNEIGHBORS 8 /** number of hits per bunch */ //#define BUNCH_NHITS 2048 #define HIT_RATE 8.0f // numbers are given as rates per ns; per-bunch numbers are then computed based // on given bunch duration /** maximum number of triplets (in future, per slice) */ //#define MAX_NTRIPLETS 4096 //#define MAX_BUNCH_NTRIPLETS 2048 #define TRIPLET_RATE HIT_RATE /** maximum number of skewlets and skewlet halfs (in future, per slice) */ //#define MAX_NSKEWLETS 16384 //#define MAX_BUNCH_NSKEWLETS 8192 #define SKEWLET_RATE (2 * HIT_RATE) /** maximum number of candidate tracks (in future, per slice) */ //#define MAX_NTRACKS 16384 //#define MAX_BUNCH_NTRACKS 8192 #define TRACK_RATE (2 * HIT_RATE) // forward definitions class BunchCtx; struct TripletFinderOptions; /** class representing the global context, i.e. one for all bunches */ class GlobalCtx { public: /** dummy default constructor */ __host__ __device__ GlobalCtx() {} /** initializing constructor for global context */ GlobalCtx(const TripletFinderOptions &opts); /** does all reconstruction work */ vector reconstruct(void); /** load the tubes from file */ void load_tubes(void); /** builds the straw map (apart from loading tubes) */ void build_straw_map(void); /** verifies some assumed properties of straw map (for each sector-layer) */ void verify_straw_map(void); /** builds the sector-layers to traverse the straw map quicker */ void build_sector_layers(void); /** builds skewlet bin layers to check against skewlets quicker */ void build_skewlet_bin_layers(void); /** loads the hits from file */ void load_hits(void); /** sets the maximum number of triplets, tracks etc. based on the number of hits */ void set_max_numbers(void); /** allocate (and reset) memory for all needs during the runtime of the algorithm */ void alloc_dev_memory(void); /** initialize bunches */ void init_bunches(void); /** find neighbors for each tube, also set pivot range id for each tube */ void find_neighbor_tubes(void); /** synchronizes bunches to device*/ void sync_bunches_to_device(void); /** bunches (i.e., groups into bunches) hits; one hit can land into 2 bunches */ void bunch_hits(void); /** bunch hits on host */ void bunch_hits_host(void); #if WITH_DYN_PAR /** initialize dynamic parallelism, so that DP launches do not incur additional overhead */ void dyn_init(void); /** calls per-bunch reconstruction with dynamic parallelism */ void reconstruct_dyn(void); #endif /** calls per-bunch reconstruction with single thread block per bunch */ void reconstruct_tblock(void); /** calls per-bunch reconstruction with 1 host stream per bunch */ void reconstruct_hstream(void); /** calls per-bunch reconstruction with 1 host thread / stream and multiple bunches per stream */ void reconstruct_hthread(void); /** collects tracks from all bunches into a single vector */ void collect_tracks(vector &tracks); /** device version of the global context */ GlobalCtx *d_this; /** triplet finder options */ TripletFinderOptions opts; /** tubes on host and device side, [ntubes] */ layptr tubes, d_tubes; /** the total number of tubes */ int ntubes; /** number of tube rows and sectors */ int nrows, nsectors; /** number of straight (non-skewed) rows */ //int nstraight_rows; /** the straw map (tube ids by row and sector) [ntubes] */ int *straw_map, *d_straw_map; /** starts of tube arrays per row and sector [nrows * nsectors] */ int *straw_map_starts, *d_straw_map_starts; /** tube counts for each specific row and sector [nrows * nsectors] */ int *straw_map_counts, *d_straw_map_counts; /** sector-layers (to traverse the straw map quicker), [nrows, nsectors] (was [nstraight_rows, nsectors] at some time) */ SectorLayer *sector_layers, *d_sector_layers; /** sector-layers for skewlet bins (to perform skewlet checks faster) [nskewlet_ranges * nsectors] */ SectorLayer *skewlet_bin_layers, *d_skewlet_bin_layers; /** the pivot tube ranges */ range *pivot_ranges, *d_pivot_ranges; /** the number of pivot ranges */ int npivot_ranges; /** the skewed ranges */ range *skewed_ranges, *d_skewed_ranges; /** the number of skewed ranges */ int nskewed_ranges; /** total number of skewlet ranges, nskewed_ranges - 1 */ int nskewlet_ranges; // maximum numbers; note that appropriate GPU-side numbers are written into // kernel memory, and thus shared between all GPUs /** maximum number of hits, per bunch */ int max_nhits; /** maximum number of skewlets, per bunch */ int max_nskewlets; /** maximum number of triplets, per bunch */ int max_ntriplets; /** maximum number of tracks, per bunch */ int max_ntracks; /** hits on host and device side, [nhits] */ Hit *hits, *d_hits; /** the total number of hits */ int nhits; /** neighbor tube id array, 2D, [MAX_NNEIGHBOR_TUBES, ntubes]*/ int *d_neighbor_tubes; /** the pitch value for tube neighbors */ int neighbor_pitch; /** tube neighbor count array, [ntubes], each element <= MAX_NNEIGHBOR_TUBES */ int *d_nneighbors; /** the list of bunch contexts */ BunchCtx *bunches, *d_bunches; /** the total number of bunches */ int nbunches; // bunch context can access all fields of the global context friend class BunchCtx; }; // class GlobalCtx #endif