28 const short int* __restrict__ gpu_paramNo_arr,
29 const unsigned int* __restrict__ gpu_nKnots_arr,
30 const float* __restrict__ gpu_coeff_many,
31 float* __restrict__ gpu_weights,
32 const cudaTextureObject_t __restrict__ text_coeff_x);
39 const float* __restrict__ gpu_coeffs_tf1,
40 const short int* __restrict__ gpu_paramNo_arr_tf1,
41 float* __restrict__ gpu_weights_tf1);
43#ifndef Weight_On_SplineBySpline_Basis
51 const float* __restrict__ gpu_weights,
52 const float* __restrict__ gpu_weights_tf1,
54 float* __restrict__ gpu_total_weights,
56 const cudaTextureObject_t __restrict__ text_nParamPerEvent,
57 const cudaTextureObject_t __restrict__ text_nParamPerEvent_TF1);
76 #ifndef Weight_On_SplineBySpline_Basis
77 float **cpu_total_weights,
80 unsigned int total_nknots,
81 unsigned int n_splines,
106 std::vector<float> cpu_many_array_TF1,
107 std::vector<short int> cpu_paramNo_arr_TF1,
108 #ifndef Weight_On_SplineBySpline_Basis
110 std::vector<unsigned int> cpu_nParamPerEvent,
112 std::vector<unsigned int> cpu_nParamPerEvent_TF1,
115 unsigned int n_splines,
116 short int spline_size,
117 unsigned int total_nknots,
152 #ifdef Weight_On_SplineBySpline_Basis
154 float* cpu_weights_tf1,
156 float* cpu_total_weights,
162 const unsigned int h_n_splines,
163 const unsigned int h_n_tf1);
172 #ifndef Weight_On_SplineBySpline_Basis
173 float *cpu_total_weights
225 #ifndef Weight_On_SplineBySpline_Basis
Contains definitions for spline coefficients and structure used in both CPU and GPU code.
Class responsible for calculating spline weight on GPU.
cudaTextureObject_t text_nParamPerEvent_TF1
KS: Map keeping track how many parameters applies to each event, we keep two numbers here {number of ...
unsigned int * gpu_nParamPerEvent
KS: GPU map keeping track how many parameters applies to each event, we keep two numbers here {number...
__host__ void CleanupGPU_SplineMonolith(float *cpu_total_weights)
This function deallocates the resources allocated for the separate {x} and {ybcd} arrays in the and T...
cudaTextureObject_t text_coeff_x
KS: Textures are L1 cache variables which are well optimised for fetching. Make texture only for vari...
short int * gpu_paramNo_TF1_arr
CW: GPU array with the number of points per TF1 object.
__host__ void InitGPU_Vals(float **vals)
Allocate memory for spline segments.
int h_n_events
Number of events living on CPU.
float * gpu_weights
GPU arrays to hold weight for each spline.
float * gpu_coeff_many
GPU arrays to hold other coefficients.
unsigned int * gpu_nParamPerEvent_TF1
KS: GPU map keeping track how many parameters applies to each event, we keep two numbers here {number...
cudaTextureObject_t text_nParamPerEvent
KS: Map keeping track how many parameters applies to each event, we keep two numbers here {number of ...
__host__ void InitGPU_SplineMonolith(float **cpu_total_weights, int n_events, unsigned int total_nknots, unsigned int n_splines, unsigned int n_tf1, int Eve_size)
Allocate memory on gpu for spline monolith.
virtual ~SMonolithGPU()
destructor
__host__ void RunGPU_SplineMonolith(float *cpu_total_weights, float *vals, short int *segment, const unsigned int h_n_splines, const unsigned int h_n_tf1)
Run the GPU code for the separate many arrays. As in separate {x}, {y,b,c,d} arrays Pass the segment ...
float * gpu_coeff_x
KS: GPU arrays to hold X coefficient.
SMonolithGPU()
constructor
short int * gpu_nPoints_arr
GPU arrays to hold number of points.
float * gpu_weights_tf1
GPU arrays to hold weight for each TF1.
__host__ void InitGPU_Segments(short int **segment)
Allocate memory for spline segments.
unsigned int * gpu_nKnots_arr
KS: GPU Number of knots per spline.
int h_n_params
Number of params living on CPU.
float * gpu_coeff_TF1_many
GPU arrays to hold TF1 coefficients.
short int * gpu_paramNo_arr
CW: GPU array with the number of points per spline (not per spline point!)
__host__ void CopyToGPU_SplineMonolith(SplineMonoStruct *cpu_spline_handler, std::vector< float > cpu_many_array_TF1, std::vector< short int > cpu_paramNo_arr_TF1, int n_events, std::vector< unsigned int > cpu_nParamPerEvent, std::vector< unsigned int > cpu_nParamPerEvent_TF1, int n_params, unsigned int n_splines, short int spline_size, unsigned int total_nknots, unsigned int n_tf1)
Copies data from CPU to GPU for the spline monolith.
float * gpu_total_weights
GPU arrays to hold weight for event.
__host__ void CleanupGPU_Segments(short int *segment, float *vals)
Clean up pinned variables at CPU.
__global__ void EvalOnGPU_TotWeight(const float *__restrict__ gpu_weights, const float *__restrict__ gpu_weights_tf1, float *__restrict__ gpu_total_weights, const cudaTextureObject_t __restrict__ text_nParamPerEvent, const cudaTextureObject_t __restrict__ text_nParamPerEvent_TF1)
KS: Evaluate the total spline event weight on the GPU, as in most cases GPU is faster,...
__global__ void EvalOnGPU_Splines(const short int *__restrict__ gpu_paramNo_arr, const unsigned int *__restrict__ gpu_nKnots_arr, const float *__restrict__ gpu_coeff_many, float *__restrict__ gpu_weights, const cudaTextureObject_t __restrict__ text_coeff_x)
Evaluate the spline on the GPU Using one {y,b,c,d} array and one {x} array Should be most efficient a...
__host__ void SynchroniseSplines()
Make sure all Cuda threads finished execution.
__global__ void EvalOnGPU_TF1(const float *__restrict__ gpu_coeffs_tf1, const short int *__restrict__ gpu_paramNo_arr_tf1, float *__restrict__ gpu_weights_tf1)
Evaluate the TF1 on the GPU Using 5th order polynomial.
Common CUDA utilities and definitions for shared GPU functionality.
KS: Struct storing information for spline monolith.