gpuSplineUtils.cuh

Go to the documentation of this file.

 
//MaCh3 included
#include "Manager/gpuUtils.cuh"
#include "Splines/SplineCommon.h"
 
__host__ void SynchroniseSplines();
 
__global__ void EvalOnGPU_Splines(
    const unsigned int gpu_n_splines,
    const short int gpu_spline_size,
    const short int* __restrict__ gpu_paramNo_arr,
    const unsigned int* __restrict__ gpu_nKnots_arr,
    const float* __restrict__ gpu_coeff_many,
    const float* __restrict__ gpu_par_val,
    const short int* __restrict__ gpu_spline_segment,
    float* __restrict__ gpu_weights,
    const cudaTextureObject_t __restrict__ text_coeff_x);
 
__global__ void EvalOnGPU_TF1(
  const unsigned int gpu_n_TF1,
  const float* __restrict__ gpu_coeffs_tf1,
  const short int* __restrict__ gpu_paramNo_arr_tf1,
  const float* __restrict__ gpu_par_val,
  float* __restrict__ gpu_weights_tf1);
 
__global__ void EvalOnGPU_TotWeight(
  const int gpu_n_events,
  const float* __restrict__ gpu_weights,
  const float* __restrict__ gpu_weights_tf1,
  M3::float_t* __restrict__ gpu_total_weights,
  const cudaTextureObject_t __restrict__ text_nParamPerEvent,
  const cudaTextureObject_t __restrict__ text_nParamPerEvent_TF1);
 
class SplineMonolithGPU
{
  public:
    SplineMonolithGPU();
    virtual ~SplineMonolithGPU();
 
    __host__ void InitGPU_SplineMonolith(
      unsigned int total_nknots,
      unsigned int n_splines,
      unsigned int n_tf1,
      int Eve_size);
 
    __host__ void InitGPU_Unbinned_SplineMonolith(
        M3::float_t **cpu_total_weights,
        int n_events);
 
    __host__ void CopyToGPU_SplineMonolith_Unbinned(
      const SplineMonoStruct* cpu_spline_handler,
 
      // TFI related now
      const std::vector<float>& cpu_many_array_TF1,
      const std::vector<short int>& cpu_paramNo_arr_TF1,
      const int n_events,
      const std::vector<unsigned int>& cpu_nParamPerEvent,
      // TFI related now
      const std::vector<unsigned int>& cpu_nParamPerEvent_TF1,
      const int n_params,
      const unsigned int n_splines,
      const short int spline_size,
      const unsigned int total_nknots,
      const unsigned int n_tf1);
 
    __host__ void CopyToGPU_SplineMonolith_Binned(
      M3::float_t *manycoeff_arr,
      M3::float_t *xcoeff_arr,
      const std::vector<short int>& uniquesplinevec_Monolith,
      const std::vector<unsigned int>& coeffindexvec,
 
      const int n_params,
      const unsigned int n_splines,
      const short int spline_size,
      const unsigned int total_nknots);
 
    __host__ void InitGPU_Segments(short int **segment);
 
    __host__ void InitGPU_Vals(float **vals);
 
 
    __host__ void RunGPU_SplineMonolith_Unbinned(
      M3::float_t* cpu_total_weights,
      // Holds the changes in parameters
      float *vals,
      // Holds the segments for parameters
      short int *segment);
 
    __host__ void RunGPU_SplineMonolith_Binned(
      M3::float_t* cpu_spline_weights,
      // Holds the changes in parameters
      float *vals,
      // Holds the segments for parameters
      short int *segment);
 
    __host__ void CleanupPinnedMemory(M3::float_t *cpu_total_weights,
                                      short int *segment, float *vals);
 
  private:
    unsigned int *gpu_nParamPerEvent;
    unsigned int *gpu_nParamPerEvent_TF1;
 
    float *gpu_coeff_x;
 
    float *gpu_coeff_many;
 
    unsigned int *gpu_nKnots_arr;
 
    short int *gpu_paramNo_arr;
 
    float *gpu_coeff_TF1_many;
    short int *gpu_nPoints_arr;
    short int *gpu_paramNo_TF1_arr;
 
    float* gpu_par_val;
    short int* gpu_spline_segment;
 
    M3::float_t *gpu_total_weights;
    float *gpu_weights;
    float *gpu_weights_tf1;
 
    short int cpu_spline_size;
    unsigned int cpu_n_splines;
    unsigned int cpu_n_TF1;
    int cpu_n_params;
    int cpu_n_events;
    #ifndef _LOW_MEMORY_STRUCTS_
    std::vector<float> cpu_tmp_weights;
    #endif
 
    // ******************************************
    // TEXTURES
    // ******************************************
    cudaTextureObject_t text_coeff_x = 0;
    cudaTextureObject_t text_nParamPerEvent = 0;
    cudaTextureObject_t text_nParamPerEvent_TF1 = 0;
};