8#ifdef AU_ROOFIT_BACKEND_CUDA
10#include <cuda_runtime.h>
13bool IsDevicePointer(
const void *ptr) {
14 cudaPointerAttributes attrs;
15 cudaError_t err = cudaPointerGetAttributes(&attrs, ptr);
16 if (err != cudaSuccess) {
21 return attrs.type == cudaMemoryTypeDevice ||
22 attrs.type == cudaMemoryTypeManaged;
34Double_t SoftPlusNeg(Double_t z) {
36 return std::log1p(std::exp(-z));
37 return -z + std::log1p(std::exp(z));
40Double_t SigmoidNeg(Double_t z) {
42 return 1.0 / (1.0 + std::exp(z));
43 return std::exp(-z) / (1.0 + std::exp(-z));
46Double_t StepAntideriv(Double_t z) {
return -SoftPlusNeg(z) + SigmoidNeg(z); }
57Double_t ExpErfc(Double_t a, Double_t b) {
59 return std::exp(a) * std::erfc(b);
60 Double_t t = 1.0 / (b * b);
61 Double_t erfcx = (1.0 / (b * std::sqrt(
M_PI))) *
62 (1.0 - 0.5 * t + 0.75 * t * t - 1.875 * t * t * t);
63 return std::exp(a - b * b) * erfcx;
71const Double_t kDensityFloor = 1e-300;
73Double_t ExpTailDensity(Double_t y, Double_t sigma, Double_t tau) {
74 Double_t sqrt2_sigma = std::sqrt(2.0) * sigma;
75 Double_t v = ExpErfc(y / tau, y / sqrt2_sigma);
76 return v > kDensityFloor ? v : kDensityFloor;
79Double_t ExpTailAntideriv(Double_t y, Double_t sigma, Double_t tau) {
80 Double_t sqrt2_sigma = std::sqrt(2.0) * sigma;
81 Double_t offset = sigma * sigma / tau;
82 Double_t gauss_corr = std::exp(sigma * sigma / (2.0 * tau * tau));
83 return tau * (ExpErfc(y / tau, y / sqrt2_sigma) +
84 gauss_corr * std::erf((y - offset) / sqrt2_sigma));
87Double_t LowLinDensity(Double_t y, Double_t sigma, Double_t slope) {
88 Double_t lin = 1.0 + slope * y;
91 Double_t sqrt2_sigma = std::sqrt(2.0) * sigma;
92 Double_t v = lin * std::erfc(y / sqrt2_sigma);
93 return v > kDensityFloor ? v : kDensityFloor;
96Double_t LowLinAntideriv(Double_t y, Double_t sigma, Double_t slope) {
97 Double_t sqrt2_sigma = std::sqrt(2.0) * sigma;
98 Double_t beta = 1.0 / sqrt2_sigma;
99 Double_t by = beta * y;
100 Double_t erfc_by = std::erfc(by);
101 Double_t erf_by = std::erf(by);
102 Double_t gauss = std::exp(-(by * by));
103 Double_t inv_beta_sqrt_pi = 1.0 / (beta * std::sqrt(
M_PI));
104 Double_t flat = y * erfc_by - gauss * inv_beta_sqrt_pi;
105 Double_t lin = 0.5 * y * y * erfc_by - y * gauss * inv_beta_sqrt_pi / 2.0 +
106 erf_by / (4.0 * beta * beta);
107 return flat + slope * lin;
110Double_t LowLinIntegral(Double_t y_lo, Double_t y_hi, Double_t sigma,
112 Double_t eff_lo = y_lo;
113 Double_t eff_hi = y_hi;
115 Double_t threshold = -1.0 / slope;
116 if (threshold > eff_lo)
118 if (eff_lo >= eff_hi)
120 }
else if (slope < -1e-10) {
121 Double_t threshold = -1.0 / slope;
122 if (threshold < eff_hi)
124 if (eff_lo >= eff_hi)
127 Double_t val = LowLinAntideriv(eff_hi, sigma, slope) -
128 LowLinAntideriv(eff_lo, sigma, slope);
129 if (!std::isfinite(val) || val < 1e-300)
137 RooAbsReal &mu, RooAbsReal &sigma)
138 : RooAbsPdf(name, title),
x_(
"x",
"x", this, x),
mu_(
"mu",
"mu", this, mu),
139 sigma_(
"sigma",
"sigma", this, sigma) {}
142 : RooAbsPdf(other, name),
x_(
"x", this, other.
x_),
146 Double_t sigma = (Double_t)
sigma_;
148 return kDensityFloor;
149 Double_t z = ((Double_t)
x_ - (Double_t)
mu_) / sigma;
150 Double_t s = SigmoidNeg(z);
152 return v > kDensityFloor ? v : kDensityFloor;
156 std::span<const double> x_vals = ctx.at(
x_);
157 std::span<double> output = ctx.output();
158 size_t n = output.size();
159 Double_t sigma = ctx.at(
sigma_)[0];
160 Double_t mu = ctx.at(
mu_)[0];
162 for (
size_t i = 0; i < n; ++i)
163 output[i] = kDensityFloor;
166 Double_t inv_sigma = 1.0 / sigma;
167#ifdef AU_ROOFIT_BACKEND_CUDA
168 if (IsDevicePointer(output.data())) {
173#pragma omp parallel for simd
174 for (
size_t i = 0; i < n; ++i) {
175 Double_t z = (x_vals[i] - mu) * inv_sigma;
176 Double_t s = SigmoidNeg(z);
178 output[i] = v > kDensityFloor ? v : kDensityFloor;
184 const char * )
const {
185 if (matchArgs(allVars, analVars,
x_))
191 const char *rangeName)
const {
194 Double_t sigma = (Double_t)
sigma_;
197 Double_t mu = (Double_t)
mu_;
198 Double_t x_lo =
x_.min(rangeName);
199 Double_t x_hi =
x_.max(rangeName);
200 Double_t z_lo = (x_lo - mu) / sigma;
201 Double_t z_hi = (x_hi - mu) / sigma;
202 Double_t val = sigma * (StepAntideriv(z_hi) - StepAntideriv(z_lo));
203 if (!std::isfinite(val) || val < 1e-300)
209 RooAbsReal &mu, RooAbsReal &sigma,
210 RooAbsReal &tau_ratio)
211 : RooAbsPdf(name, title),
x_(
"x",
"x", this, x),
mu_(
"mu",
"mu", this, mu),
212 sigma_(
"sigma",
"sigma", this, sigma),
213 tau_ratio_(
"tau_ratio",
"tau_ratio", this, tau_ratio) {}
216 : RooAbsPdf(other, name),
x_(
"x", this, other.
x_),
221 Double_t sigma = (Double_t)
sigma_;
223 if (sigma <= 0 || tau_ratio <= 0)
225 Double_t tau = tau_ratio * sigma;
226 Double_t y = (Double_t)
x_ - (Double_t)
mu_;
227 return ExpTailDensity(y, sigma, tau);
231 std::span<const double> x_vals = ctx.at(
x_);
232 std::span<double> output = ctx.output();
233 size_t n = output.size();
234 Double_t sigma = ctx.at(
sigma_)[0];
235 Double_t mu = ctx.at(
mu_)[0];
237 if (sigma <= 0 || tau_ratio <= 0) {
238 for (
size_t i = 0; i < n; ++i)
239 output[i] = kDensityFloor;
242 Double_t tau = tau_ratio * sigma;
243 Double_t inv_tau = 1.0 / tau;
244 Double_t inv_sqrt2_sigma = 1.0 / (std::sqrt(2.0) * sigma);
245#ifdef AU_ROOFIT_BACKEND_CUDA
246 if (IsDevicePointer(output.data())) {
252#pragma omp parallel for simd
253 for (
size_t i = 0; i < n; ++i) {
254 Double_t y = x_vals[i] - mu;
255 Double_t v = ExpErfc(y * inv_tau, y * inv_sqrt2_sigma);
256 output[i] = v > kDensityFloor ? v : kDensityFloor;
262 const char * )
const {
263 if (matchArgs(allVars, analVars,
x_))
269 const char *rangeName)
const {
272 Double_t sigma = (Double_t)
sigma_;
274 if (sigma <= 0 || tau_ratio <= 0)
276 Double_t tau = tau_ratio * sigma;
277 Double_t mu = (Double_t)
mu_;
278 Double_t x_lo =
x_.min(rangeName);
279 Double_t x_hi =
x_.max(rangeName);
280 Double_t y_lo = x_lo - mu;
281 Double_t y_hi = x_hi - mu;
283 ExpTailAntideriv(y_hi, sigma, tau) - ExpTailAntideriv(y_lo, sigma, tau);
284 if (!std::isfinite(val) || val < 1e-300)
290 RooAbsReal &mu, RooAbsReal &sigma,
292 : RooAbsPdf(name, title),
x_(
"x",
"x", this, x),
mu_(
"mu",
"mu", this, mu),
293 sigma_(
"sigma",
"sigma", this, sigma),
294 slope_(
"slope",
"slope", this, slope) {}
297 : RooAbsPdf(other, name),
x_(
"x", this, other.
x_),
302 Double_t sigma = (Double_t)
sigma_;
305 Double_t y = (Double_t)
x_ - (Double_t)
mu_;
306 Double_t slope = (Double_t)
slope_;
307 return LowLinDensity(y, sigma, slope);
311 std::span<const double> x_vals = ctx.at(
x_);
312 std::span<double> output = ctx.output();
313 size_t n = output.size();
314 Double_t sigma = ctx.at(
sigma_)[0];
315 Double_t mu = ctx.at(
mu_)[0];
316 Double_t slope = ctx.at(
slope_)[0];
318 for (
size_t i = 0; i < n; ++i)
319 output[i] = kDensityFloor;
322 Double_t inv_sqrt2_sigma = 1.0 / (std::sqrt(2.0) * sigma);
323#ifdef AU_ROOFIT_BACKEND_CUDA
324 if (IsDevicePointer(output.data())) {
330#pragma omp parallel for simd
331 for (
size_t i = 0; i < n; ++i) {
332 Double_t y = x_vals[i] - mu;
333 Double_t lin = 1.0 + slope * y;
336 Double_t v = lin * std::erfc(y * inv_sqrt2_sigma);
337 output[i] = v > kDensityFloor ? v : kDensityFloor;
343 const char * )
const {
344 if (matchArgs(allVars, analVars,
x_))
350 const char *rangeName)
const {
353 Double_t sigma = (Double_t)
sigma_;
356 Double_t mu = (Double_t)
mu_;
357 Double_t slope = (Double_t)
slope_;
358 Double_t x_lo =
x_.min(rangeName);
359 Double_t x_hi =
x_.max(rangeName);
360 Double_t y_lo = x_lo - mu;
361 Double_t y_hi = x_hi - mu;
362 return LowLinIntegral(y_lo, y_hi, sigma, slope);
366 RooAbsReal &x, RooAbsReal &mu, RooAbsReal &sigma,
367 RooAbsReal &tau_ratio)
368 : RooAbsPdf(name, title),
x_(
"x",
"x", this, x),
mu_(
"mu",
"mu", this, mu),
369 sigma_(
"sigma",
"sigma", this, sigma),
370 tau_ratio_(
"tau_ratio",
"tau_ratio", this, tau_ratio) {}
373 : RooAbsPdf(other, name),
x_(
"x", this, other.
x_),
378 Double_t sigma = (Double_t)
sigma_;
380 if (sigma <= 0 || tau_ratio <= 0)
382 Double_t tau = tau_ratio * sigma;
383 Double_t z = (Double_t)
mu_ - (Double_t)
x_;
384 return ExpTailDensity(z, sigma, tau);
388 std::span<const double> x_vals = ctx.at(
x_);
389 std::span<double> output = ctx.output();
390 size_t n = output.size();
391 Double_t sigma = ctx.at(
sigma_)[0];
392 Double_t mu = ctx.at(
mu_)[0];
394 if (sigma <= 0 || tau_ratio <= 0) {
395 for (
size_t i = 0; i < n; ++i)
396 output[i] = kDensityFloor;
399 Double_t tau = tau_ratio * sigma;
400 Double_t inv_tau = 1.0 / tau;
401 Double_t inv_sqrt2_sigma = 1.0 / (std::sqrt(2.0) * sigma);
402#ifdef AU_ROOFIT_BACKEND_CUDA
403 if (IsDevicePointer(output.data())) {
409#pragma omp parallel for simd
410 for (
size_t i = 0; i < n; ++i) {
411 Double_t z = mu - x_vals[i];
412 Double_t v = ExpErfc(z * inv_tau, z * inv_sqrt2_sigma);
413 output[i] = v > kDensityFloor ? v : kDensityFloor;
419 const char * )
const {
420 if (matchArgs(allVars, analVars,
x_))
426 const char *rangeName)
const {
429 Double_t sigma = (Double_t)
sigma_;
431 if (sigma <= 0 || tau_ratio <= 0)
433 Double_t tau = tau_ratio * sigma;
434 Double_t mu = (Double_t)
mu_;
435 Double_t x_lo =
x_.min(rangeName);
436 Double_t x_hi =
x_.max(rangeName);
437 Double_t z_lo = mu - x_lo;
438 Double_t z_hi = mu - x_hi;
440 ExpTailAntideriv(z_lo, sigma, tau) - ExpTailAntideriv(z_hi, sigma, tau);
441 if (!std::isfinite(val) || val < 1e-300)
Host-callable launch wrappers for the photopeak PDF CUDA kernels.
void RooLowLinTail_launchKernel(double *output, const double *x_vals, double mu, double slope, double inv_sqrt2_sigma, size_t n)
Launch the low-energy linear tail kernel.
void RooHighExpTail_launchKernel(double *output, const double *x_vals, double mu, double inv_tau, double inv_sqrt2_sigma, size_t n)
Launch the high-energy exponential tail kernel.
void RooLowExpTail_launchKernel(double *output, const double *x_vals, double mu, double inv_tau, double inv_sqrt2_sigma, size_t n)
Launch the low-energy exponential tail kernel.
void RooStepShelf_launchKernel(double *output, const double *x_vals, double mu, double inv_sigma, size_t n)
Launch the resolution-smeared step-shelf kernel.
Custom RooFit PDF components for gamma-ray photopeak fitting.
Exponential tail above the photopeak, convolved with the resolution.
Int_t getAnalyticalIntegral(RooArgSet &allVars, RooArgSet &analVars, const char *rangeName=nullptr) const override
Advertise the analytic integral over the observable.
RooHighExpTail()
Default constructor for ROOT I/O; leaves proxies unbound.
Double_t evaluate() const override
Scalar evaluation for a single observable value.
Double_t analyticalIntegral(Int_t code, const char *rangeName=nullptr) const override
Evaluate the analytic integral advertised above.
void doEval(RooFit::EvalContext &ctx) const override
Batched evaluation over every event in one pass.
Exponential tail below the photopeak, convolved with the resolution.
Double_t evaluate() const override
Scalar evaluation for a single observable value.
RooLowExpTail()
Default constructor for ROOT I/O; leaves proxies unbound.
Int_t getAnalyticalIntegral(RooArgSet &allVars, RooArgSet &analVars, const char *rangeName=nullptr) const override
Advertise the analytic integral over the observable.
Double_t analyticalIntegral(Int_t code, const char *rangeName=nullptr) const override
Evaluate the analytic integral advertised above.
void doEval(RooFit::EvalContext &ctx) const override
Batched evaluation over every event in one pass.
Linear tail below the photopeak, convolved with the resolution.
Double_t analyticalIntegral(Int_t code, const char *rangeName=nullptr) const override
Evaluate the analytic integral advertised above.
RooLowLinTail()
Default constructor for ROOT I/O; leaves proxies unbound.
void doEval(RooFit::EvalContext &ctx) const override
Batched evaluation over every event in one pass.
Int_t getAnalyticalIntegral(RooArgSet &allVars, RooArgSet &analVars, const char *rangeName=nullptr) const override
Advertise the analytic integral over the observable.
Double_t evaluate() const override
Scalar evaluation for a single observable value.
Resolution-smeared step on the low side of the photopeak.
Double_t analyticalIntegral(Int_t code, const char *rangeName=nullptr) const override
Evaluate the analytic integral advertised above.
void doEval(RooFit::EvalContext &ctx) const override
Batched evaluation over every event in one pass.
Int_t getAnalyticalIntegral(RooArgSet &allVars, RooArgSet &analVars, const char *rangeName=nullptr) const override
Advertise the analytic integral over the observable.
RooStepShelf()
Default constructor for ROOT I/O; leaves proxies unbound.
Double_t evaluate() const override
Scalar evaluation for a single observable value.