Skip to content

Commit 0d868fd

Browse files
jzmaddockjzmaddock
authored
Merge pull request #1349 from JacobHass8/log-lower-incomplete-gamma
Implementation of the log of the lower incomplete gamma function
2 parents d8597d3 + 76c9c2b commit 0d868fd

File tree

9 files changed

+346
-2
lines changed

9 files changed

+346
-2
lines changed

‎doc/sf/igamma.qbk‎

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -26,6 +26,12 @@
2626
template <class T1, class T2, class ``__Policy``>
2727
BOOST_MATH_GPU_ENABLED ``__sf_result`` lgamma_q(T1 a, T2 z, const ``__Policy``&);
2828

29+
template <class T1, class T2>
30+
BOOST_MATH_GPU_ENABLED ``__sf_result`` lgamma_p(T1 a, T2 z);
31+
32+
template <class T1, class T2, class ``__Policy``>
33+
BOOST_MATH_GPU_ENABLED ``__sf_result`` lgamma_p(T1 a, T2 z, const ``__Policy``&);
34+
2935
template <class T1, class T2>
3036
BOOST_MATH_GPU_ENABLED ``__sf_result`` tgamma_lower(T1 a, T2 z);
3137

@@ -72,6 +78,15 @@ This function changes rapidly from 0 to 1 around the point z == a:
7278

7379
[graph gamma_p]
7480

81+
template <class T1, class T2>
82+
BOOST_MATH_GPU_ENABLED ``__sf_result`` lgamma_p(T1 a, T2 z);
83+
84+
template <class T1, class T2, class ``__Policy``>
85+
BOOST_MATH_GPU_ENABLED ``__sf_result`` lgamma_p(T1 a, T2 z, const ``__Policy``&);
86+
87+
Returns the natural log of the normalized lower incomplete gamma function
88+
of a and z.
89+
7590
template <class T1, class T2>
7691
BOOST_MATH_GPU_ENABLED ``__sf_result`` gamma_q(T1 a, T2 z);
7792

‎include/boost/math/special_functions/gamma.hpp‎

Lines changed: 63 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1290,7 +1290,6 @@ BOOST_MATH_GPU_ENABLED T incomplete_tgamma_large_x(const T& a, const T& x, const
12901290
return result;
12911291
}
12921292

1293-
12941293
//
12951294
// Main incomplete gamma entry point, handles all four incomplete gamma's:
12961295
//
@@ -1813,6 +1812,46 @@ BOOST_MATH_GPU_ENABLED T lgamma_incomplete_imp(T a, T x, const Policy& pol)
18131812
return log(gamma_q(a, x, pol));
18141813
}
18151814

