MaCh3 2.2.1
Reference Guide
Loading...
Searching...
No Matches
gpuSplineUtils.cuh
Go to the documentation of this file.
1
11
12//MaCh3 included
13#include "Manager/gpuUtils.cuh"
15
17__host__ void SynchroniseSplines();
18
27__global__ void EvalOnGPU_Splines(
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);
33
38__global__ void EvalOnGPU_TF1(
39 const float* __restrict__ gpu_coeffs_tf1,
40 const short int* __restrict__ gpu_paramNo_arr_tf1,
41 float* __restrict__ gpu_weights_tf1);
42
43#ifndef Weight_On_SplineBySpline_Basis
50__global__ void EvalOnGPU_TotWeight(
51 const float* __restrict__ gpu_weights,
52 const float* __restrict__ gpu_weights_tf1,
53
54 float* __restrict__ gpu_total_weights,
55
56 const cudaTextureObject_t __restrict__ text_nParamPerEvent,
57 const cudaTextureObject_t __restrict__ text_nParamPerEvent_TF1);
58#endif
59
62{
63 public:
67 virtual ~SMonolithGPU();
68
75 __host__ void InitGPU_SplineMonolith(
76 #ifndef Weight_On_SplineBySpline_Basis
77 float **cpu_total_weights,
78 int n_events,
79 #endif
80 unsigned int total_nknots,
81 unsigned int n_splines,
82 unsigned int n_tf1,
83 int Eve_size);
84
102 __host__ void CopyToGPU_SplineMonolith(
103 SplineMonoStruct* cpu_spline_handler,
104
105 // TFI related now
106 std::vector<float> cpu_many_array_TF1,
107 std::vector<short int> cpu_paramNo_arr_TF1,
108 #ifndef Weight_On_SplineBySpline_Basis
109 int n_events,
110 std::vector<unsigned int> cpu_nParamPerEvent,
111 // TFI related now
112 std::vector<unsigned int> cpu_nParamPerEvent_TF1,
113 #endif
114 int n_params,
115 unsigned int n_splines,
116 short int spline_size,
117 unsigned int total_nknots,
118 unsigned int n_tf1);
119
122 __host__ void InitGPU_Segments(short int **segment);
123
126 __host__ void InitGPU_Vals(float **vals);
127
131
132
133
151 __host__ void RunGPU_SplineMonolith(
152 #ifdef Weight_On_SplineBySpline_Basis
153 float* cpu_weights,
154 float* cpu_weights_tf1,
155 #else
156 float* cpu_total_weights,
157 #endif
158 // Holds the changes in parameters
159 float *vals,
160 // Holds the segments for parameters
161 short int *segment,
162 const unsigned int h_n_splines,
163 const unsigned int h_n_tf1);
164
171 __host__ void CleanupGPU_SplineMonolith(
172 #ifndef Weight_On_SplineBySpline_Basis
173 float *cpu_total_weights
174 #endif
175 );
176
180 __host__ void CleanupGPU_Segments(short int *segment, float *vals);
181
182 private:
184 unsigned int *gpu_nParamPerEvent;
187
190
193
195 unsigned int *gpu_nKnots_arr;
196
198 short int *gpu_paramNo_arr;
199
203 short int *gpu_nPoints_arr;
206
213
214 // h_NAME declares HOST constants (live on CPU)
219
220 // ******************************************
221 // TEXTURES
222 // ******************************************
224 cudaTextureObject_t text_coeff_x = 0;
225 #ifndef Weight_On_SplineBySpline_Basis
227 cudaTextureObject_t text_nParamPerEvent = 0;
229 cudaTextureObject_t text_nParamPerEvent_TF1 = 0;
230 #endif
231};
232
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.
Definition: SplineCommon.h:30