mirror of
https://github.com/boostorg/math.git
synced 2025-05-11 21:33:52 +00:00
1324401c92
Add SYCL testing of ariy functions Add CUDA testing of airy functions Add NVRTC testing of airy functions Add GPU support to ellint rc Add GPU support to ellint rd Add GPU support to ellint rf Add GPU support to ellint rg Add GPU support to ellint rj Add GPU support to ellint d Add GPU support to ellint_1 Markup forward and add ellint_3 return type def for NVRTC platform Add CUDA testing of ellint 1 NVRTC fixes Add NVRTC testing of ellint_1 Add GPU support to ellint_2 Add CUDA testing of ellint_2 Fix NVRTC errors Add NVRTC testing of ellint_2 Add GPU support to atanh Add GPU support to ellint_3 Add NVRTC testing of ellint_3 Add CUDA testing of ellint_3 Replace use of static const char* Add SYCL testing of ellint_1 Add SYCL testing of ellint 2 with slight tolerance bump Remove recursion from ellint_rj Add ellint_d CUDA testing Add NVRTC testing of ellint_d Add SYCL testing of ellint_d Remove SYCL ellint_3 support Update docs Add GPU support to jacobi zeta Add CUDA testing of jacobi zeta Add NVRTC testing of jacobi zeta Add SYCL testing of jacobi zeta Add GPU support to heuman_lambda Add NVRTC testing of heuman lambda Add CUDA testing of heuman_lambda Add SYCL testing of heuman lambda Add markers to docs Add marker for CUDA only functions in the docs
101 lines
2.8 KiB
Plaintext
101 lines
2.8 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 <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 *in, float_type *out, int numElements)
|
|
{
|
|
using std::cos;
|
|
int i = blockDim.x * blockIdx.x + threadIdx.x;
|
|
|
|
if (i < numElements)
|
|
{
|
|
out[i] = boost::math::airy_ai(in[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_vector(numElements);
|
|
|
|
// Allocate the managed output vector C
|
|
cuda_managed_ptr<float_type> output_vector(numElements);
|
|
|
|
// Initialize the input vectors
|
|
for (int i = 0; i < numElements; ++i)
|
|
{
|
|
input_vector[i] = rand()/(float_type)RAND_MAX;
|
|
}
|
|
|
|
// 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_vector.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 vectorAdd 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::airy_ai(input_vector[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]) > 10)
|
|
{
|
|
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;
|
|
}
|