1815+
// Calculate log of incomplete gamma function
1816+
template <class T, class Policy>
1817+
BOOST_MATH_GPU_ENABLED T lgamma_incomplete_lower_imp(T a, T x, const Policy& pol)
1818+
{
1819+
using namespace boost::math; // temporary until we're in the right namespace
1820+
1821+
BOOST_MATH_STD_USING_CORE
1822+
1823+
// Check for invalid inputs (a < 0 or x < 0)
1824+
constexpr auto function = "boost::math::lgamma_p<%1%>(%1%, %1%)";
1825+
if(a <= 0)
1826+
return policies::raise_domain_error<T>(function, "Argument a to the incomplete gamma function must be greater than zero (got a=%1%).", a, pol);
1827+
if(x < 0)
1828+
return policies::raise_domain_error<T>(function, "Argument x to the incomplete gamma function must be >= 0 (got x=%1%).", x, pol);
1829+
1830+
// Taken from conditions on Line 1709. There are also
1831+
// conditions on Line 1368, but didn't implement that one here.
1832+
// The second condition ensures floats do not return -inf for small
1833+
// values of x.
1834+
if (((a > 4 * x) && (a >= max_factorial<T>::value)) || ((T(-0.4) / log(x) < a) && (x < T(1.0)))){
1835+
return log(detail::lower_gamma_series(a, x, pol)) - log(a) + a * log(x) - x - lgamma(a, pol);
1836+
}
1837+
1838+
//
1839+
// Can't do better than taking the log of P, but...
1840+
//
1841+
// Figure out whether we need P or Q, since if we calculate P and it's too close to unity
1842+
// we will lose precision in the result, selection logic here is extracted from gamma_incomplete_imp_final:
1843+
//
1844+
bool need_p = false;
1845+
if ((x < 1.1) && (x >= 0.5) && (x * 0.75f < a))
1846+
need_p = true;
1847+
else if ((x < a) && (x >= 1.1))
1848+
need_p = true;
1849+
1850+
if (need_p)
1851+
return log(gamma_p(a, x, pol));
1852+
return log1p(-gamma_q(a, x, pol), pol);
1853+
}
1854+
18161855
//
18171856
// Ratios of two gamma functions:
18181857
//
@@ -2454,6 +2493,29 @@ BOOST_MATH_GPU_ENABLED inline tools::promote_args_t<T1, T2> lgamma_q(T1 a, T2 z)
24542493
{
24552494
return lgamma_q(a, z, policies::policy<>());
24562495
}
2496+
2497+
template <class T1, class T2, class Policy>
2498+
BOOST_MATH_GPU_ENABLED inline tools::promote_args_t<T1, T2> lgamma_p(T1 a, T2 z, const Policy& /* pol */)
2499+
{
2500+
typedef tools::promote_args_t<T1, T2> result_type;
2501+
typedef typename policies::evaluation<result_type, Policy>::type value_type;
2502+
typedef typename policies::normalise<
2503+
Policy,
2504+
policies::promote_float<false>,
2505+
policies::promote_double<false>,
2506+
policies::discrete_quantile<>,
2507+
policies::assert_undefined<> >::type forwarding_policy;
2508+
2509+
return policies::checked_narrowing_cast<result_type, forwarding_policy>(
2510+
detail::lgamma_incomplete_lower_imp(static_cast<value_type>(a),
2511+
static_cast<value_type>(z), forwarding_policy()), "lgamma_p<%1%>(%1%, %1%)");
2512+
}
2513+
2514+
template <class T1, class T2>
2515+
BOOST_MATH_GPU_ENABLED inline tools::promote_args_t<T1, T2> lgamma_p(T1 a, T2 z)
2516+
{
2517+
return lgamma_p(a, z, policies::policy<>());
2518+
}
24572519
//
24582520
// Regularised lower incomplete gamma:
24592521
//

‎include/boost/math/special_functions/math_fwd.hpp‎

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -567,6 +567,12 @@ namespace boost
567567
template <class RT1, class RT2, class Policy>
568568
BOOST_MATH_GPU_ENABLED tools::promote_args_t<RT1, RT2> lgamma_q(RT1 a, RT2 z, const Policy&);
569569

570+
template <class RT1, class RT2>
571+
BOOST_MATH_GPU_ENABLED tools::promote_args_t<RT1, RT2> lgamma_p(RT1 a, RT2 z);
572+
573+
template <class RT1, class RT2, class Policy>
574+
BOOST_MATH_GPU_ENABLED tools::promote_args_t<RT1, RT2> lgamma_p(RT1 a, RT2 z, const Policy&);
575+
570576
template <class RT1, class RT2>
571577
BOOST_MATH_GPU_ENABLED tools::promote_args_t<RT1, RT2> gamma_p(RT1 a, RT2 z);
572578

@@ -1525,6 +1531,9 @@ namespace boost
15251531
\
15261532
template <class RT1, class RT2>\
15271533
BOOST_MATH_GPU_ENABLED inline boost::math::tools::promote_args_t<RT1, RT2> lgamma_q(RT1 a, RT2 z){ return boost::math::lgamma_q(a, z, Policy()); }\
1534+
\
1535+
template <class RT1, class RT2>\
1536+
BOOST_MATH_GPU_ENABLED inline boost::math::tools::promote_args_t<RT1, RT2> lgamma_p(RT1 a, RT2 z){ return boost::math::lgamma_p(a, z, Policy()); }\
15281537
\
15291538
template <class RT1, class RT2>\
15301539
BOOST_MATH_GPU_ENABLED inline boost::math::tools::promote_args_t<RT1, RT2> gamma_p(RT1 a, RT2 z){ return boost::math::gamma_p(a, z, Policy()); }\

