mirror of
https://github.com/boostorg/math.git
synced 2025-05-11 21:33:52 +00:00
e4a01104d0
Fix igamma_large support on device Add GPU support to toms748 Add GPU support to igamma_inv Add GPU markers to gamma_inva Add GPU Markers to lgamma_small Remove STL usage from gamma Remove NVRTC workaround Fix fraction use of STL headers Mark gamma functions in fwd Disable declval on all GPU platforms Disable more unneeded code on device Add forward decl for NVRTC tgamma Disable unneeded items for all GPU Change workaround for missing overloads Rearrange definition location Add include path to cuda now that workaround is removed Fix NVRTC incompatibility with recursion and forward decls Add tgamma_ratio CUDA and NVRTC testing Fix NVRTC handling of gamma_p_derivative Add gamma_p_derivative CUDA and NVRTC testing Remove recursion from gamma_incomplete_imp Add SYCL testing of igamma, igamma_inv, and igamma_inva Ignore literal-range warnings Remove use of static const char* for function name Fix missing CUDA header Remove calls under NVRTC to fwd decl Add more nvrtc workarounds Use builtin erfc instead of header cycle Add CUDA and NVRTC testing of gamma_p_inv Adjust tolerances Add GPU support to chi squared dist Fix static local variable Add chi squared dist SYCL testing Add chi squared dist CUDA testing Add chi squared dist NVRTC testing Add GPU support to weibull dist Add weibull dist SYCL testing Add weibull dist CUDA testing Add weibull dist NVRTC testing
108 lines
3.1 KiB
Plaintext
108 lines
3.1 KiB
Plaintext
|
|
// Copyright John Maddock 2016.
|
|
// Copyright Matt Borland 2024.
|
|
// Use, modification and distribution are subject to the
|
|
// Boost Software License, Version 1.0. (See accompanying file
|
|
// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
|
|
|
|
#include <iostream>
|
|
#include <iomanip>
|
|
#include <vector>
|
|
#include <random>
|
|
#include <boost/math/special_functions.hpp>
|
|
#include "cuda_managed_ptr.hpp"
|
|
#include "stopwatch.hpp"
|
|
|
|
// For the CUDA runtime routines (prefixed with "cuda_")
|
|
#include <cuda_runtime.h>
|
|
|
|
typedef float float_type;
|
|
|
|
/**
|
|
* CUDA Kernel Device code
|
|
*
|
|
*/
|
|
__global__ void cuda_test(const float_type *in1, const float_type *in2, float_type *out, int numElements)
|
|
{
|
|
using std::cos;
|
|
int i = blockDim.x * blockIdx.x + threadIdx.x;
|
|
|
|
if (i < numElements)
|
|
{
|
|
out[i] = boost::math::gamma_p_inv(in1[i], in2[i]);
|
|
}
|
|
}
|
|
|
|
/**
|
|
* Host main routine
|
|
*/
|
|
int main(void)
|
|
{
|
|
// Error code to check return values for CUDA calls
|
|
cudaError_t err = cudaSuccess;
|
|
|
|
// Print the vector length to be used, and compute its size
|
|
int numElements = 50000;
|
|
std::cout << "[Vector operation on " << numElements << " elements]" << std::endl;
|
|
|
|
// Allocate the managed input vector A
|
|
cuda_managed_ptr<float_type> input_vector1(numElements);
|
|
|
|
// Allocate the managed input vector B
|
|
cuda_managed_ptr<float_type> input_vector2(numElements);
|
|
|
|
// Allocate the managed output vector C
|
|
cuda_managed_ptr<float_type> output_vector(numElements);
|
|
|
|
// Initialize the input vectors
|
|
std::mt19937_64 gen(42);
|
|
std::uniform_real_distribution<float_type> dist(0, 1);
|
|
for (int i = 0; i < numElements; ++i)
|
|
{
|
|
input_vector1[i] = dist(gen);
|
|
input_vector2[i] = dist(gen);
|
|
}
|
|
|
|
// Launch the Vector Add CUDA Kernel
|
|
int threadsPerBlock = 256;
|
|
int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;
|
|
std::cout << "CUDA kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock << " threads" << std::endl;
|
|
|
|
watch w;
|
|
|
|
cuda_test<<<blocksPerGrid, threadsPerBlock>>>(input_vector1.get(), input_vector2.get(), output_vector.get(), numElements);
|
|
cudaDeviceSynchronize();
|
|
|
|
std::cout << "CUDA kernal done in: " << w.elapsed() << "s" << std::endl;
|
|
|
|
err = cudaGetLastError();
|
|
|
|
if (err != cudaSuccess)
|
|
{
|
|
std::cerr << "Failed to launch CUDA kernel (error code " << cudaGetErrorString(err) << ")!" << std::endl;
|
|
return EXIT_FAILURE;
|
|
}
|
|
|
|
// Verify that the result vector is correct
|
|
std::vector<float_type> results;
|
|
results.reserve(numElements);
|
|
w.reset();
|
|
for(int i = 0; i < numElements; ++i)
|
|
results.push_back(boost::math::gamma_p_inv(input_vector1[i], input_vector2[i]));
|
|
double t = w.elapsed();
|
|
// check the results
|
|
for(int i = 0; i < numElements; ++i)
|
|
{
|
|
if (boost::math::epsilon_difference(output_vector[i], results[i]) > 100)
|
|
{
|
|
std::cerr << "Result verification failed at element " << i << "!" << std::endl;
|
|
return EXIT_FAILURE;
|
|
}
|
|
}
|
|
|
|
std::cout << "Test PASSED, normal calculation time: " << t << "s" << std::endl;
|
|
std::cout << "Done\n";
|
|
|
|
return 0;
|
|
}
|