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 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);
 
__global__ void EvalOnGPU_TF1(
  const float* __restrict__ gpu_coeffs_tf1,
  const short int* __restrict__ gpu_paramNo_arr_tf1,
  float* __restrict__ gpu_weights_tf1);
 
#ifndef Weight_On_SplineBySpline_Basis
__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);
#endif
 
class SMonolithGPU
{
  public:
    SMonolithGPU();
    virtual ~SMonolithGPU();
 
    __host__ void InitGPU_SplineMonolith(
      #ifndef Weight_On_SplineBySpline_Basis
      float **cpu_total_weights,
      int n_events,
      #endif
      unsigned int total_nknots,
      unsigned int n_splines,
      unsigned int n_tf1,
      int Eve_size);
 
    __host__ void CopyToGPU_SplineMonolith(
      SplineMonoStruct* cpu_spline_handler,
 
      // TFI related now
      std::vector<float> cpu_many_array_TF1,
      std::vector<short int> cpu_paramNo_arr_TF1,
      #ifndef Weight_On_SplineBySpline_Basis
      int n_events,
      std::vector<unsigned int> cpu_nParamPerEvent,
      // TFI related now
      std::vector<unsigned int> cpu_nParamPerEvent_TF1,
      #endif
      int n_params,
      unsigned int n_splines,
      short int spline_size,
      unsigned int total_nknots,
      unsigned int n_tf1);
 
    __host__ void InitGPU_Segments(short int **segment);
 
    __host__ void InitGPU_Vals(float **vals);
 
 
 
 
    __host__ void RunGPU_SplineMonolith(
      #ifdef Weight_On_SplineBySpline_Basis
      float* cpu_weights,
      float* cpu_weights_tf1,
      #else
      float* cpu_total_weights,
      #endif
      // Holds the changes in parameters
      float *vals,
      // Holds the segments for parameters
      short int *segment,
      const unsigned int h_n_splines,
      const unsigned int h_n_tf1);
 
    __host__ void CleanupGPU_SplineMonolith(
      #ifndef Weight_On_SplineBySpline_Basis
      float *cpu_total_weights
      #endif
      );
 
    __host__ void CleanupGPU_Segments(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_total_weights;
    float *gpu_weights;
    float *gpu_weights_tf1;
 
    // h_NAME declares HOST constants (live on CPU)
    int h_n_params;
    int h_n_events;
 
    // ******************************************
    // TEXTURES
    // ******************************************
    cudaTextureObject_t text_coeff_x = 0;
    #ifndef Weight_On_SplineBySpline_Basis
    cudaTextureObject_t text_nParamPerEvent = 0;
    cudaTextureObject_t text_nParamPerEvent_TF1 = 0;
    #endif
};