‎test/compile_test/instantiate.hpp‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -264,6 +264,7 @@ void instantiate(RealType)
264264
boost::math::gamma_p(v1, v2);
265265
boost::math::gamma_q(v1, v2);
266266
boost::math::lgamma_q(v1, v2);
267+
boost::math::lgamma_p(v1, v2);
267268
boost::math::gamma_p_inv(v1, v2);
268269
boost::math::gamma_q_inv(v1, v2);
269270
boost::math::gamma_p_inva(v1, v2);
@@ -544,6 +545,7 @@ void instantiate(RealType)
544545
boost::math::gamma_p(v1 * 1, v2 + 0);
545546
boost::math::gamma_q(v1 * 1, v2 + 0);
546547
boost::math::lgamma_q(v1 * 1, v2 + 0);
548+
boost::math::lgamma_p(v1 * 1, v2 + 0);
547549
boost::math::gamma_p_inv(v1 * 1, v2 + 0);
548550
boost::math::gamma_q_inv(v1 * 1, v2 + 0);
549551
boost::math::gamma_p_inva(v1 * 1, v2 + 0);
@@ -796,6 +798,7 @@ void instantiate(RealType)
796798
boost::math::gamma_p(v1, v2, pol);
797799
boost::math::gamma_q(v1, v2, pol);
798800
boost::math::lgamma_q(v1, v2, pol);
801+
boost::math::lgamma_p(v1, v2, pol);
799802
boost::math::gamma_p_inv(v1, v2, pol);
800803
boost::math::gamma_q_inv(v1, v2, pol);
801804
boost::math::gamma_p_inva(v1, v2, pol);
@@ -1074,6 +1077,7 @@ void instantiate(RealType)
10741077
test::gamma_p(v1, v2);
10751078
test::gamma_q(v1, v2);
10761079
test::lgamma_q(v1, v2);
1080+
test::lgamma_p(v1, v2);
10771081
test::gamma_p_inv(v1, v2);
10781082
test::gamma_q_inv(v1, v2);
10791083
test::gamma_p_inva(v1, v2);
@@ -1356,6 +1360,7 @@ void instantiate_mixed(RealType)
13561360
boost::math::gamma_p(fr, lr);
13571361
boost::math::gamma_q(i, s);
13581362
boost::math::lgamma_q(i, s);
1363+
boost::math::lgamma_p(i, s);
13591364
boost::math::gamma_q(fr, lr);
13601365
boost::math::gamma_p_inv(i, fr);
13611366
boost::math::gamma_q_inv(s, fr);
@@ -1572,6 +1577,7 @@ void instantiate_mixed(RealType)
15721577
boost::math::gamma_p(fr, lr, pol);
15731578
boost::math::gamma_q(i, s, pol);
15741579
boost::math::lgamma_q(i, s, pol);
1580+
boost::math::lgamma_p(i, s, pol);
15751581
boost::math::gamma_q(fr, lr, pol);
15761582
boost::math::gamma_p_inv(i, fr, pol);
15771583
boost::math::gamma_q_inv(s, fr, pol);
@@ -1784,8 +1790,10 @@ void instantiate_mixed(RealType)
17841790
test::gamma_p(fr, lr);
17851791
test::gamma_q(i, s);
17861792
test::lgamma_q(i, s);
1793+
test::lgamma_p(i, s);
17871794
test::gamma_q(fr, lr);
17881795
test::lgamma_q(fr, lr);
1796+
test::lgamma_p(fr, lr);
17891797
test::gamma_p_inv(i, fr);
17901798
test::gamma_q_inv(s, fr);
17911799
test::gamma_p_inva(i, lr);

‎test/compile_test/sf_gamma_incl_test.cpp‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -45,6 +45,12 @@ void compile_and_link_test()
4545
check_result<long double>(boost::math::lgamma_q<long double>(l, l));
4646
#endif
4747

