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