-
Notifications
You must be signed in to change notification settings - Fork 13
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
- Loading branch information
1 parent
59e3eb3
commit 355b776
Showing
9 changed files
with
192 additions
and
30 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,61 @@ | ||
| /** \file sphericart_cuda.hpp | ||
| * Defines the CUDA API for `sphericart`. | ||
| */ | ||
|
|
||
| #ifndef SPHERICART_CUDA_HPP | ||
| #define SPHERICART_CUDA_HPP | ||
|
|
||
| #include "sphericart.hpp" | ||
|
|
||
| namespace sphericart { | ||
|
|
||
| namespace cuda { | ||
| /** | ||
| * A spherical harmonics calculator. | ||
| * | ||
| * It handles initialization of the prefactors upon initialization and it | ||
| * stores the buffers that are necessary to compute the spherical harmonics | ||
| * efficiently. | ||
| */ | ||
| template <typename T> class SphericalHarmonics { | ||
| public: | ||
| /** Initialize the SphericalHarmonics class setting maximum degree and | ||
| * normalization | ||
| * | ||
| * @param l_max | ||
| * The maximum degree of the spherical harmonics to be calculated. | ||
| * @param normalized | ||
| * If `false` (default) computes the scaled spherical harmonics, which | ||
| * are homogeneous polynomials in the Cartesian coordinates of the input | ||
| * points. If `true`, computes the normalized spherical harmonics that are | ||
| * evaluated on the unit sphere. In practice, this simply computes the | ||
| * scaled harmonics at the normalized coordinates \f$(x/r, y/r, z/r)\f$, and | ||
| * adapts the derivatives accordingly. | ||
| */ | ||
| SphericalHarmonics(size_t l_max, bool normalized = false); | ||
|
|
||
| /* @cond */ | ||
| ~SphericalHarmonics(); | ||
| /* @endcond */ | ||
|
|
||
| /** Computes the spherical harmonics for one or more 3D points, using | ||
| * pre-allocated device-side pointers | ||
| * | ||
| * @param xyz todo docs | ||
| * @param sph todo docs | ||
| */ | ||
| void compute(const T *xyz, size_t nsamples, bool compute_with_gradients, | ||
| bool compute_with_hessian, size_t GRID_DIM_X, | ||
| size_t GRID_DIM_Y, T *sph, T *dsph = nullptr, | ||
| T *ddsph = nullptr); | ||
|
|
||
| private: | ||
| size_t l_max; // maximum l value computed by this class | ||
| bool normalized; // should we normalize the input vectors? | ||
| T *prefactors_cpu; // host prefactors buffer | ||
| T *prefactors_cuda; // storage space for prefactors | ||
| }; | ||
|
|
||
| } // namespace cuda | ||
| } // namespace sphericart | ||
| #endif |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,88 @@ | ||
| #include <cuda.h> | ||
| #include <cuda_runtime.h> | ||
| #include <stdexcept> | ||
| #include <iostream> | ||
|
|
||
| #define _SPHERICART_INTERNAL_IMPLEMENTATION | ||
| #include "sphericart_cuda.hpp" | ||
| #include "cuda_base.hpp" | ||
|
|
||
|
|
||
| /*host macro that checks for errors in CUDA calls, and prints the file + line | ||
| * and error string if one occurs | ||
| */ | ||
| #define CUDA_CHECK(call) \ | ||
| do { \ | ||
| cudaError_t cudaStatus = (call); \ | ||
| if (cudaStatus != cudaSuccess) { \ | ||
| std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ \ | ||
| << " - " << cudaGetErrorString(cudaStatus) << std::endl; \ | ||
| cudaDeviceReset(); \ | ||
| exit(EXIT_FAILURE); \ | ||
| } \ | ||
| } while (0) | ||
|
|
||
| using namespace sphericart::cuda; | ||
|
|
||
| template <typename T> | ||
| SphericalHarmonics<T>::SphericalHarmonics(size_t l_max, bool normalized) { | ||
| /* | ||
| This is the constructor of the SphericalHarmonics class. It initizlizes | ||
| buffer space, compute prefactors, and sets the function pointers that are | ||
| used for the actual calls | ||
| */ | ||
|
|
||
| this->l_max = (int)l_max; | ||
| this->nprefactors = (int)(l_max + 1) * (l_max + 2); | ||
| this->normalized = normalized; | ||
| this->prefactors_cpu = new T[this->nprefactors]; | ||
|
|
||
| // compute prefactors on host first | ||
| compute_sph_prefactors<T>((int)l_max, this->prefactors_cpu); | ||
| // allocate them on device and copy to device | ||
| CUDA_CHECK( | ||
| cudaMalloc((void **)this->prefactors_cuda, this->nprefactors * sizeof(T))); | ||
| CUDA_CHECK(cudaMemcpy(this->prefactors_cpu, this->prefactors_cuda, | ||
| this->nprefactors * sizeof(T), | ||
| cudaMemcpyHostToDevice)); | ||
| } | ||
|
|
||
| template <typename T> SphericalHarmonics<T>::~SphericalHarmonics() { | ||
| // Destructor, frees the prefactors | ||
| delete[] (this->prefactors_cpu); | ||
| CUDA_CHECK(cudaFree(this->prefactors_cuda)); | ||
| } | ||
| template <typename T> | ||
| void SphericalHarmonics<T>::compute(const T *xyz, const size_t nsamples, | ||
| bool compute_with_gradients, | ||
| bool compute_with_hessian, | ||
| size_t GRID_DIM_X, size_t GRID_DIM_Y, | ||
| T *sph, T *dsph, | ||
| T *ddsph) { | ||
|
|
||
| if (sph == nullptr) { | ||
| throw std::runtime_error( | ||
| "sphericart::cuda::SphericalHarmonics::compute expected " | ||
| "sph ptr initialised, instead nullptr found. Initialise " | ||
| "sph with cudaMalloc."); | ||
| } | ||
|
|
||
| if (compute_with_gradients && dsph == nullptr) { | ||
| throw std::runtime_error( | ||
| "sphericart::cuda::SphericalHarmonics::compute expected " | ||
| "dsph != nullptr since compute_with_gradients = true. " | ||
| "initialise dsph with cudaMalloc."); | ||
| } | ||
|
|
||
| if (compute_with_hessian && ddsph == nullptr) { | ||
| throw std::runtime_error( | ||
| "sphericart::cuda::SphericalHarmonics::compute expected " | ||
| "ddsph != nullptr since compute_with_hessian = true. " | ||
| "initialise ddsph with cudaMalloc."); | ||
| } | ||
|
|
||
| sphericart::cuda::spherical_harmonics_cuda_base<T>( | ||
| xyz, nsamples, this->prefactors_cuda, this->nprefactors, this->l_max, | ||
| this->normalized, GRID_DIM_X, GRID_DIM_Y, compute_with_gradients, | ||
| compute_with_hessian, sph, dsph, ddsph); | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,19 @@ | ||
| #include "sphericart_cuda.hpp" | ||
| #include <cuda.h> | ||
| #include <cuda_runtime.h> | ||
|
|
||
| using namespace sphericart::cuda; | ||
|
|
||
| template <typename T> | ||
| SphericalHarmonics<T>::SphericalHarmonics(size_t l_max, bool normalized) {} | ||
|
|
||
| template <typename T> SphericalHarmonics<T>::~SphericalHarmonics() {} | ||
|
|
||
| void SphericalHarmonics<T>::compute(const T *xyz, const size_t nsamples, | ||
| bool compute_with_gradients, | ||
| bool compute_with_hessian, | ||
| size_t GRID_DIM_X, size_t GRID_DIM_Y, | ||
| T *sph, T *dsph = nullptr, | ||
| T *ddsph = nullptr) { | ||
| throw std::runtime_error("sphericart was not compiled with CUDA support"); | ||
| } |