48+
check_result<float>(boost::math::lgamma_p<float>(f, f));
49+
check_result<double>(boost::math::lgamma_p<double>(d, d));
50+
#ifndef BOOST_MATH_NO_LONG_DOUBLE_MATH_FUNCTIONS
51+
check_result<long double>(boost::math::lgamma_p<long double>(l, l));
52+
#endif
53+
4854
check_result<float>(boost::math::gamma_p_inv<float>(f, f));
4955
check_result<double>(boost::math::gamma_p_inv<double>(d, d));
5056
#ifndef BOOST_MATH_NO_LONG_DOUBLE_MATH_FUNCTIONS

‎test/cuda_jamfile‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -371,6 +371,8 @@ run test_gamma_p_inv_double.cu ;
371371
run test_gamma_p_inv_float.cu ;
372372
run test_lgamma_q_double.cu ;
373373
run test_lgamma_q_float.cu ;
374+
run test_lgamma_p_double.cu ;
375+
run test_lgamma_p_float.cu ;
374376

375377
run test_log1p_double.cu ;
376378
run test_log1p_float.cu ;

‎test/test_igamma.hpp‎

Lines changed: 39 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -263,7 +263,19 @@ void test_spots(T, const char* name = nullptr)
263263
BOOST_CHECK_CLOSE(::boost::math::lgamma_q(static_cast<T>(501.25), static_cast<T>(2000)), static_cast<T>(-810.2453406781655559126505101822969531699112391075198076300675402L), tolerance);
264264
BOOST_CHECK_CLOSE(::boost::math::lgamma_q(static_cast<T>(20), static_cast<T>(0.25)), static_cast<T>(-2.946458104491857816330873290969917497748067639461638294404e-31L), tolerance);
265265
BOOST_CHECK_CLOSE(::boost::math::lgamma_q(static_cast<T>(40), static_cast<T>(0.75)), static_cast<T>(-5.930604927955460343652485525435087275997461623988991819824e-54L), tolerance);
266-
#if defined(__CYGWIN__) || defined(__MINGW32__)
266+
267+
//
268+
// Check that lgamma_q returns correct values with spot values calculated via wolframalpha log(P[a, x])
269+
// This is calculated using: N[Log[GammaRegularized[a,0, z]],64]
270+
//
271+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(500), static_cast<T>(10)), static_cast<T>(-1470.017750815998931281954666549641187420649099004671023115157832L), tolerance);
272+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(100), static_cast<T>(0.25)), static_cast<T>(-502.6163334118978895536207514636026023439623265152862757105793000L), tolerance);
273+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(20), static_cast<T>(10.25)), static_cast<T>(-5.404004887981642339930593767572610169901594898478031307722239712L), tolerance);
274+
// Small "a" produce larger errors
275+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(0.25), static_cast<T>(100)), static_cast<T>(-3.220751038854414755009496530271388459559061551701603447517040280e-46L), tolerance);
276+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(0.25), static_cast<T>(10)), static_cast<T>(-2.083032578160285760530530498275075010777428544413918699832758176e-6L), tolerance);
277+
278+
#if defined(__CYGWIN__) || defined(__MINGW32__)
267279
T gcc_win_mul = 2;
268280
#else
269281
T gcc_win_mul = 1;
@@ -287,6 +299,24 @@ void test_spots(T, const char* name = nullptr)
287299
BOOST_CHECK_CLOSE(::boost::math::lgamma_q(static_cast<T>(1200), static_cast<T>(1250.25)), static_cast<T>(-2.591934862117586205519309712218581885256650074210410262843591453L), tolerance * ((std::numeric_limits<T>::max_digits10 >= 36 || std::is_same<T, boost::math::concepts::real_concept>::value) ? 750 : (std::is_same<T, float>::value ? 1 : 50))); // Test fails on ARM64 and s390x long doubles and real_concept types unless tolerance is adjusted
288300
BOOST_CHECK_CLOSE(::boost::math::lgamma_q(static_cast<T>(2200), static_cast<T>(2249.75)), static_cast<T>(-1.933779894897391651410597618307863427927461116308937004149240320L), tolerance * (std::is_floating_point<T>::value ? 1 : 10));
289301
BOOST_CHECK_CLOSE(::boost::math::lgamma_q(static_cast<T>(2200), static_cast<T>(2250.25)), static_cast<T>(-1.950346484067948344620463026377077515919992808640737320057812268L), tolerance * (std::is_same<T, float>::value ? 1 : (std::is_floating_point<T>::value ? 100 : 200)));
302+
303+
// Long double and real_concept types need increased precision
304+
T real_concept_tol = 1;
305+
if (std::is_same<T, boost::math::concepts::real_concept>::value || std::is_same<T, long double>::value){
306+
real_concept_tol = 3;
307+
}
308+
// Pair of tests that bisect the crossover condition in our code at double and then quad precision
309+
// Oddly, the crossover condition is smaller for quad precision. This is because max_factorial is 100
310+
// for boost::multiprecision::cpp_bin_float_quad and 170 for doubles.
311+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(169.75), static_cast<T>(0.75)), static_cast<T>(-754.8681912874632573100058312311927462406154378562940316233389406L), tolerance * real_concept_tol);
312+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(170.25), static_cast<T>(0.75)), static_cast<T>(-757.5814133895304434271729579978676692688834086380018151200693572L), tolerance * real_concept_tol);
313+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(99.75), static_cast<T>(0.75)), static_cast<T>(-392.0259615581237826290999388631292473247947826682978959914359465L), tolerance * real_concept_tol);
314+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(100.25), static_cast<T>(0.75)), static_cast<T>(-394.4749200332583219473980963811639065003421270272773619742710832L), tolerance * real_concept_tol);
315+
316+
// Check large a, x values. Precision just isn't great here.
317+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(1450.25), static_cast<T>(1500.75)), static_cast<T>(-0.09812447528127799786140178403478691390753413399549580160096975713L), tolerance * (std::is_same<T, boost::math::concepts::real_concept>::value ? 16 : 1));
318+
BOOST_CHECK_CLOSE(::boost::math::lgamma_p(static_cast<T>(2000), static_cast<T>(1900)), static_cast<T>(-4.448523733381445722945397105917814000790587922314824687065050805L), tolerance * gcc_win_mul * (std::is_same<T, boost::math::concepts::real_concept>::value ? 8 : 1));
319+
290320
//
291321
// Coverage:
292322
//
@@ -302,6 +332,10 @@ void test_spots(T, const char* name = nullptr)
302332
BOOST_CHECK_THROW(boost::math::lgamma_q(static_cast<T>(1), static_cast<T>(-2)), std::domain_error);
303333
BOOST_CHECK_THROW(boost::math::lgamma_q(static_cast<T>(0), static_cast<T>(2)), std::domain_error);
304334

