00001
00006 #ifndef __NLL_H__
00007 #define __NLL_H__
00008
00009 #include <cuda.h>
00010 #include <hemi/hemi.h>
00011
00012 #ifdef __CUDACC__
00013 #include <curand_kernel.h>
00014 #endif
00015
00016 #ifndef __HEMI_ARRAY_H__
00017 #define __HEMI_ARRAY_H__
00018 #include <hemi/array.h>
00019 #endif
00020
00025 #ifdef __CUDACC__
00026 typedef curandStateXORWOW RNGState;
00027 #else
00028 typedef int RNGState;
00029 #endif
00030
00031 class TNtuple;
00032
00033 #ifdef __CUDACC__
00034
00043 __global__ void init_device_rngs(int nthreads, unsigned long long seed,
00044 curandState* state);
00045 #endif
00046
00047
00059 HEMI_KERNEL(pick_new_vector)(int nthreads, RNGState* rng, float sigma,
00060 const float* current_vector,
00061 float* proposed_vector);
00062
00082 HEMI_KERNEL(jump_decider)(RNGState* rng, double* nll_current,
00083 const double* nll_proposed, float* v_current,
00084 const float* v_proposed, unsigned ns, int* counter,
00085 float* jump_buffer);
00086
00098 HEMI_KERNEL(nll_event_chunks)(const float* lut, const float* pars,
00099 const size_t ne, const size_t ns,
00100 double* sums);
00101
00102
00112 HEMI_KERNEL(nll_event_reduce)(const size_t nthreads, const double* sums,
00113 double* total_sum);
00114
00115
00129 HEMI_KERNEL(nll_total)(const size_t ns, const float* pars,
00130 const float* expectations,
00131 const float* constraints,
00132 const double* events_total,
00133 double* nll);
00134
00135 #endif // __NLL_H__
00136