Add proof of concept

[mborland]

Closes: #1
Based off @jzmaddock original PR here: boostorg/math#127

The rules seem to be more relaxed than in the original PR. assert and static_assert are allowed. Static variables are still not allowed, but I think that can be worked around using C++17's inline variables for global variables. The runner is using NVIDIA Tesla T4 (arch 70), and the CI provides version 12.5.0 of the NVIDIA toolkit which is the latest available. Locally I am using RHEL 9.4 with an RTX 3060 (arch 85) with the same version of the NVIDIA toolkit.

CC: @NAThompson, @ckormanyos

[jzmaddock]

Just checking... do we need these changes to the tests, or is it just a side effect of merging in an old branch?

[mborland]

I was having compilation issues with BOOST_MATH_CHECK_THROW but the version from Boost.Test worked just fine. Since the tests are already using Boost.test for everything else I figured it wasn't a big change.

[jzmaddock]

Thanks for this Matt!

Some random thoughts that have occurred to me:

• It's probably possible to build with b2 if we wanted to (though for a handful of CUDA tests it makes little difference), if we assume the compiler is command line compatible with say gcc, then a user-config.jam that looks like:

using gcc : some-cuda-name :  "some-path/nvcc"  : <cxxflags>-fspecial-cuda-option ;

would do the trick, same for msvc.

• It is possible to verify things build (but not run) on a regular github runner, see for example: https://github.com/boostorg/config/blob/5734e160e08b8df898c7f747000f27a3aafb7b2b/.github/workflows/ci.yml#L565
• I notice that OpenCL/Sycl now has a fancy new clang based compiler: https://www.intel.com/content/www/us/en/developer/tools/oneapi/data-parallel-c-plus-plus.html, the advantage of Sycl is that the code can be tested on a regular Intel Core processor, or via an emulator on a regular Github job (I believe this is what Boost.Compute are doing). I haven't gotten around to downloading Intel's Sycl implementation to try it yet though...
• Thinking out loud here, on reflection, assert's aren't really that useful in a device context - I would presume folks would want everything optimised anyway, plus what does an assert on the GPU actually do?

[mborland]

Thanks for this Matt!

Some random thoughts that have occurred to me:

• It's probably possible to build with b2 if we wanted to (though for a handful of CUDA tests it makes little difference), if we assume the compiler is command line compatible with say gcc, then a user-config.jam that looks like:

using gcc : some-cuda-name :  "some-path/nvcc"  : <cxxflags>-fspecial-cuda-option ;

would do the trick, same for msvc.

Could be worth a try. Peter also pointed me to: https://github.com/tee3/boost-build-nvcc.

• It is possible to verify things build (but not run) on a regular github runner, see for example: https://github.com/boostorg/config/blob/5734e160e08b8df898c7f747000f27a3aafb7b2b/.github/workflows/ci.yml#L565

I didn't realize you could do that without the hardware installed.

• I notice that OpenCL/Sycl now has a fancy new clang based compiler: https://www.intel.com/content/www/us/en/developer/tools/oneapi/data-parallel-c-plus-plus.html, the advantage of Sycl is that the code can be tested on a regular Intel Core processor, or via an emulator on a regular Github job (I believe this is what Boost.Compute are doing). I haven't gotten around to downloading Intel's Sycl implementation to try it yet though...

I have this installed on my local machine as well. ChatGPT says I should just be able to replace __host__ __device__ with SYCL_EXTERNAL when SYCL_LANGUAGE_VERSION is defined. I'll have to pull out the intel docs at some point to validate if that's true or not. Boost.Compute was OpenCL, but since their last commit is over 5 years ago it only supports obsolete versions which is unfortunate.

• Thinking out loud here, on reflection, assert's aren't really that useful in a device context - I would presume folks would want everything optimised anyway, plus what does an assert on the GPU actually do?

The CUDA threads are similar to CPU fibers so only the thread gets terminated not the entire process on an assertion. Since all of the assertions are using our own macro it would be easy to replace them with nothing if __CUDACC__ is defined.

[mborland]

I have the one test running fine locally with CUDA:

1: Test command: /home/mborland/Documents/boost/build/stage/bin/boost_cuda_math-test_arcsine_cdf_double
1: Working Directory: /home/mborland/Documents/boost/build/libs/cuda-math/test
1: Test timeout computed to be: 1500
1: [Vector operation on 50000 elements]
1: CUDA kernel launch with 98 blocks of 512 threads
1: CUDA kernal done in 0.020311s
1: Test PASSED with calculation time: 0.000759265s
1: Done
1: Failed to deinitialize the device! error=driver shutting down
1/1 Test #1: run-boost_cuda_math-test_arcsine_cdf_double ...   Passed    0.32 sec

I'll see if changing some macros around enables support for SYCL.

[mborland]

The beta function is now green with both SYCL and CUDA

[mborland]

Merging this since it's generally synchronized with the linked PR in math