335+
BOOST_CHECK_THROW(boost::math::lgamma_p(static_cast<T>(-1), static_cast<T>(2)), std::domain_error);
336+
BOOST_CHECK_THROW(boost::math::lgamma_p(static_cast<T>(1), static_cast<T>(-2)), std::domain_error);
337+
BOOST_CHECK_THROW(boost::math::lgamma_p(static_cast<T>(0), static_cast<T>(2)), std::domain_error);
338+
305339
BOOST_CHECK_THROW(boost::math::gamma_p_derivative(static_cast<T>(-1), static_cast<T>(2)), std::domain_error);
306340
BOOST_CHECK_THROW(boost::math::gamma_p_derivative(static_cast<T>(1), static_cast<T>(-2)), std::domain_error);
307341
BOOST_CHECK_THROW(boost::math::gamma_p_derivative(static_cast<T>(0), static_cast<T>(2)), std::domain_error);
@@ -317,6 +351,10 @@ void test_spots(T, const char* name = nullptr)
317351
BOOST_CHECK((boost::math::isnan)(boost::math::lgamma_q(static_cast<T>(1), static_cast<T>(-2))));
318352
BOOST_CHECK((boost::math::isnan)(boost::math::lgamma_q(static_cast<T>(0), static_cast<T>(2))));
319353

354+
BOOST_CHECK((boost::math::isnan)(boost::math::lgamma_p(static_cast<T>(-1), static_cast<T>(2))));
355+
BOOST_CHECK((boost::math::isnan)(boost::math::lgamma_p(static_cast<T>(1), static_cast<T>(-2))));
356+
BOOST_CHECK((boost::math::isnan)(boost::math::lgamma_p(static_cast<T>(0), static_cast<T>(2))));
357+
320358
BOOST_CHECK((boost::math::isnan)(boost::math::gamma_p_derivative(static_cast<T>(-1), static_cast<T>(2))));
321359
BOOST_CHECK((boost::math::isnan)(boost::math::gamma_p_derivative(static_cast<T>(1), static_cast<T>(-2))));
322360
BOOST_CHECK((boost::math::isnan)(boost::math::gamma_p_derivative(static_cast<T>(0), static_cast<T>(2))));

‎test/test_lgamma_p_double.cu‎

Lines changed: 102 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,102 @@
1+
2+
// Copyright John Maddock 2016.
3+
// Copyright Matt Borland 2024.
4+
// Use, modification and distribution are subject to the
5+
// Boost Software License, Version 1.0. (See accompanying file
6+
// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
7+
8+
#define BOOST_MATH_PROMOTE_DOUBLE_POLICY false
9+
10+
#include <iostream>
11+
#include <iomanip>
12+
#include <vector>
13+
#include <boost/math/special_functions.hpp>
14+
#include "cuda_managed_ptr.hpp"
15+
#include "stopwatch.hpp"
16+
17+
// For the CUDA runtime routines (prefixed with "cuda_")
18+
#include <cuda_runtime.h>
19+
20+
typedef double float_type;
21+
22+
/**
23+
* CUDA Kernel Device code
24+
*
25+
*/
26+
__global__ void cuda_test(const float_type *in, float_type *out, int numElements)
27+
{
28+
using std::cos;
29+
int i = blockDim.x * blockIdx.x + threadIdx.x;
30+
31+
if (i < numElements)
32+
{
33+
out[i] = boost::math::lgamma_p(in[i], in[i]);
34+
}
35+
}
36+
37+
/**
38+
* Host main routine
39+
*/
40+
int main(void)
41+
{
42+
// Error code to check return values for CUDA calls
43+
cudaError_t err = cudaSuccess;
44+
45+
// Print the vector length to be used, and compute its size
46+
int numElements = 50000;
47+
std::cout << "[Vector operation on " << numElements << " elements]" << std::endl;
48+
49+
// Allocate the managed input vector A
50+
cuda_managed_ptr<float_type> input_vector(numElements);
51+
52+
// Allocate the managed output vector C
53+
cuda_managed_ptr<float_type> output_vector(numElements);
54+
55+
// Initialize the input vectors
56+
for (int i = 0; i < numElements; ++i)
57+
{
58+
input_vector[i] = rand()/(float_type)RAND_MAX;
59+
}
60+
61+
// Launch the Vector Add CUDA Kernel
62+
int threadsPerBlock = 1024;
63+
int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;
64+
std::cout << "CUDA kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock << " threads" << std::endl;
65+
66+
watch w;
67+
68+
cuda_test<<<blocksPerGrid, threadsPerBlock>>>(input_vector.get(), output_vector.get(), numElements);
69+
cudaDeviceSynchronize();
70+
71+
std::cout << "CUDA kernal done in: " << w.elapsed() << "s" << std::endl;
72+
73+
err = cudaGetLastError();
74+
75+
if (err != cudaSuccess)
76+
{
77+
std::cerr << "Failed to launch vectorAdd kernel (error code " << cudaGetErrorString(err) << ")!" << std::endl;
78+
return EXIT_FAILURE;
79+
}
80+
81+
// Verify that the result vector is correct
82+
std::vector<float_type> results;
83+
results.reserve(numElements);
84+
w.reset();
85+
for(int i = 0; i < numElements; ++i)
86+
results.push_back(boost::math::lgamma_p(input_vector[i], input_vector[i]));
87+
double t = w.elapsed();
88+
// check the results
89+
for(int i = 0; i < numElements; ++i)
90+
{
91+
if (boost::math::epsilon_difference(output_vector[i], results[i]) > 10)
92+
{
93+
std::cerr << "Result verification failed at element " << i << "!" << std::endl;
94+
return EXIT_FAILURE;
95+
}
96+
}
97+
98+
std::cout << "Test PASSED, normal calculation time: " << t << "s" << std::endl;
99+
std::cout << "Done\n";
100+
101+
return 0;
102+
}

0 commit comments

Comments
 (0)