diff --git a/cpp/daal/include/algorithms/engines/mrg32k3a/mrg32k3a.h b/cpp/daal/include/algorithms/engines/mrg32k3a/mrg32k3a.h new file mode 100644 index 00000000000..a70c1853e1a --- /dev/null +++ b/cpp/daal/include/algorithms/engines/mrg32k3a/mrg32k3a.h @@ -0,0 +1,183 @@ +/* file: mrg32k3a.h */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +/* +//++ +// Implementation of the MRG32k3a engine: a 32-bit combined multiple recursive generator +// with two components of order 3, optimized for batch processing. +//-- +*/ + +#ifndef __MRG32K3A_H__ +#define __MRG32K3A_H__ + +#include "algorithms/engines/mrg32k3a/mrg32k3a_types.h" +#include "algorithms/engines/engine.h" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace mrg32k3a +{ +/** + * @defgroup engines_mrg32k3a_batch Batch + * @ingroup engines_mrg32k3a + * @{ + */ +namespace interface1 +{ +/** + * + * \brief Provides methods to run implementations of the mrg32k3a engine. + * This class is associated with the \ref mrg32k3a::interface1::Batch "mrg32k3a::Batch" class + * and supports the method of mrg32k3a engine computation in the batch processing mode + * + * \tparam algorithmFPType Data type to use in intermediate computations of mrg32k3a engine, double or float + * \tparam method Computation method of the engine, mrg32k3a::Method + * \tparam cpu Version of the cpu-specific implementation of the engine, daal::CpuType + */ +template +class BatchContainer : public daal::algorithms::AnalysisContainerIface +{ +public: + /** + * Constructs a container for the mrg32k3a engine with a specified environment + * in the batch processing mode + * \param[in] daalEnv Environment object + */ + BatchContainer(daal::services::Environment::env * daalEnv); + ~BatchContainer(); + /** + * Computes the result of the mrg32k3a engine in the batch processing mode + * + * \return Status of computations + */ + services::Status compute() DAAL_C11_OVERRIDE; +}; + +/** + * + * \brief Provides methods for mrg32k3a engine computations in the batch processing mode + * + * \tparam algorithmFPType Data type to use in intermediate computations of mrg32k3a engine, double or float + * \tparam method Computation method of the engine, mrg32k3a::Method + * + * \par Enumerations + * - mrg32k3a::Method Computation methods for the mrg32k3a engine + * + * \par References + * - \ref engines::interface1::Input "engines::Input" class + * - \ref engines::interface1::Result "engines::Result" class + */ +template +class DAAL_EXPORT Batch : public engines::BatchBase +{ +public: + typedef engines::BatchBase super; + + typedef typename super::InputType InputType; + typedef typename super::ResultType ResultType; + + /** + * Creates mrg32k3a engine + * \param[in] seed Initial condition for mrg32k3a engine + * + * \return Pointer to mrg32k3a engine + */ + static services::SharedPtr > create(size_t seed = 777); + + /** + * Returns method of the engine + * \return Method of the engine + */ + virtual int getMethod() const DAAL_C11_OVERRIDE { return (int)method; } + + /** + * Returns the structure that contains results of mrg32k3a engine + * \return Structure that contains results of mrg32k3a engine + */ + ResultPtr getResult() { return _result; } + + /** + * Registers user-allocated memory to store results of mrg32k3a engine + * \param[in] result Structure to store results of mrg32k3a engine + * + * \return Status of computations + */ + services::Status setResult(const ResultPtr & result) + { + DAAL_CHECK(result, services::ErrorNullResult) + _result = result; + _res = _result.get(); + return services::Status(); + } + + /** + * Returns a pointer to the newly allocated mrg32k3a engine + * with a copy of input objects and parameters of this mrg32k3a engine + * \return Pointer to the newly allocated engine + */ + services::SharedPtr > clone() const { return services::SharedPtr >(cloneImpl()); } + + /** + * Allocates memory to store the result of the mrg32k3a engine + * + * \return Status of computations + */ + virtual services::Status allocateResult() DAAL_C11_OVERRIDE + { + services::Status s = this->_result->template allocate(&(this->input), NULL, (int)method); + this->_res = this->_result.get(); + return s; + } + +protected: + Batch(size_t seed = 777) { initialize(); } + + Batch(const Batch & other) : super(other) { initialize(); } + + virtual Batch * cloneImpl() const DAAL_C11_OVERRIDE { return new Batch(*this); } + + void initialize() + { + Analysis::_ac = new __DAAL_ALGORITHM_CONTAINER(batch, BatchContainer, algorithmFPType, method)(&_env); + _in = &input; + _result.reset(new ResultType()); + } + +private: + ResultPtr _result; + + Batch & operator=(const Batch &); +}; +typedef services::SharedPtr > mrg32k3aPtr; +typedef services::SharedPtr > mrg32k3aConstPtr; + +} // namespace interface1 +using interface1::BatchContainer; +using interface1::Batch; +using interface1::mrg32k3aPtr; +using interface1::mrg32k3aConstPtr; +/** @} */ +} // namespace mrg32k3a +} // namespace engines +} // namespace algorithms +} // namespace daal +#endif diff --git a/cpp/daal/include/algorithms/engines/mrg32k3a/mrg32k3a_types.h b/cpp/daal/include/algorithms/engines/mrg32k3a/mrg32k3a_types.h new file mode 100644 index 00000000000..8fdc58b98c8 --- /dev/null +++ b/cpp/daal/include/algorithms/engines/mrg32k3a/mrg32k3a_types.h @@ -0,0 +1,65 @@ +/* file: mrg32k3a_types.h */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +/* +//++ +// Implementation of the MRG32k3a engine: a 32-bit combined multiple recursive generator +// with two components of order 3, optimized for batch processing. +//-- +*/ + +#ifndef __MRG32K3A_TYPES_H__ +#define __MRG32K3A_TYPES_H__ + +#include "algorithms/algorithm.h" +#include "services/daal_defines.h" +#include "data_management/data/numeric_table.h" +#include "data_management/data/homogen_numeric_table.h" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +/** + * @defgroup engines_mrg32k3a mrg32k3a Engine + * \copydoc daal::algorithms::engines::mrg32k3a + * @ingroup engines + * @{ + */ +/** + * \brief Contains classes for mrg32k3a engine + */ +namespace mrg32k3a +{ +/** + * + * Available methods to compute mrg32k3a engine + */ +enum Method +{ + defaultDense = 0 /*!< Default: performance-oriented method. */ +}; + +} // namespace mrg32k3a +/** @} */ +} // namespace engines +} // namespace algorithms +} // namespace daal + +#endif diff --git a/cpp/daal/include/algorithms/engines/philox4x32x10/philox4x32x10.h b/cpp/daal/include/algorithms/engines/philox4x32x10/philox4x32x10.h new file mode 100644 index 00000000000..3a5d0e33180 --- /dev/null +++ b/cpp/daal/include/algorithms/engines/philox4x32x10/philox4x32x10.h @@ -0,0 +1,183 @@ +/* file: philox4x32x10.h */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +/* +//++ +// Implementation of the Philox4x32-10 engine: a counter-based pseudorandom number generator (PRNG) +// that uses 4x32-bit keys and performs 10 rounds of mixing to produce high-quality randomness. +//-- +*/ + +#ifndef __PHILOX4X32X10_H__ +#define __PHILOX4X32X10_H__ + +#include "algorithms/engines/philox4x32x10/philox4x32x10_types.h" +#include "algorithms/engines/engine.h" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace philox4x32x10 +{ +/** + * @defgroup engines_philox4x32x10_batch Batch + * @ingroup engines_philox4x32x10 + * @{ + */ +namespace interface1 +{ +/** + * + * \brief Provides methods to run implementations of the philox4x32x10 engine. + * This class is associated with the \ref philox4x32x10::interface1::Batch "philox4x32x10::Batch" class + * and supports the method of philox4x32x10 engine computation in the batch processing mode + * + * \tparam algorithmFPType Data type to use in intermediate computations of philox4x32x10 engine, double or float + * \tparam method Computation method of the engine, philox4x32x10::Method + * \tparam cpu Version of the cpu-specific implementation of the engine, daal::CpuType + */ +template +class BatchContainer : public daal::algorithms::AnalysisContainerIface +{ +public: + /** + * Constructs a container for the philox4x32x10 engine with a specified environment + * in the batch processing mode + * \param[in] daalEnv Environment object + */ + BatchContainer(daal::services::Environment::env * daalEnv); + ~BatchContainer(); + /** + * Computes the result of the philox4x32x10 engine in the batch processing mode + * + * \return Status of computations + */ + services::Status compute() DAAL_C11_OVERRIDE; +}; + +/** + * + * \brief Provides methods for philox4x32x10 engine computations in the batch processing mode + * + * \tparam algorithmFPType Data type to use in intermediate computations of philox4x32x10 engine, double or float + * \tparam method Computation method of the engine, philox4x32x10::Method + * + * \par Enumerations + * - philox4x32x10::Method Computation methods for the philox4x32x10 engine + * + * \par References + * - \ref engines::interface1::Input "engines::Input" class + * - \ref engines::interface1::Result "engines::Result" class + */ +template +class DAAL_EXPORT Batch : public engines::BatchBase +{ +public: + typedef engines::BatchBase super; + + typedef typename super::InputType InputType; + typedef typename super::ResultType ResultType; + + /** + * Creates philox4x32x10 engine + * \param[in] seed Initial condition for philox4x32x10 engine + * + * \return Pointer to philox4x32x10 engine + */ + static services::SharedPtr > create(size_t seed = 777); + + /** + * Returns method of the engine + * \return Method of the engine + */ + virtual int getMethod() const DAAL_C11_OVERRIDE { return (int)method; } + + /** + * Returns the structure that contains results of philox4x32x10 engine + * \return Structure that contains results of philox4x32x10 engine + */ + ResultPtr getResult() { return _result; } + + /** + * Registers user-allocated memory to store results of philox4x32x10 engine + * \param[in] result Structure to store results of philox4x32x10 engine + * + * \return Status of computations + */ + services::Status setResult(const ResultPtr & result) + { + DAAL_CHECK(result, services::ErrorNullResult) + _result = result; + _res = _result.get(); + return services::Status(); + } + + /** + * Returns a pointer to the newly allocated philox4x32x10 engine + * with a copy of input objects and parameters of this philox4x32x10 engine + * \return Pointer to the newly allocated engine + */ + services::SharedPtr > clone() const { return services::SharedPtr >(cloneImpl()); } + + /** + * Allocates memory to store the result of the philox4x32x10 engine + * + * \return Status of computations + */ + virtual services::Status allocateResult() DAAL_C11_OVERRIDE + { + services::Status s = this->_result->template allocate(&(this->input), NULL, (int)method); + this->_res = this->_result.get(); + return s; + } + +protected: + Batch(size_t seed = 777) { initialize(); } + + Batch(const Batch & other) : super(other) { initialize(); } + + virtual Batch * cloneImpl() const DAAL_C11_OVERRIDE { return new Batch(*this); } + + void initialize() + { + Analysis::_ac = new __DAAL_ALGORITHM_CONTAINER(batch, BatchContainer, algorithmFPType, method)(&_env); + _in = &input; + _result.reset(new ResultType()); + } + +private: + ResultPtr _result; + + Batch & operator=(const Batch &); +}; +typedef services::SharedPtr > philox4x32x10Ptr; +typedef services::SharedPtr > philox4x32x10ConstPtr; + +} // namespace interface1 +using interface1::BatchContainer; +using interface1::Batch; +using interface1::philox4x32x10Ptr; +using interface1::philox4x32x10ConstPtr; +/** @} */ +} // namespace philox4x32x10 +} // namespace engines +} // namespace algorithms +} // namespace daal +#endif diff --git a/cpp/daal/include/algorithms/engines/philox4x32x10/philox4x32x10_types.h b/cpp/daal/include/algorithms/engines/philox4x32x10/philox4x32x10_types.h new file mode 100644 index 00000000000..0c0a92c9b3a --- /dev/null +++ b/cpp/daal/include/algorithms/engines/philox4x32x10/philox4x32x10_types.h @@ -0,0 +1,65 @@ +/* file: philox4x32x10_types.h */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +/* +//++ +// Implementation of the Philox4x32-10 engine: a counter-based pseudorandom number generator (PRNG) +// that uses 4x32-bit keys and performs 10 rounds of mixing to produce high-quality randomness. +//-- +*/ + +#ifndef __PHILOX4X32X10_TYPES_H__ +#define __PHILOX4X32X10_TYPES_H__ + +#include "algorithms/algorithm.h" +#include "services/daal_defines.h" +#include "data_management/data/numeric_table.h" +#include "data_management/data/homogen_numeric_table.h" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +/** + * @defgroup engines_philox4x32x10 philox4x32x10 Engine + * \copydoc daal::algorithms::engines::philox4x32x10 + * @ingroup engines + * @{ + */ +/** + * \brief Contains classes for philox4x32x10 engine + */ +namespace philox4x32x10 +{ +/** + * + * Available methods to compute philox4x32x10 engine + */ +enum Method +{ + defaultDense = 0 /*!< Default: performance-oriented method. */ +}; + +} // namespace philox4x32x10 +/** @} */ +} // namespace engines +} // namespace algorithms +} // namespace daal + +#endif diff --git a/cpp/daal/include/daal.h b/cpp/daal/include/daal.h index 881a6c39fbe..9e91570ac56 100755 --- a/cpp/daal/include/daal.h +++ b/cpp/daal/include/daal.h @@ -304,6 +304,10 @@ #include "algorithms/engines/engine_family.h" #include "algorithms/engines/mt2203/mt2203.h" #include "algorithms/engines/mt2203/mt2203_types.h" +#include "algorithms/engines/mrg32k3a/mrg32k3a.h" +#include "algorithms/engines/mrg32k3a/mrg32k3a_types.h" +#include "algorithms/engines/philox4x32x10/philox4x32x10.h" +#include "algorithms/engines/philox4x32x10/philox4x32x10_types.h" #include "algorithms/dbscan/dbscan_types.h" #include "algorithms/dbscan/dbscan_batch.h" #include "algorithms/dbscan/dbscan_distributed.h" diff --git a/cpp/daal/include/daal_win.h b/cpp/daal/include/daal_win.h index e17eff16796..6e86076e275 100755 --- a/cpp/daal/include/daal_win.h +++ b/cpp/daal/include/daal_win.h @@ -316,6 +316,10 @@ #include "algorithms/engines/engine_family.h" #include "algorithms/engines/mt2203/mt2203.h" #include "algorithms/engines/mt2203/mt2203_types.h" +#include "algorithms/engines/mrg32k3a/mrg32k3a.h" +#include "algorithms/engines/mrg32k3a/mrg32k3a_types.h" +#include "algorithms/engines/philox4x32x10/philox4x32x10.h" +#include "algorithms/engines/philox4x32x10/philox4x32x10_types.h" #include "algorithms/dbscan/dbscan_types.h" #include "algorithms/dbscan/dbscan_batch.h" #include "algorithms/dbscan/dbscan_distributed.h" diff --git a/cpp/daal/src/algorithms/engines/mcg59/mcg59_batch_impl.h b/cpp/daal/src/algorithms/engines/mcg59/mcg59_batch_impl.h index 6c3040da615..0a9933b57dd 100644 --- a/cpp/daal/src/algorithms/engines/mcg59/mcg59_batch_impl.h +++ b/cpp/daal/src/algorithms/engines/mcg59/mcg59_batch_impl.h @@ -26,9 +26,6 @@ #include "src/externals/service_rng.h" #include "src/data_management/service_numeric_table.h" -static const int leapfrogMethodErrcode = -1002; -static const int skipAheadMethodErrcode = -1003; - namespace daal { namespace algorithms @@ -67,7 +64,7 @@ class BatchImpl : public algorithms::engines::mcg59::interface1::Batch +SharedPtr > Batch::create(size_t seed) +{ + SharedPtr > engPtr; + +#define DAAL_CREATE_ENGINE_CPU(cpuId, ...) engPtr.reset(new BatchImpl(__VA_ARGS__)); + + DAAL_DISPATCH_FUNCTION_BY_CPU(DAAL_CREATE_ENGINE_CPU, seed); + +#undef DAAL_CREATE_ENGINE_CPU + return engPtr; +} + +template class Batch; +template class Batch; + +} // namespace interface1 +} // namespace mrg32k3a +} // namespace engines +} // namespace algorithms +} // namespace daal diff --git a/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_batch_container.h b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_batch_container.h new file mode 100644 index 00000000000..c925b3c4483 --- /dev/null +++ b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_batch_container.h @@ -0,0 +1,68 @@ +/* file: mrg32k3a_batch_container.h */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +/* +//++ +// Implementation of mrg32k3a calculation algorithm container. +//-- +*/ + +#ifndef __MRG32K3A_BATCH_CONTAINER_H__ +#define __MRG32K3A_BATCH_CONTAINER_H__ + +#include "algorithms/engines/mrg32k3a/mrg32k3a.h" +#include "src/algorithms/engines/mrg32k3a/mrg32k3a_kernel.h" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace mrg32k3a +{ +namespace interface1 +{ +template +BatchContainer::BatchContainer(daal::services::Environment::env * daalEnv) : AnalysisContainerIface(daalEnv) +{ + __DAAL_INITIALIZE_KERNELS(internal::mrg32k3aKernel, algorithmFPType, method); +} + +template +BatchContainer::~BatchContainer() +{ + __DAAL_DEINITIALIZE_KERNELS(); +} + +template +services::Status BatchContainer::compute() +{ + daal::services::Environment::env & env = *_env; + engines::Result * result = static_cast(_res); + NumericTable * resultTable = result->get(engines::randomNumbers).get(); + + __DAAL_CALL_KERNEL(env, internal::mrg32k3aKernel, __DAAL_KERNEL_ARGUMENTS(algorithmFPType, method), compute, resultTable); +} + +} // namespace interface1 +} // namespace mrg32k3a +} // namespace engines +} // namespace algorithms +} // namespace daal + +#endif diff --git a/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_batch_impl.h b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_batch_impl.h new file mode 100644 index 00000000000..0b7048fa95a --- /dev/null +++ b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_batch_impl.h @@ -0,0 +1,114 @@ +/* file: mrg32k3a_batch_impl.h */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +/* +//++ +// Implementation of the class defining the mrg32k3a engine. +//-- +*/ + +#include "algorithms/engines/mrg32k3a/mrg32k3a.h" +#include "src/algorithms/engines/engine_batch_impl.h" +#include "src/externals/service_rng.h" +#include "src/data_management/service_numeric_table.h" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace mrg32k3a +{ +namespace internal +{ +template +class BatchImpl : public algorithms::engines::mrg32k3a::interface1::Batch, + public algorithms::engines::internal::BatchBaseImpl +{ +public: + typedef algorithms::engines::mrg32k3a::interface1::Batch super1; + typedef algorithms::engines::internal::BatchBaseImpl super2; + BatchImpl(size_t seed = 777) : baseRng(seed, __DAAL_BRNG_MRG32K3A), super2(seed) {} + + void * getState() DAAL_C11_OVERRIDE { return baseRng.getState(); } + + int getStateSize() const DAAL_C11_OVERRIDE { return baseRng.getStateSize(); } + + services::Status saveStateImpl(byte * dest) const DAAL_C11_OVERRIDE + { + DAAL_CHECK(!baseRng.saveState((void *)dest), ErrorIncorrectErrorcodeFromGenerator); + return services::Status(); + } + + services::Status loadStateImpl(const byte * src) DAAL_C11_OVERRIDE + { + DAAL_CHECK(!baseRng.loadState((const void *)src), ErrorIncorrectErrorcodeFromGenerator); + return services::Status(); + } + + services::Status leapfrogImpl(size_t threadNum, size_t nThreads) DAAL_C11_OVERRIDE + { + int errcode = baseRng.leapfrog(threadNum, nThreads); + services::Status s; + if (errcode == __DAAL_LEAPFROG_METHOD_ERRCODE) + s.add(ErrorLeapfrogUnsupported); + else if (errcode) + s.add(ErrorIncorrectErrorcodeFromGenerator); + return s; + } + + services::Status skipAheadImpl(size_t nSkip) DAAL_C11_OVERRIDE + { + int errcode = baseRng.skipAhead(nSkip); + services::Status s; + if (errcode == __DAAL_SKIP_AHEAD_METHOD_ERRCODE) + s.add(ErrorSkipAheadUnsupported); + else if (errcode) + s.add(ErrorIncorrectErrorcodeFromGenerator); + return s; + } + + virtual BatchImpl * cloneImpl() const DAAL_C11_OVERRIDE + { + return new BatchImpl(*this); + } + + bool hasSupport(engines::internal::ParallelizationTechnique technique) const DAAL_C11_OVERRIDE + { + switch (technique) + { + case engines::internal::family: return false; + case engines::internal::skipahead: return true; + case engines::internal::leapfrog: return false; + } + return false; + } + + ~BatchImpl() {} + +protected: + BatchImpl(const BatchImpl & other) : super1(other), super2(other), baseRng(other.baseRng) {} + + daal::internal::BaseRNGsInst baseRng; +}; + +} // namespace internal +} // namespace mrg32k3a +} // namespace engines +} // namespace algorithms +} // namespace daal diff --git a/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_dense_default_batch_fpt_cpu.cpp b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_dense_default_batch_fpt_cpu.cpp new file mode 100644 index 00000000000..529c4af2635 --- /dev/null +++ b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_dense_default_batch_fpt_cpu.cpp @@ -0,0 +1,47 @@ +/* file: mrg32k3a_dense_default_batch_fpt_cpu.cpp */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +//++ +// Implementation of mrg32k3a calculation functions. +//-- + +#include "src/algorithms/engines/mrg32k3a/mrg32k3a_batch_container.h" +#include "src/algorithms/engines/mrg32k3a/mrg32k3a_kernel.h" +#include "src/algorithms/engines/mrg32k3a/mrg32k3a_impl.i" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace mrg32k3a +{ +namespace interface1 +{ +template class BatchContainer; +} // namespace interface1 + +namespace internal +{ +template class mrg32k3aKernel; +} // namespace internal + +} // namespace mrg32k3a +} // namespace engines +} // namespace algorithms +} // namespace daal diff --git a/cpp/oneapi/dal/backend/primitives/rng/partial_shuffle.hpp b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_dense_default_batch_fpt_dispatcher.cpp similarity index 60% rename from cpp/oneapi/dal/backend/primitives/rng/partial_shuffle.hpp rename to cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_dense_default_batch_fpt_dispatcher.cpp index 6861a44a3f7..fd78108df73 100644 --- a/cpp/oneapi/dal/backend/primitives/rng/partial_shuffle.hpp +++ b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_dense_default_batch_fpt_dispatcher.cpp @@ -1,5 +1,6 @@ +/* file: mrg32k3a_dense_default_batch_fpt_dispatcher.cpp */ /******************************************************************************* -* Copyright 2021 Intel Corporation +* Copyright contributors to the oneDAL project * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,15 +15,16 @@ * limitations under the License. *******************************************************************************/ -#pragma once +//++ +// Implementation of mrg32k3a calculation algorithm dispatcher. +//-- -#include "oneapi/dal/backend/primitives/ndarray.hpp" +#include "src/algorithms/engines/mrg32k3a/mrg32k3a_batch_container.h" -namespace oneapi::dal::backend::primitives { - -void partial_fisher_yates_shuffle(ndview& result_array, - std::int64_t top, - std::int64_t seed); -void partial_fisher_yates_shuffle(ndview& result_array, std::int64_t top); - -} // namespace oneapi::dal::backend::primitives +namespace daal +{ +namespace algorithms +{ +__DAAL_INSTANTIATE_DISPATCH_CONTAINER(engines::mrg32k3a::BatchContainer, batch, DAAL_FPTYPE, engines::mrg32k3a::defaultDense) +} // namespace algorithms +} // namespace daal diff --git a/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_impl.i b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_impl.i new file mode 100644 index 00000000000..f8f12b2deea --- /dev/null +++ b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_impl.i @@ -0,0 +1,49 @@ +/* file: mrg32k3a_impl.i */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +/* +//++ +// Implementation of mrg32k3a algorithm. +//-- +*/ + +#ifndef __MRG32K3A_IMPL_I__ +#define __MRG32K3A_IMPL_I__ + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace mrg32k3a +{ +namespace internal +{ +template +Status mrg32k3aKernel::compute(NumericTable * resultTensor) +{ + return Status(); +} + +} // namespace internal +} // namespace mrg32k3a +} // namespace engines +} // namespace algorithms +} // namespace daal + +#endif diff --git a/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_kernel.h b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_kernel.h new file mode 100644 index 00000000000..80c9fbe44d9 --- /dev/null +++ b/cpp/daal/src/algorithms/engines/mrg32k3a/mrg32k3a_kernel.h @@ -0,0 +1,58 @@ +/* file: mrg32k3a_kernel.h */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +//++ +// Declaration of a template function for calculating values using the MRG32k3a generator. +//-- + +#ifndef __MRG32K3A_KERNEL_H__ +#define __MRG32K3A_KERNEL_H__ + +#include "algorithms/engines/mrg32k3a/mrg32k3a.h" +#include "src/algorithms/kernel.h" +#include "data_management/data/numeric_table.h" + +using namespace daal::services; +using namespace daal::data_management; + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace mrg32k3a +{ +namespace internal +{ +/** + * \brief Kernel for mrg32k3a calculation + */ +template +class mrg32k3aKernel : public Kernel +{ +public: + Status compute(NumericTable * resultTable); +}; + +} // namespace internal +} // namespace mrg32k3a +} // namespace engines +} // namespace algorithms +} // namespace daal + +#endif diff --git a/cpp/daal/src/algorithms/engines/mt19937/mt19937_batch_impl.h b/cpp/daal/src/algorithms/engines/mt19937/mt19937_batch_impl.h index e92d0e46612..b5959d5bfe6 100644 --- a/cpp/daal/src/algorithms/engines/mt19937/mt19937_batch_impl.h +++ b/cpp/daal/src/algorithms/engines/mt19937/mt19937_batch_impl.h @@ -26,9 +26,6 @@ #include "src/externals/service_rng.h" #include "src/data_management/service_numeric_table.h" -static const int leapfrogMethodErrcode = -1002; -static const int skipAheadMethodErrcode = -1003; - namespace daal { namespace algorithms @@ -67,7 +64,7 @@ class BatchImpl : public algorithms::engines::mt19937::interface1::Batchleapfrog(threadNum, nThreads); services::Status s; - if (errcode == leapfrogMethodErrcode) + if (errcode == __DAAL_LEAPFROG_METHOD_ERRCODE) s.add(ErrorLeapfrogUnsupported); else if (errcode) s.add(ErrorIncorrectErrorcodeFromGenerator); @@ -199,7 +196,7 @@ class BatchImpl : public algorithms::engines::mt2203::interface1::BatchskipAhead(nSkip); services::Status s; - if (errcode == skipAheadMethodErrcode) + if (errcode == __DAAL_SKIP_AHEAD_METHOD_ERRCODE) s.add(ErrorSkipAheadUnsupported); else if (errcode) s.add(ErrorIncorrectErrorcodeFromGenerator); diff --git a/cpp/daal/src/algorithms/engines/mt2203/mt2203_kernel.h b/cpp/daal/src/algorithms/engines/mt2203/mt2203_kernel.h index b7de119367f..e588a02c8fb 100644 --- a/cpp/daal/src/algorithms/engines/mt2203/mt2203_kernel.h +++ b/cpp/daal/src/algorithms/engines/mt2203/mt2203_kernel.h @@ -19,8 +19,8 @@ // Declaration of template function that calculate mt2203s. //-- -#ifndef __MCG59_KERNEL_H__ -#define __MCG59_KERNEL_H__ +#ifndef __MT2203_KERNEL_H__ +#define __MT2203_KERNEL_H__ #include "algorithms/engines/mt2203/mt2203.h" #include "src/algorithms/kernel.h" diff --git a/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10.cpp b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10.cpp new file mode 100644 index 00000000000..47fb7dae70f --- /dev/null +++ b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10.cpp @@ -0,0 +1,59 @@ +/* file: philox4x32x10.cpp */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +//++ +// Implementation of the Philox4x32-10 engine: a counter-based pseudorandom number generator (PRNG) +// that uses 4x32-bit keys and performs 10 rounds of mixing to produce high-quality randomness. +//-- + +#include "algorithms/engines/philox4x32x10/philox4x32x10.h" +#include "src/externals/service_dispatch.h" +#include "src/algorithms/engines/philox4x32x10/philox4x32x10_batch_impl.h" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace philox4x32x10 +{ +namespace interface1 +{ +using namespace daal::services; +using namespace philox4x32x10::internal; + +template +SharedPtr > Batch::create(size_t seed) +{ + SharedPtr > engPtr; +#define DAAL_CREATE_ENGINE_CPU(cpuId, ...) engPtr.reset(new BatchImpl(__VA_ARGS__)); + + DAAL_DISPATCH_FUNCTION_BY_CPU(DAAL_CREATE_ENGINE_CPU, seed); + +#undef DAAL_CREATE_ENGINE_CPU + return engPtr; +} + +template class Batch; +template class Batch; + +} // namespace interface1 +} // namespace philox4x32x10 +} // namespace engines +} // namespace algorithms +} // namespace daal diff --git a/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_batch_container.h b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_batch_container.h new file mode 100644 index 00000000000..9cb747e95a8 --- /dev/null +++ b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_batch_container.h @@ -0,0 +1,68 @@ +/* file: philox4x32x10_batch_container.h */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +/* +//++ +// Implementation of philox4x32x10 calculation algorithm container. +//-- +*/ + +#ifndef __PHILOX4X32X10_BATCH_CONTAINER_H__ +#define __PHILOX4X32X10_BATCH_CONTAINER_H__ + +#include "algorithms/engines/philox4x32x10/philox4x32x10.h" +#include "src/algorithms/engines/philox4x32x10/philox4x32x10_kernel.h" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace philox4x32x10 +{ +namespace interface1 +{ +template +BatchContainer::BatchContainer(daal::services::Environment::env * daalEnv) : AnalysisContainerIface(daalEnv) +{ + __DAAL_INITIALIZE_KERNELS(internal::philox4x32x10Kernel, algorithmFPType, method); +} + +template +BatchContainer::~BatchContainer() +{ + __DAAL_DEINITIALIZE_KERNELS(); +} + +template +services::Status BatchContainer::compute() +{ + daal::services::Environment::env & env = *_env; + engines::Result * result = static_cast(_res); + NumericTable * resultTable = result->get(engines::randomNumbers).get(); + + __DAAL_CALL_KERNEL(env, internal::philox4x32x10Kernel, __DAAL_KERNEL_ARGUMENTS(algorithmFPType, method), compute, resultTable); +} + +} // namespace interface1 +} // namespace philox4x32x10 +} // namespace engines +} // namespace algorithms +} // namespace daal + +#endif diff --git a/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_batch_impl.h b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_batch_impl.h new file mode 100644 index 00000000000..ebf6b7469ec --- /dev/null +++ b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_batch_impl.h @@ -0,0 +1,114 @@ +/* file: philox4x32x10_batch_impl.h */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +/* +//++ +// Implementation of the class defining the philox4x32x10 engine +//-- +*/ + +#include "algorithms/engines/philox4x32x10/philox4x32x10.h" +#include "src/algorithms/engines/engine_batch_impl.h" +#include "src/externals/service_rng.h" +#include "src/data_management/service_numeric_table.h" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace philox4x32x10 +{ +namespace internal +{ +template +class BatchImpl : public algorithms::engines::philox4x32x10::interface1::Batch, + public algorithms::engines::internal::BatchBaseImpl +{ +public: + typedef algorithms::engines::philox4x32x10::interface1::Batch super1; + typedef algorithms::engines::internal::BatchBaseImpl super2; + BatchImpl(size_t seed = 777) : baseRng(seed, __DAAL_BRNG_PHILOX4X32X10), super2(seed) {} + + void * getState() DAAL_C11_OVERRIDE { return baseRng.getState(); } + + int getStateSize() const DAAL_C11_OVERRIDE { return baseRng.getStateSize(); } + + services::Status saveStateImpl(byte * dest) const DAAL_C11_OVERRIDE + { + DAAL_CHECK(!baseRng.saveState((void *)dest), ErrorIncorrectErrorcodeFromGenerator); + return services::Status(); + } + + services::Status loadStateImpl(const byte * src) DAAL_C11_OVERRIDE + { + DAAL_CHECK(!baseRng.loadState((const void *)src), ErrorIncorrectErrorcodeFromGenerator); + return services::Status(); + } + + services::Status leapfrogImpl(size_t threadNum, size_t nThreads) DAAL_C11_OVERRIDE + { + int errcode = baseRng.leapfrog(threadNum, nThreads); + services::Status s; + if (errcode == __DAAL_LEAPFROG_METHOD_ERRCODE) + s.add(ErrorLeapfrogUnsupported); + else if (errcode) + s.add(ErrorIncorrectErrorcodeFromGenerator); + return s; + } + + services::Status skipAheadImpl(size_t nSkip) DAAL_C11_OVERRIDE + { + int errcode = baseRng.skipAhead(nSkip); + services::Status s; + if (errcode == __DAAL_SKIP_AHEAD_METHOD_ERRCODE) + s.add(ErrorSkipAheadUnsupported); + else if (errcode) + s.add(ErrorIncorrectErrorcodeFromGenerator); + return s; + } + + virtual BatchImpl * cloneImpl() const DAAL_C11_OVERRIDE + { + return new BatchImpl(*this); + } + + bool hasSupport(engines::internal::ParallelizationTechnique technique) const DAAL_C11_OVERRIDE + { + switch (technique) + { + case engines::internal::family: return false; + case engines::internal::skipahead: return true; + case engines::internal::leapfrog: return false; + } + return false; + } + + ~BatchImpl() {} + +protected: + BatchImpl(const BatchImpl & other) : super1(other), super2(other), baseRng(other.baseRng) {} + + daal::internal::BaseRNGsInst baseRng; +}; + +} // namespace internal +} // namespace philox4x32x10 +} // namespace engines +} // namespace algorithms +} // namespace daal diff --git a/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_dense_default_batch_fpt_cpu.cpp b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_dense_default_batch_fpt_cpu.cpp new file mode 100644 index 00000000000..946517c1d9c --- /dev/null +++ b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_dense_default_batch_fpt_cpu.cpp @@ -0,0 +1,47 @@ +/* file: philox4x32x10_dense_default_batch_fpt_cpu.cpp */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +//++ +// Implementation of philox4x32x10 calculation functions. +//-- + +#include "src/algorithms/engines/philox4x32x10/philox4x32x10_batch_container.h" +#include "src/algorithms/engines/philox4x32x10/philox4x32x10_kernel.h" +#include "src/algorithms/engines/philox4x32x10/philox4x32x10_impl.i" + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace philox4x32x10 +{ +namespace interface1 +{ +template class BatchContainer; +} // namespace interface1 + +namespace internal +{ +template class philox4x32x10Kernel; +} // namespace internal + +} // namespace philox4x32x10 +} // namespace engines +} // namespace algorithms +} // namespace daal diff --git a/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_dense_default_batch_fpt_dispatcher.cpp b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_dense_default_batch_fpt_dispatcher.cpp new file mode 100644 index 00000000000..1640fc4ec12 --- /dev/null +++ b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_dense_default_batch_fpt_dispatcher.cpp @@ -0,0 +1,30 @@ +/* file: philox4x32x10_dense_default_batch_fpt_dispatcher.cpp */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +//++ +// Implementation of philox4x32x10 calculation algorithm dispatcher. +//-- + +#include "src/algorithms/engines/philox4x32x10/philox4x32x10_batch_container.h" + +namespace daal +{ +namespace algorithms +{ +__DAAL_INSTANTIATE_DISPATCH_CONTAINER(engines::philox4x32x10::BatchContainer, batch, DAAL_FPTYPE, engines::philox4x32x10::defaultDense) +} // namespace algorithms +} // namespace daal diff --git a/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_impl.i b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_impl.i new file mode 100644 index 00000000000..5aa5addc22b --- /dev/null +++ b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_impl.i @@ -0,0 +1,49 @@ +/* file: philox4x32x10_impl.i */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +/* +//++ +// Implementation of philox4x32x10 algorithm. +//-- +*/ + +#ifndef __PHILOX4X32X10_IMPL_I__ +#define __PHILOX4X32X10_IMPL_I__ + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace philox4x32x10 +{ +namespace internal +{ +template +Status philox4x32x10Kernel::compute(NumericTable * resultTensor) +{ + return Status(); +} + +} // namespace internal +} // namespace philox4x32x10 +} // namespace engines +} // namespace algorithms +} // namespace daal + +#endif diff --git a/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_kernel.h b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_kernel.h new file mode 100644 index 00000000000..5870d781abd --- /dev/null +++ b/cpp/daal/src/algorithms/engines/philox4x32x10/philox4x32x10_kernel.h @@ -0,0 +1,58 @@ +/* file: philox4x32x10_kernel.h */ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +//++ +// Declaration of a template function for generating values using the Philox4x32-10 engine. +//-- + +#ifndef __PHILOX4X32X10_KERNEL_H__ +#define __PHILOX4X32X10_KERNEL_H__ + +#include "algorithms/engines/philox4x32x10/philox4x32x10.h" +#include "src/algorithms/kernel.h" +#include "data_management/data/numeric_table.h" + +using namespace daal::services; +using namespace daal::data_management; + +namespace daal +{ +namespace algorithms +{ +namespace engines +{ +namespace philox4x32x10 +{ +namespace internal +{ +/** + * \brief Kernel for philox4x32x10 calculation + */ +template +class philox4x32x10Kernel : public Kernel +{ +public: + Status compute(NumericTable * resultTable); +}; + +} // namespace internal +} // namespace philox4x32x10 +} // namespace engines +} // namespace algorithms +} // namespace daal + +#endif diff --git a/cpp/daal/src/externals/service_rng_mkl.h b/cpp/daal/src/externals/service_rng_mkl.h index b2dcd81b78b..14f6fcc8c06 100644 --- a/cpp/daal/src/externals/service_rng_mkl.h +++ b/cpp/daal/src/externals/service_rng_mkl.h @@ -32,6 +32,8 @@ #define __DAAL_BRNG_MT2203 VSL_BRNG_MT2203 #define __DAAL_BRNG_MT19937 VSL_BRNG_MT19937 #define __DAAL_BRNG_MCG59 VSL_BRNG_MCG59 +#define __DAAL_BRNG_MRG32K3A VSL_BRNG_MRG32K3A +#define __DAAL_BRNG_PHILOX4X32X10 VSL_BRNG_PHILOX4X32X10 #define __DAAL_RNG_METHOD_UNIFORM_STD VSL_RNG_METHOD_UNIFORM_STD #define __DAAL_RNG_METHOD_UNIFORMBITS32_STD 0 #define __DAAL_RNG_METHOD_BERNOULLI_ICDF VSL_RNG_METHOD_BERNOULLI_ICDF @@ -39,6 +41,10 @@ #define __DAAL_RNG_METHOD_GAUSSIAN_BOXMULLER2 VSL_RNG_METHOD_GAUSSIAN_BOXMULLER2 #define __DAAL_RNG_METHOD_GAUSSIAN_ICDF VSL_RNG_METHOD_GAUSSIAN_ICDF +// Errors +#define __DAAL_LEAPFROG_METHOD_ERRCODE VSL_RNG_ERROR_LEAPFROG_UNSUPPORTED +#define __DAAL_SKIP_AHEAD_METHOD_ERRCODE VSL_RNG_ERROR_SKIPAHEAD_UNSUPPORTED + namespace daal { namespace internal diff --git a/cpp/daal/src/externals/service_rng_openrng.h b/cpp/daal/src/externals/service_rng_openrng.h index dd70c644606..d8109ccc366 100644 --- a/cpp/daal/src/externals/service_rng_openrng.h +++ b/cpp/daal/src/externals/service_rng_openrng.h @@ -25,6 +25,8 @@ #define __DAAL_BRNG_MT2203 VSL_BRNG_MT2203 #define __DAAL_BRNG_MT19937 VSL_BRNG_MT19937 #define __DAAL_BRNG_MCG59 VSL_BRNG_MCG59 +#define __DAAL_BRNG_MRG32K3A VSL_BRNG_MRG32K3A +#define __DAAL_BRNG_PHILOX4X32X10 VSL_BRNG_PHILOX4X32X10 #define __DAAL_RNG_METHOD_UNIFORM_STD VSL_RNG_METHOD_UNIFORM_STD #define __DAAL_RNG_METHOD_UNIFORMBITS32_STD 0 #define __DAAL_RNG_METHOD_BERNOULLI_ICDF VSL_RNG_METHOD_BERNOULLI_ICDF @@ -32,6 +34,10 @@ #define __DAAL_RNG_METHOD_GAUSSIAN_BOXMULLER2 VSL_RNG_METHOD_GAUSSIAN_BOXMULLER2 #define __DAAL_RNG_METHOD_GAUSSIAN_ICDF VSL_RNG_METHOD_GAUSSIAN_ICDF +// Erros +#define __DAAL_LEAPFROG_METHOD_ERRCODE VSL_RNG_ERROR_LEAPFROG_UNSUPPORTED +#define __DAAL_SKIP_AHEAD_METHOD_ERRCODE VSL_RNG_ERROR_SKIPAHEAD_UNSUPPORTED + namespace daal { namespace internal diff --git a/cpp/daal/src/externals/service_rng_ref.h b/cpp/daal/src/externals/service_rng_ref.h index fc56fcf6205..3957a895226 100644 --- a/cpp/daal/src/externals/service_rng_ref.h +++ b/cpp/daal/src/externals/service_rng_ref.h @@ -36,16 +36,22 @@ #include // RNGs - #define __DAAL_BRNG_MT2203 (1 << 20) * 9 //VSL_BRNG_MT2203 - #define __DAAL_BRNG_MT19937 (1 << 20) * 8 //VSL_BRNG_MT19937 - #define __DAAL_BRNG_MCG59 (1 << 20) * 4 //VSL_BRNG_MCG59 + #define __DAAL_BRNG_MT2203 (1 << 20) * 9 //VSL_BRNG_MT2203 + #define __DAAL_BRNG_MT19937 (1 << 20) * 8 //VSL_BRNG_MT19937 + #define __DAAL_BRNG_MCG59 (1 << 20) * 4 //VSL_BRNG_MCG59 + #define __DAAL_BRNG_MRG32K3A (1 << 20) * 3 //VSL_BRNG_MRG32K3A + #define __DAAL_BRNG_PHILOX4X32X10 (1 << 20) * 16 //VSL_BRNG_PHILOX4X32X10 #define __DAAL_RNG_METHOD_UNIFORM_STD 0 //VSL_RNG_METHOD_UNIFORM_STD #define __DAAL_RNG_METHOD_UNIFORMBITS32_STD 4 - #define __DAAL_RNG_METHOD_BERNOULLI_ICDF 0 //VSL_RNG_METHOD_BERNOULLI_ICDF - #define __DAAL_RNG_METHOD_GAUSSIAN_BOXMULLER 0 //VSL_RNG_METHOD_GAUSSIAN_BOXMULLER - #define __DAAL_RNG_METHOD_GAUSSIAN_BOXMULLER2 1 //VSL_RNG_METHOD_GAUSSIAN_BOXMULLER2 - #define __DAAL_RNG_METHOD_GAUSSIAN_ICDF 2 //VSL_RNG_METHOD_GAUSSIAN_ICDF + #define __DAAL_RNG_METHOD_BERNOULLI_ICDF 0 //VSL_RNG_METHOD_BERNOULLI_ICDF + #define __DAAL_RNG_METHOD_GAUSSIAN_BOXMULLER 0 //VSL_RNG_METHOD_GAUSSIAN_BOXMULLER + #define __DAAL_RNG_METHOD_GAUSSIAN_BOXMULLER2 1 //VSL_RNG_METHOD_GAUSSIAN_BOXMULLER2 + #define __DAAL_RNG_METHOD_GAUSSIAN_ICDF 2 //VSL_RNG_METHOD_GAUSSIAN_ICDF + + // Errors + #define __DAAL_LEAPFROG_METHOD_ERRCODE -1002 // VSL_RNG_ERROR_LEAPFROG_UNSUPPORTED + #define __DAAL_SKIP_AHEAD_METHOD_ERRCODE -1003 // VSL_RNG_ERROR_SKIPAHEAD_UNSUPPORTED namespace daal { diff --git a/cpp/oneapi/dal/algo/connected_components/backend/cpu/vertex_partitioning_default_kernel.hpp b/cpp/oneapi/dal/algo/connected_components/backend/cpu/vertex_partitioning_default_kernel.hpp index 4da1866e277..497e98427e2 100644 --- a/cpp/oneapi/dal/algo/connected_components/backend/cpu/vertex_partitioning_default_kernel.hpp +++ b/cpp/oneapi/dal/algo/connected_components/backend/cpu/vertex_partitioning_default_kernel.hpp @@ -24,7 +24,7 @@ #include "oneapi/dal/backend/memory.hpp" #include "oneapi/dal/backend/interop/common.hpp" #include "oneapi/dal/table/homogen.hpp" -#include "oneapi/dal/backend/primitives/rng/rng_engine.hpp" +#include "oneapi/dal/backend/primitives/rng/host_engine.hpp" #include "oneapi/dal/detail/threading.hpp" namespace oneapi::dal::preview::connected_components::backend { @@ -90,9 +90,12 @@ std::int32_t most_frequent_element(const std::atomic *components, const std::int64_t &samples_count = 1024) { std::int32_t *rnd_vertex_ids = allocate(vertex_allocator, samples_count); - dal::backend::primitives::engine eng; - dal::backend::primitives::rng rn_gen; - rn_gen.uniform(samples_count, rnd_vertex_ids, eng.get_state(), 0, vertex_count); + dal::backend::primitives::host_engine eng; + dal::backend::primitives::uniform(samples_count, + rnd_vertex_ids, + eng, + 0, + vertex_count); std::int32_t *root_sample_counts = allocate(vertex_allocator, vertex_count); diff --git a/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_kernel_hist_impl.hpp b/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_kernel_hist_impl.hpp index 9dfe252e849..64dbae4c084 100644 --- a/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_kernel_hist_impl.hpp +++ b/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_kernel_hist_impl.hpp @@ -21,7 +21,7 @@ #include "oneapi/dal/backend/primitives/utils.hpp" #include "oneapi/dal/algo/decision_forest/train_types.hpp" -#include "oneapi/dal/backend/primitives/rng/rng_engine_collection.hpp" +#include "oneapi/dal/backend/primitives/rng/host_engine_collection.hpp" #include "oneapi/dal/algo/decision_forest/backend/gpu/train_misc_structs.hpp" #include "oneapi/dal/algo/decision_forest/backend/gpu/train_impurity_data.hpp" @@ -50,7 +50,7 @@ class train_kernel_hist_impl { using model_manager_t = train_model_manager; using train_context_t = train_context; using imp_data_t = impurity_data; - using rng_engine_t = pr::engine; + using rng_engine_t = pr::host_engine; using rng_engine_list_t = std::vector; using msg = dal::detail::error_messages; using comm_t = bk::communicator; diff --git a/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_kernel_hist_impl_dpc.cpp b/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_kernel_hist_impl_dpc.cpp index fc875683784..a396f11c048 100644 --- a/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_kernel_hist_impl_dpc.cpp +++ b/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_kernel_hist_impl_dpc.cpp @@ -396,14 +396,13 @@ sycl::event train_kernel_hist_impl::gen_initial_tree_or Index* const node_list_ptr = node_list_host.get_mutable_data(); for (Index node_idx = 0; node_idx < node_count; ++node_idx) { - pr::rng rn_gen; Index* gen_row_idx_global_ptr = selected_row_global_ptr + ctx.selected_row_total_count_ * node_idx; - rn_gen.uniform(ctx.selected_row_total_count_, - gen_row_idx_global_ptr, - rng_engine_list[engine_offset + node_idx].get_state(), - 0, - ctx.row_total_count_); + pr::uniform(ctx.selected_row_total_count_, + gen_row_idx_global_ptr, + rng_engine_list[engine_offset + node_idx], + 0, + ctx.row_total_count_); if (ctx.distr_mode_) { Index* node_ptr = node_list_ptr + node_idx * impl_const_t::node_prop_count_; @@ -483,15 +482,14 @@ train_kernel_hist_impl::gen_feature_list( auto node_vs_tree_map_list_host = node_vs_tree_map_list.to_host(queue_); - pr::rng rn_gen; auto tree_map_ptr = node_vs_tree_map_list_host.get_mutable_data(); if (ctx.selected_ftr_count_ != ctx.column_count_) { for (Index node = 0; node < node_count; ++node) { - rn_gen.uniform_without_replacement( + pr::uniform_without_replacement( ctx.selected_ftr_count_, selected_features_host_ptr + node * ctx.selected_ftr_count_, selected_features_host_ptr + (node + 1) * ctx.selected_ftr_count_, - rng_engine_list[tree_map_ptr[node]].get_state(), + rng_engine_list[tree_map_ptr[node]], 0, ctx.column_count_); } @@ -524,7 +522,6 @@ train_kernel_hist_impl::gen_random_thresholds( auto node_vs_tree_map_list_host = node_vs_tree_map.to_host(queue_); - pr::rng rn_gen; auto tree_map_ptr = node_vs_tree_map_list_host.get_mutable_data(); // Create arrays for random generated bins @@ -537,11 +534,11 @@ train_kernel_hist_impl::gen_random_thresholds( // Generate random bins for selected features for (Index node = 0; node < node_count; ++node) { - rn_gen.uniform(ctx.selected_ftr_count_, - random_bins_host_ptr + node * ctx.selected_ftr_count_, - rng_engine_list[tree_map_ptr[node]].get_state(), - 0.0f, - 1.0f); + pr::uniform(ctx.selected_ftr_count_, + random_bins_host_ptr + node * ctx.selected_ftr_count_, + rng_engine_list[tree_map_ptr[node]], + 0.0f, + 1.0f); } auto event_rnd_generate = random_bins_com.assign_from_host(queue_, random_bins_host_ptr, random_bins_com.get_count()); @@ -1660,12 +1657,10 @@ sycl::event train_kernel_hist_impl::compute_results( const Float div1 = Float(1) / Float(built_tree_count + tree_idx_in_block + 1); - pr::rng rn_gen; - for (Index column_idx = 0; column_idx < ctx.column_count_; ++column_idx) { - rn_gen.shuffle(oob_row_count, - permutation_ptr, - engine_arr[built_tree_count + tree_idx_in_block].get_state()); + pr::shuffle(oob_row_count, + permutation_ptr, + engine_arr[built_tree_count + tree_idx_in_block]); const Float oob_err_perm = compute_oob_error_perm(ctx, model_manager, data_host, @@ -1858,7 +1853,7 @@ train_result train_kernel_hist_impl::operator()( de::check_mul_overflow((ctx.tree_count_ - 1), skip_num); - pr::engine_collection collection(ctx.tree_count_, desc.get_seed()); + pr::host_engine_collection collection(ctx.tree_count_, desc.get_seed()); rng_engine_list_t engine_arr = collection([&](std::size_t i, std::size_t& skip) { skip = i * skip_num; }); diff --git a/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_splitter_impl.hpp b/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_splitter_impl.hpp index 2cadc5f72ae..ed26b7cd491 100644 --- a/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_splitter_impl.hpp +++ b/cpp/oneapi/dal/algo/decision_forest/backend/gpu/train_splitter_impl.hpp @@ -19,7 +19,7 @@ #include "oneapi/dal/table/common.hpp" #include "oneapi/dal/backend/primitives/ndarray.hpp" #include "oneapi/dal/backend/primitives/utils.hpp" -#include "oneapi/dal/backend/primitives/rng/rng_engine_collection.hpp" + #include "oneapi/dal/algo/decision_forest/train_types.hpp" #include "oneapi/dal/algo/decision_forest/backend/gpu/train_misc_structs.hpp" diff --git a/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernel_distr.hpp b/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernel_distr.hpp index be51ce18880..bceb0e3f4d9 100644 --- a/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernel_distr.hpp +++ b/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernel_distr.hpp @@ -18,6 +18,7 @@ #include "oneapi/dal/algo/kmeans_init/compute_types.hpp" #include "oneapi/dal/backend/dispatcher.hpp" +#include "oneapi/dal/backend/primitives/rng/host_engine.hpp" namespace oneapi::dal::kmeans_init::backend { diff --git a/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernel_distr_random_dense_dpc.cpp b/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernel_distr_random_dense_dpc.cpp index 642e937b562..b6a5c1395d6 100644 --- a/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernel_distr_random_dense_dpc.cpp +++ b/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernel_distr_random_dense_dpc.cpp @@ -75,7 +75,7 @@ ids_arr_t generate_random_indices(std::int64_t count, std::int64_t scount, std:: ids_arr_t result = ids_arr_t::empty(count); auto ndres = pr::ndview::wrap(result.get_mutable_data(), { count }); ONEDAL_ASSERT(count < scount); - partial_fisher_yates_shuffle(ndres, scount, seed); + pr::partial_fisher_yates_shuffle(ndres, scount, seed); return result; } diff --git a/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernels_impl.hpp b/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernels_impl.hpp index 1a37a0a963e..70ec2dc2e6d 100644 --- a/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernels_impl.hpp +++ b/cpp/oneapi/dal/algo/kmeans_init/backend/gpu/compute_kernels_impl.hpp @@ -18,7 +18,6 @@ #include "oneapi/dal/backend/common.hpp" #include "oneapi/dal/backend/primitives/ndarray.hpp" -#include "oneapi/dal/backend/primitives/rng/partial_shuffle.hpp" #include "oneapi/dal/backend/primitives/selection/select_indexed_rows.hpp" #include "oneapi/dal/algo/kmeans_init/common.hpp" diff --git a/cpp/oneapi/dal/algo/louvain/backend/cpu/louvain_data.hpp b/cpp/oneapi/dal/algo/louvain/backend/cpu/louvain_data.hpp index d21de8c9627..b41ef0e23a6 100644 --- a/cpp/oneapi/dal/algo/louvain/backend/cpu/louvain_data.hpp +++ b/cpp/oneapi/dal/algo/louvain/backend/cpu/louvain_data.hpp @@ -17,7 +17,7 @@ #pragma once #include "oneapi/dal/backend/memory.hpp" -#include "oneapi/dal/backend/primitives/rng/rng_engine.hpp" +#include "oneapi/dal/backend/primitives/rng/host_engine.hpp" namespace oneapi::dal::preview::louvain::backend { using namespace oneapi::dal::preview::detail; @@ -123,8 +123,7 @@ struct louvain_data { // Total link weight in the network value_type m; - engine eng; - rng rn_gen; + host_engine eng; const std::int64_t vertex_count; const std::int64_t edge_count; diff --git a/cpp/oneapi/dal/algo/louvain/backend/cpu/vertex_partitioning_default_kernel.hpp b/cpp/oneapi/dal/algo/louvain/backend/cpu/vertex_partitioning_default_kernel.hpp index 79e294e9f47..e287c3f2f66 100644 --- a/cpp/oneapi/dal/algo/louvain/backend/cpu/vertex_partitioning_default_kernel.hpp +++ b/cpp/oneapi/dal/algo/louvain/backend/cpu/vertex_partitioning_default_kernel.hpp @@ -206,7 +206,7 @@ inline Float move_nodes(const dal::preview::detail::topology& t, ld.random_order[index] = index; } // random shuffle - ld.rn_gen.uniform(t._vertex_count, ld.index, ld.eng.get_state(), 0, t._vertex_count); + uniform(t._vertex_count, ld.index, ld.eng, 0, t._vertex_count); for (std::int64_t index = 0; index < t._vertex_count; ++index) { std::swap(ld.random_order[index], ld.random_order[ld.index[index]]); } diff --git a/cpp/oneapi/dal/backend/primitives/objective_function/test/fixture.hpp b/cpp/oneapi/dal/backend/primitives/objective_function/test/fixture.hpp index d22a2dde0a1..8eaee870d3a 100644 --- a/cpp/oneapi/dal/backend/primitives/objective_function/test/fixture.hpp +++ b/cpp/oneapi/dal/backend/primitives/objective_function/test/fixture.hpp @@ -25,7 +25,7 @@ #include "oneapi/dal/table/csr_accessor.hpp" #include "oneapi/dal/detail/debug.hpp" -#include "oneapi/dal/backend/primitives/rng/rng_engine.hpp" +#include "oneapi/dal/backend/primitives/rng/host_engine.hpp" namespace oneapi::dal::backend::primitives::test { @@ -572,13 +572,12 @@ class logloss_test : public te::float_algo_fixture rn_gen; auto vec_host = ndarray::empty(this->get_queue(), { dim }, sycl::usm::alloc::host); for (std::int32_t ij = 0; ij < num_checks; ++ij) { - primitives::engine eng(2007 + dim * num_checks + ij); - rn_gen.uniform(dim, vec_host.get_mutable_data(), eng.get_state(), -1.0, 1.0); + primitives::host_engine eng(2007 + dim * num_checks + ij); + primitives::uniform(dim, vec_host.get_mutable_data(), eng, -1.0, 1.0); auto vec_gpu = vec_host.to_device(this->get_queue()); auto out_vector = ndarray::empty(this->get_queue(), { dim }, sycl::usm::alloc::device); diff --git a/cpp/oneapi/dal/backend/primitives/objective_function/test/spmd_fixture.hpp b/cpp/oneapi/dal/backend/primitives/objective_function/test/spmd_fixture.hpp index e902dd452e1..63ab0a07c13 100644 --- a/cpp/oneapi/dal/backend/primitives/objective_function/test/spmd_fixture.hpp +++ b/cpp/oneapi/dal/backend/primitives/objective_function/test/spmd_fixture.hpp @@ -100,12 +100,12 @@ class logloss_spmd_test : public logloss_test { std::int64_t num_checks = 5; std::vector> vecs_host(num_checks), vecs_gpu(num_checks); - rng rn_gen; + for (std::int64_t ij = 0; ij < num_checks; ++ij) { - engine eng(2007 + dim * num_checks + ij); + host_engine eng(2007 + dim * num_checks + ij); vecs_host[ij] = (ndarray::empty(this->get_queue(), { dim }, sycl::usm::alloc::host)); - rn_gen.uniform(dim, vecs_host[ij].get_mutable_data(), eng.get_state(), -1.0, 1.0); + uniform(dim, vecs_host[ij].get_mutable_data(), eng, -1.0, 1.0); vecs_gpu[ij] = vecs_host[ij].to_device(this->get_queue()); } diff --git a/cpp/oneapi/dal/backend/primitives/optimizers/test/cg_solver_dpc.cpp b/cpp/oneapi/dal/backend/primitives/optimizers/test/cg_solver_dpc.cpp index ea320f690a2..6ac19ef1cbe 100644 --- a/cpp/oneapi/dal/backend/primitives/optimizers/test/cg_solver_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/optimizers/test/cg_solver_dpc.cpp @@ -20,7 +20,7 @@ #include "oneapi/dal/test/engine/common.hpp" #include "oneapi/dal/test/engine/fixtures.hpp" #include "oneapi/dal/table/row_accessor.hpp" -#include "oneapi/dal/backend/primitives/rng/rng_engine.hpp" +#include "oneapi/dal/backend/primitives/rng/host_engine.hpp" #include namespace oneapi::dal::backend::primitives::test { @@ -43,9 +43,8 @@ class cg_solver_test : public te::float_algo_fixture { x_host_ = ndarray::empty(this->get_queue(), { n_ }, sycl::usm::alloc::host); b_host_ = ndarray::empty(this->get_queue(), { n_ }, sycl::usm::alloc::host); - primitives::rng rn_gen; - primitives::engine eng(4014 + n_); - rn_gen.uniform(n_, x_host_.get_mutable_data(), eng.get_state(), -1.0, 1.0); + primitives::host_engine eng(4014 + n_); + primitives::uniform(n_, x_host_.get_mutable_data(), eng, -1.0, 1.0); create_stable_matrix(this->get_queue(), A_host_); diff --git a/cpp/oneapi/dal/backend/primitives/optimizers/test/fixture.hpp b/cpp/oneapi/dal/backend/primitives/optimizers/test/fixture.hpp index a6b87b2dcc1..a397246a9a7 100644 --- a/cpp/oneapi/dal/backend/primitives/optimizers/test/fixture.hpp +++ b/cpp/oneapi/dal/backend/primitives/optimizers/test/fixture.hpp @@ -21,7 +21,7 @@ #include "oneapi/dal/backend/primitives/ndarray.hpp" #include "oneapi/dal/test/engine/common.hpp" #include "oneapi/dal/test/engine/fixtures.hpp" -#include "oneapi/dal/backend/primitives/rng/rng_engine.hpp" +#include "oneapi/dal/backend/primitives/rng/host_engine.hpp" #include "oneapi/dal/backend/primitives/blas/gemv.hpp" #include "oneapi/dal/backend/primitives/element_wise.hpp" @@ -133,11 +133,10 @@ void create_stable_matrix(sycl::queue& queue, ONEDAL_ASSERT(A.get_dimension(1) == n); auto J = ndarray::empty(queue, { n, n }, sycl::usm::alloc::host); auto eigen_values = ndarray::empty(queue, { n }, sycl::usm::alloc::host); - primitives::rng rn_gen; - primitives::engine eng(2007 + n); + primitives::host_engine eng(2007 + n); - rn_gen.uniform(n * n, J.get_mutable_data(), eng.get_state(), -1.0, 1.0); - rn_gen.uniform(n, eigen_values.get_mutable_data(), eng.get_state(), bottom_eig, top_eig); + primitives::uniform(n * n, J.get_mutable_data(), eng, -1.0, 1.0); + primitives::uniform(n, eigen_values.get_mutable_data(), eng, bottom_eig, top_eig); // orthogonalize matrix J gram_schmidt(J); diff --git a/cpp/oneapi/dal/backend/primitives/optimizers/test/newton_cg_dpc.cpp b/cpp/oneapi/dal/backend/primitives/optimizers/test/newton_cg_dpc.cpp index f473dddf1f7..ad2faabeff8 100644 --- a/cpp/oneapi/dal/backend/primitives/optimizers/test/newton_cg_dpc.cpp +++ b/cpp/oneapi/dal/backend/primitives/optimizers/test/newton_cg_dpc.cpp @@ -22,7 +22,7 @@ #include "oneapi/dal/test/engine/common.hpp" #include "oneapi/dal/test/engine/fixtures.hpp" #include "oneapi/dal/table/row_accessor.hpp" -#include "oneapi/dal/backend/primitives/rng/rng_engine.hpp" +#include "oneapi/dal/backend/primitives/rng/host_engine.hpp" #include #include "oneapi/dal/backend/primitives/objective_function.hpp" @@ -56,10 +56,10 @@ class newton_cg_test : public te::float_algo_fixture { ndarray::empty(this->get_queue(), { n_ + 1 }, sycl::usm::alloc::host); auto params_host = ndarray::empty(this->get_queue(), { p_ + 1 }, sycl::usm::alloc::host); - primitives::rng rn_gen; - primitives::engine eng(2007 + n); - rn_gen.uniform(n_ * p_, X_host.get_mutable_data(), eng.get_state(), -10.0, 10.0); - rn_gen.uniform(p_ + 1, params_host.get_mutable_data(), eng.get_state(), -5.0, 5.0); + + primitives::host_engine eng(2007 + n); + primitives::uniform(n_ * p_, X_host.get_mutable_data(), eng, -10.0, 10.0); + primitives::uniform(p_ + 1, params_host.get_mutable_data(), eng, -5.0, 5.0); for (std::int64_t i = 0; i < n_; ++i) { float_t val = 0; for (std::int64_t j = 0; j < p_; ++j) { @@ -142,9 +142,8 @@ class newton_cg_test : public te::float_algo_fixture { ndarray::empty(this->get_queue(), { n_, n_ }, sycl::usm::alloc::host); solution_ = ndarray::empty(this->get_queue(), { n_ }, sycl::usm::alloc::host); auto b_host = ndarray::empty(this->get_queue(), { n_ }, sycl::usm::alloc::host); - primitives::rng rn_gen; - primitives::engine eng(4014 + n_); - rn_gen.uniform(n_, solution_.get_mutable_data(), eng.get_state(), -1.0, 1.0); + primitives::host_engine eng(4014 + n_); + uniform(n_, solution_.get_mutable_data(), eng, -1.0, 1.0); create_stable_matrix(this->get_queue(), A_host, float_t(0.1), float_t(5.0)); @@ -164,7 +163,7 @@ class newton_cg_test : public te::float_algo_fixture { auto buffer = ndarray::empty(this->get_queue(), { n_ }, sycl::usm::alloc::host); for (std::int32_t test_num = 0; test_num < 5; ++test_num) { - rn_gen.uniform(n_, x_host.get_mutable_data(), eng.get_state(), -1.0, 1.0); + uniform(n_, x_host.get_mutable_data(), eng, -1.0, 1.0); auto x_gpu = x_host.to_device(this->get_queue()); auto compute_event_vec = func_->update_x(x_gpu, true, {}); wait_or_pass(compute_event_vec).wait_and_throw(); diff --git a/cpp/oneapi/dal/backend/primitives/rng/device_engine.hpp b/cpp/oneapi/dal/backend/primitives/rng/device_engine.hpp new file mode 100644 index 00000000000..7acb09812e5 --- /dev/null +++ b/cpp/oneapi/dal/backend/primitives/rng/device_engine.hpp @@ -0,0 +1,465 @@ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +#pragma once + +#include "oneapi/dal/backend/primitives/ndarray.hpp" + +#include "oneapi/dal/backend/primitives/rng/utils.hpp" +#include "oneapi/dal/backend/primitives/rng/rng_types.hpp" + +#include +#include +#include +#include +#include + +#include + +namespace mkl = oneapi::mkl; +namespace oneapi::dal::backend::primitives { + +#ifdef ONEDAL_DATA_PARALLEL + +/// Abstract base class for all random number generators (RNGs). +/// It defines a common interface for working with different types of RNGs, including methods +/// for retrieving the engine method and skipping ahead in the random number sequence. +class gen_base { +public: + virtual ~gen_base() = default; + + /// Method to retrieve the engine method. + /// @return The engine method as an enum value of `engine_type`. + virtual engine_type get_engine_type() const = 0; + + /// Method to skip ahead in the random number sequence. + /// @param[in] nSkip The number of steps to skip in the generator sequence. + virtual void skip_ahead_gpu(std::int64_t nSkip) = 0; +}; + +/// Implementation of the mt2203 random number generator for GPU. +/// This class encapsulates the mt2203 engine and provides an interface to use it +/// on devices supporting SYCL. +class gen_mt2203 : public gen_base { +public: + explicit gen_mt2203() = delete; + + /// Constructor that initializes the mt2203 generator for use on the GPU. + /// @param[in] queue The SYCL queue to manage device operations. + /// @param[in] seed The initial seed for the generator. + gen_mt2203(sycl::queue queue, std::int64_t seed, std::int64_t engine_idx = 0) + : _gen(queue, seed, engine_idx) {} + + /// Returns the engine method for mt2203. + /// @return The `mt2203` engine method as an enum value of `engine_type`. + engine_type get_engine_type() const override { + return engine_type::mt2203; + } + + /// Skips ahead in the random number sequence for mt2203 on the GPU. + /// Currently, the skip functionality is not implemented. + /// @param[in] nSkip The number of steps to skip in the sequence. + void skip_ahead_gpu(std::int64_t nSkip) override { + //skip; + } + + /// Retrieves a pointer to the underlying mt2203 generator. + /// @return A pointer to the `mt2203` RNG. + oneapi::mkl::rng::mt2203* get() { + return &_gen; + } + +protected: + oneapi::mkl::rng::mt2203 _gen; +}; + +/// Implementation of the philox4x32x10 random number generator for GPU. +/// This class encapsulates the philox4x32x10 engine and provides an interface to use it +/// on devices supporting SYCL. +class gen_philox : public gen_base { +public: + explicit gen_philox() = delete; + + /// Constructor that initializes the philox4x32x10 generator for use on the GPU. + /// @param[in] queue The SYCL queue to manage device operations. + /// @param[in] seed The initial seed for the generator. + gen_philox(sycl::queue queue, std::int64_t seed) : _gen(queue, seed) {} + + /// Returns the engine method for philox4x32x10. + /// @return The `philox4x32x10` engine method as an enum value of `engine_type`. + engine_type get_engine_type() const override { + return engine_type::philox4x32x10; + } + + /// Skips ahead in the random number sequence for philox4x32x10 on the GPU. + /// @param[in] nSkip The number of steps to skip in the sequence. + void skip_ahead_gpu(std::int64_t nSkip) override { + skip_ahead(_gen, nSkip); + } + + /// Retrieves a pointer to the underlying philox4x32x10 generator. + /// @return A pointer to the `philox4x32x10` RNG. + oneapi::mkl::rng::philox4x32x10* get() { + return &_gen; + } + +protected: + oneapi::mkl::rng::philox4x32x10 _gen; +}; + +/// Implementation of the mrg32k3a random number generator for GPU. +/// This class encapsulates the mrg32k3a engine and provides an interface to use it +/// on devices supporting SYCL. +class gen_mrg32k : public gen_base { +public: + explicit gen_mrg32k() = delete; + + /// Constructor that initializes the mrg32k3a generator for use on the GPU. + /// @param[in] queue The SYCL queue to manage device operations. + /// @param[in] seed The initial seed for the generator. + gen_mrg32k(sycl::queue queue, std::int64_t seed) : _gen(queue, seed) {} + + /// Returns the engine method for mrg32k3a. + /// @return The `mrg32k3a` engine method as an enum value of `engine_type`. + engine_type get_engine_type() const override { + return engine_type::mrg32k3a; + } + + /// Skips ahead in the random number sequence for mrg32k3a on the GPU. + /// @param[in] nSkip The number of steps to skip in the sequence. + void skip_ahead_gpu(std::int64_t nSkip) override { + skip_ahead(_gen, nSkip); + } + + /// Retrieves a pointer to the underlying mrg32k3a generator. + /// @return A pointer to the `mrg32k3a` RNG. + oneapi::mkl::rng::mrg32k3a* get() { + return &_gen; + } + +protected: + oneapi::mkl::rng::mrg32k3a _gen; +}; + +/// Implementation of the mt19937 random number generator for GPU. +/// This class encapsulates the mt19937 engine and provides an interface to use it +/// on devices supporting SYCL. +class gen_mt19937 : public gen_base { +public: + explicit gen_mt19937() = delete; + + /// Constructor that initializes the mt19937 generator for use on the GPU. + /// @param[in] queue The SYCL queue to manage device operations. + /// @param[in] seed The initial seed for the generator. + gen_mt19937(sycl::queue queue, std::int64_t seed) : _gen(queue, seed) {} + + /// Returns the engine method for mt19937. + /// @return The `mt19937` engine method as an enum value of `engine_type`. + engine_type get_engine_type() const override { + return engine_type::mt19937; + } + + /// Skips ahead in the random number sequence for mt19937 on the GPU. + /// @param[in] nSkip The number of steps to skip in the sequence. + void skip_ahead_gpu(std::int64_t nSkip) override { + skip_ahead(_gen, nSkip); + } + + /// Retrieves a pointer to the underlying mt19937 generator. + /// @return A pointer to the `mt19937` RNG. + oneapi::mkl::rng::mt19937* get() { + return &_gen; + } + +protected: + oneapi::mkl::rng::mt19937 _gen; +}; + +/// Implementation of the mcg59 random number generator for GPU. +/// This class encapsulates the mcg59 engine and provides an interface to use it +/// on devices supporting SYCL. +class gen_mcg59 : public gen_base { +public: + explicit gen_mcg59() = delete; + + /// Constructor that initializes the mcg59 generator for use on the GPU. + /// @param[in] queue The SYCL queue to manage device operations. + /// @param[in] seed The initial seed for the generator. + gen_mcg59(sycl::queue queue, std::int64_t seed) : _gen(queue, seed) {} + + /// Returns the engine method for mcg59. + /// @return The `mcg59` engine method as an enum value of `engine_type`. + engine_type get_engine_type() const override { + return engine_type::mcg59; + } + + /// Skips ahead in the random number sequence for mcg59 on the GPU. + /// @param[in] nSkip The number of steps to skip in the sequence. + void skip_ahead_gpu(std::int64_t nSkip) override { + skip_ahead(_gen, nSkip); + } + + /// Retrieves a pointer to the underlying mcg59 generator. + /// @return A pointer to the `mcg59` RNG. + oneapi::mkl::rng::mcg59* get() { + return &_gen; + } + +protected: + oneapi::mkl::rng::mcg59 _gen; +}; + +/// A class that provides a unified interface for random number generation on both CPU and GPU devices. +/// +/// This class serves as a wrapper for random number generators (RNGs) that supports different engine types, +/// enabling efficient random number generation on heterogeneous platforms using SYCL. It integrates a host +/// (CPU) engine and a device (GPU) engine, allowing operations to be executed seamlessly on the appropriate +/// device. +/// +/// The class provides functionality to skip ahead in the RNG sequence, retrieve engine states, and +/// manage host and device engines independently. Support for `skip_ahead` on GPU is currently limited for +/// some engine types. +class device_engine { +public: + /// @param[in] queue The SYCL queue used to manage device operations. + /// @param[in] seed The initial seed for the random number generator. Defaults to `777`. + /// @param[in] method The engine method. Defaults to `engine_type::mt2203`. + device_engine(sycl::queue& queue, + std::int64_t seed = 777, + engine_type method = engine_type::mt2203, + std::int64_t idx = 0) + : q(queue) { + switch (method) { + case engine_type::mt2203: + host_engine_ = daal::algorithms::engines::mt2203::Batch<>::create(seed); + dpc_engine_ = std::make_shared(queue, seed, idx); + break; + case engine_type::mcg59: + host_engine_ = daal::algorithms::engines::mcg59::Batch<>::create(seed); + dpc_engine_ = std::make_shared(queue, seed); + break; + case engine_type::mrg32k3a: + host_engine_ = daal::algorithms::engines::mrg32k3a::Batch<>::create(seed); + dpc_engine_ = std::make_shared(queue, seed); + break; + case engine_type::philox4x32x10: + host_engine_ = daal::algorithms::engines::philox4x32x10::Batch<>::create(seed); + dpc_engine_ = std::make_shared(queue, seed); + break; + case engine_type::mt19937: + host_engine_ = daal::algorithms::engines::mt19937::Batch<>::create(seed); + dpc_engine_ = std::make_shared(queue, seed); + break; + default: throw std::invalid_argument("Unsupported engine type 1"); + } + impl_ = + dynamic_cast(host_engine_.get()); + if (!impl_) { + throw std::domain_error("RNG engine is not supported"); + } + } + + // Copy constructor + device_engine(const device_engine&) = default; + + // Copy assignment operator + device_engine& operator=(const device_engine&) = default; + + // Move constructor + device_engine(device_engine&& other) noexcept = default; + + // Move assignment operator + device_engine& operator=(device_engine&& other) noexcept = default; + + /// Destructor. + ~device_engine() = default; + + /// Retrieves the state of the host rng engine(DAAL). + /// @return Pointer to the host engine state. + void* get_host_engine_state() const { + return impl_->getState(); + } + + /// Retrieves the base pointer of the device rng engine. + /// @return Shared pointer to the device rng engine base. + auto get_device_engine_base_ptr() { + return dpc_engine_; + } + + /// Advances the rng sequence on the CPU by a specified number of steps. + /// @param[in] nSkip The number of steps to skip. + void skip_ahead_cpu(size_t nSkip) { + host_engine_->skipAhead(nSkip); + } + + /// Advances the rng sequence on the GPU by a specified number of steps. + /// @param[in] nSkip The number of steps to skip. + void skip_ahead_gpu(size_t nSkip) { + dpc_engine_->skip_ahead_gpu(nSkip); + } + + /// Advances the rng sequence on both CPU and GPU by a specified number of steps. + /// @param[in] nSkip The number of steps to skip. + void skip_ahead(size_t nSkip) { + skip_ahead_cpu(nSkip); + skip_ahead_gpu(nSkip); + } + + /// Retrieves the SYCL queue associated with this rng engine. + /// @return Reference to the SYCL queue. + sycl::queue& get_queue() { + return q; + } + +private: + sycl::queue q; + daal::algorithms::engines::EnginePtr host_engine_; + std::shared_ptr dpc_engine_; + daal::algorithms::engines::internal::BatchBaseImpl* impl_; +}; + +/// Generates uniformly distributed random numbers on the CPU. +/// @tparam Type The data type of the generated numbers. +/// @param[in] count The number of random numbers to generate. +/// @param[out] dst Pointer to the output buffer. +/// @param[in] engine_ Reference to the device engine. +/// @param[in] a The lower bound of the uniform distribution. +/// @param[in] b The upper bound of the uniform distribution. +template +void uniform(std::int64_t count, Type* dst, device_engine& engine_, Type a, Type b) { + if (sycl::get_pointer_type(dst, engine_.get_queue().get_context()) == + sycl::usm::alloc::device) { + throw domain_error(dal::detail::error_messages::unsupported_data_type()); + } + auto state = engine_.get_host_engine_state(); + uniform_dispatcher::uniform_by_cpu(count, dst, state, a, b); + engine_.skip_ahead_gpu(count); +} + +/// Generates a random permutation of elements without replacement. +/// @tparam Type The data type of the elements. +/// @param[in] count The number of elements to generate. +/// @param[out] dst Pointer to the output buffer. +/// @param[out] buffer Temporary buffer used for computations. +/// @param[in] engine_ Reference to the device engine. +/// @param[in] a The lower bound of the range. +/// @param[in] b The upper bound of the range. +template +void uniform_without_replacement(std::int64_t count, + Type* dst, + Type* buffer, + device_engine& engine_, + Type a, + Type b) { + if (sycl::get_pointer_type(dst, engine_.get_queue().get_context()) == + sycl::usm::alloc::device) { + throw domain_error(dal::detail::error_messages::unsupported_data_type()); + } + void* state = engine_.get_host_engine_state(); + uniform_dispatcher::uniform_without_replacement_by_cpu(count, dst, buffer, state, a, b); + engine_.skip_ahead_gpu(count); +} + +/// Shuffles an array using random swaps. +/// @tparam Type The data type of the array elements. +/// @param[in] count The number of elements to shuffle. +/// @param[in, out] dst Pointer to the array to be shuffled. +/// @param[in] engine_ Reference to the device engine. +template >> +void shuffle(std::int64_t count, Type* dst, device_engine& engine_) { + if (sycl::get_pointer_type(dst, engine_.get_queue().get_context()) == + sycl::usm::alloc::device) { + throw domain_error(dal::detail::error_messages::unsupported_data_type()); + } + Type idx[2]; + void* state = engine_.get_host_engine_state(); + for (std::int64_t i = 0; i < count; ++i) { + uniform_dispatcher::uniform_by_cpu(2, idx, state, 0, count); + std::swap(dst[idx[0]], dst[idx[1]]); + } + engine_.skip_ahead_gpu(count); +} + +/// Generates uniformly distributed random numbers on the GPU. +/// @tparam Type The data type of the generated numbers. +/// @param[in] queue The SYCL queue for device execution. +/// @param[in] count The number of random numbers to generate. +/// @param[out] dst Pointer to the output buffer. +/// @param[in] engine_ Reference to the device engine. +/// @param[in] a The lower bound of the uniform distribution. +/// @param[in] b The upper bound of the uniform distribution. +/// @param[in] deps Dependencies for the SYCL event. +template +sycl::event uniform(sycl::queue& queue, + std::int64_t count, + Type* dst, + device_engine& engine_, + Type a, + Type b, + const event_vector& deps = {}); + +/// Generates a random permutation of elements without replacement on the GPU. +/// @tparam Type The data type of the elements. +/// @param[in] queue The SYCL queue for device execution. +/// @param[in] count The number of elements to generate. +/// @param[out] dst Pointer to the output buffer. +/// @param[out] buffer Temporary buffer used for computations. +/// @param[in] engine_ Reference to the device engine. +/// @param[in] a The lower bound of the range. +/// @param[in] b The upper bound of the range. +/// @param[in] deps Dependencies for the SYCL event. +template +sycl::event uniform_without_replacement(sycl::queue& queue, + std::int64_t count, + Type* dst, + Type* buffer, + device_engine& engine_, + Type a, + Type b, + const event_vector& deps = {}); + +/// Shuffles an array using random swaps on the GPU. +/// @tparam Type The data type of the array elements. +/// @param[in] queue The SYCL queue for device execution. +/// @param[in] count The number of elements to shuffle. +/// @param[in, out] dst Pointer to the array to be shuffled. +/// @param[in] engine_ Reference to the device engine. +/// @param[in] deps Dependencies for the SYCL event. +template +sycl::event shuffle(sycl::queue& queue, + std::int64_t count, + Type* dst, + device_engine& engine_, + const event_vector& deps = {}); + +/// Partially shuffles the first `top` elements of an array using the Fisher-Yates algorithm. +/// @tparam Type The data type of the array elements. +/// @param[in] queue_ The SYCL queue for device execution. +/// @param[in, out] result_array The array to be partially shuffled. +/// @param[in] top The number of elements to shuffle. +/// @param[in] seed The seed for the engine. +/// @param[in] method The rng engine type. Defaults to `mt19937`. +/// @param[in] deps Dependencies for the SYCL event. +template +sycl::event partial_fisher_yates_shuffle(sycl::queue& queue_, + ndview& result_array, + std::int64_t top, + std::int64_t seed, + engine_type method = engine_type::mt19937, + const event_vector& deps = {}); +#endif + +} // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/rng/device_engine_collection.hpp b/cpp/oneapi/dal/backend/primitives/rng/device_engine_collection.hpp new file mode 100644 index 00000000000..d76d0c656e1 --- /dev/null +++ b/cpp/oneapi/dal/backend/primitives/rng/device_engine_collection.hpp @@ -0,0 +1,63 @@ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +#pragma once + +#include "oneapi/dal/backend/primitives/rng/rng_dpc.hpp" +#include "oneapi/dal/backend/primitives/ndarray.hpp" +#include + +#include "oneapi/dal/backend/primitives/rng/utils.hpp" +#include "oneapi/dal/backend/primitives/rng/rng_types.hpp" +#include "oneapi/dal/table/common.hpp" + +namespace oneapi::dal::backend::primitives { + +#ifdef ONEDAL_DATA_PARALLEL + +class device_engine_collection { +public: + device_engine_collection(sycl::queue& queue, + std::int64_t count, + std::int64_t seed = 777, + engine_type method = engine_type::mt2203) + : count_(count), + base_seed_(seed) { + engines_.reserve(count_); + if (method == engine_type::mt2203) { + for (std::int64_t i = 0; i < count_; ++i) { + engines_.push_back(device_engine(queue, base_seed_, i, method)); + } + } + else { + for (std::int64_t i = 0; i < count_; ++i) { + engines_.push_back(device_engine(queue, base_seed_ + i, method)); + } + } + } + + std::vector> get_engines() const { + return engines_; + } + +private: + std::int64_t count_; + std::int64_t base_seed_; + std::vector engines_; +}; + +#endif +} // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/rng/device_engine_dpc.cpp b/cpp/oneapi/dal/backend/primitives/rng/device_engine_dpc.cpp new file mode 100644 index 00000000000..845d058ca30 --- /dev/null +++ b/cpp/oneapi/dal/backend/primitives/rng/device_engine_dpc.cpp @@ -0,0 +1,242 @@ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +#include "oneapi/dal/backend/primitives/rng/device_engine.hpp" +#include "oneapi/dal/backend/primitives/ndarray.hpp" + +#include + +namespace oneapi::dal::backend::primitives { + +namespace bk = oneapi::dal::backend; + +template +sycl::event generate_rng(Distribution& distr, + device_engine& engine_, + std::int64_t count, + Type* dst, + const event_vector& deps) { + switch (engine_.get_device_engine_base_ptr()->get_engine_type()) { + case engine_type::mt2203: { + auto& device_engine = + *(dynamic_cast(engine_.get_device_engine_base_ptr().get()))->get(); + return oneapi::mkl::rng::generate(distr, device_engine, count, dst, deps); + } + case engine_type::mcg59: { + auto& device_engine = + *(dynamic_cast(engine_.get_device_engine_base_ptr().get()))->get(); + return oneapi::mkl::rng::generate(distr, device_engine, count, dst, deps); + } + case engine_type::mrg32k3a: { + auto& device_engine = + *(dynamic_cast(engine_.get_device_engine_base_ptr().get()))->get(); + return oneapi::mkl::rng::generate(distr, device_engine, count, dst, deps); + } + case engine_type::philox4x32x10: { + auto& device_engine = + *(dynamic_cast(engine_.get_device_engine_base_ptr().get()))->get(); + return oneapi::mkl::rng::generate(distr, device_engine, count, dst, deps); + } + case engine_type::mt19937: { + auto& device_engine = + *(dynamic_cast(engine_.get_device_engine_base_ptr().get()))->get(); + return oneapi::mkl::rng::generate(distr, device_engine, count, dst, deps); + } + default: throw std::runtime_error("Unsupported engine type in generate_rng"); + } +} + +/// Generates uniformly distributed random numbers on the GPU. +/// @tparam Type The data type of the generated numbers. +/// @param[in] queue The SYCL queue for device execution. +/// @param[in] count The number of random numbers to generate. +/// @param[out] dst Pointer to the output buffer. +/// @param[in] engine_ Reference to the device engine. +/// @param[in] a The lower bound of the uniform distribution. +/// @param[in] b The upper bound of the uniform distribution. +/// @param[in] deps Dependencies for the SYCL event. +template +sycl::event uniform(sycl::queue& queue, + std::int64_t count, + Type* dst, + device_engine& engine_, + Type a, + Type b, + const event_vector& deps) { + if (sycl::get_pointer_type(dst, engine_.get_queue().get_context()) == sycl::usm::alloc::host) { + throw domain_error(dal::detail::error_messages::unsupported_data_type()); + } + oneapi::mkl::rng::uniform distr(a, b); + engine_.skip_ahead_cpu(count); + auto event = generate_rng(distr, engine_, count, dst, deps); + return event; +} + +/// Generates a random permutation of elements without replacement on the GPU. +/// @tparam Type The data type of the elements. +/// @param[in] queue The SYCL queue for device execution. +/// @param[in] count The number of elements to generate. +/// @param[out] dst Pointer to the output buffer. +/// @param[out] buffer Temporary buffer used for computations. +/// @param[in] engine_ Reference to the device engine. +/// @param[in] a The lower bound of the range. +/// @param[in] b The upper bound of the range. +/// @param[in] deps Dependencies for the SYCL event. +template +sycl::event uniform_without_replacement(sycl::queue& queue, + std::int64_t count, + Type* dst, + Type* buffer, + device_engine& engine_, + Type a, + Type b, + const event_vector& deps) { + if (sycl::get_pointer_type(dst, engine_.get_queue().get_context()) == + sycl::usm::alloc::device) { + throw domain_error(dal::detail::error_messages::unsupported_data_type()); + } + void* state = engine_.get_host_engine_state(); + engine_.skip_ahead_gpu(count); + uniform_dispatcher::uniform_without_replacement_by_cpu(count, dst, buffer, state, a, b); + auto event = queue.submit([&](sycl::handler& h) { + h.depends_on(deps); + }); + return event; +} + +/// Shuffles an array using random swaps on the GPU. +/// @tparam Type The data type of the array elements. +/// @param[in] queue The SYCL queue for device execution. +/// @param[in] count The number of elements to shuffle. +/// @param[in, out] dst Pointer to the array to be shuffled. +/// @param[in] engine_ Reference to the device engine. +/// @param[in] deps Dependencies for the SYCL event. +template +sycl::event shuffle(sycl::queue& queue, + std::int64_t count, + Type* dst, + device_engine& engine_, + const event_vector& deps) { + Type idx[2]; + if (sycl::get_pointer_type(dst, engine_.get_queue().get_context()) == + sycl::usm::alloc::device) { + throw domain_error(dal::detail::error_messages::unsupported_data_type()); + } + void* state = engine_.get_host_engine_state(); + engine_.skip_ahead_gpu(count); + + for (std::int64_t i = 0; i < count; ++i) { + uniform_dispatcher::uniform_by_cpu(2, idx, state, 0, count); + std::swap(dst[idx[0]], dst[idx[1]]); + } + auto event = queue.submit([&](sycl::handler& h) { + h.depends_on(deps); + }); + return event; +} + +/// Partially shuffles the first `top` elements of an array using the Fisher-Yates algorithm. +/// @tparam Type The data type of the array elements. +/// @param[in] queue_ The SYCL queue for device execution. +/// @param[in, out] result_array The array to be partially shuffled. +/// @param[in] top The number of elements to shuffle. +/// @param[in] seed The seed for the engine. +/// @param[in] method The rng engine type. Defaults to `mt19937`. +/// @param[in] deps Dependencies for the SYCL event. +template +sycl::event partial_fisher_yates_shuffle(sycl::queue& queue_, + ndview& result_array, + std::int64_t top, + std::int64_t seed, + engine_type method, + const event_vector& deps) { + device_engine eng_ = device_engine(queue_, seed, method); + const auto casted_top = dal::detail::integral_cast(top); + const std::int64_t count = result_array.get_count(); + const auto casted_count = dal::detail::integral_cast(count); + ONEDAL_ASSERT(casted_count < casted_top); + auto indices_ptr = result_array.get_mutable_data(); + + std::int64_t k = 0; + std::size_t value = 0; + auto state = eng_.get_host_engine_state(); + for (std::size_t i = 0; i < casted_count; i++) { + uniform_dispatcher::uniform_by_cpu(1, &value, state, i, casted_top); + for (std::size_t j = i; j > 0; j--) { + if (value == dal::detail::integral_cast(indices_ptr[j - 1])) { + value = j - 1; + } + } + if (value >= casted_top) + continue; + indices_ptr[i] = dal::detail::integral_cast(value); + k++; + } + ONEDAL_ASSERT(k == count); + auto event = queue_.submit([&](sycl::handler& h) { + h.depends_on(deps); + }); + return event; +} + +#define INSTANTIATE_UNIFORM(F) \ + template ONEDAL_EXPORT sycl::event uniform(sycl::queue& queue, \ + std::int64_t count_, \ + F* dst, \ + device_engine& engine_, \ + F a, \ + F b, \ + const event_vector& deps); + +INSTANTIATE_UNIFORM(float) +INSTANTIATE_UNIFORM(double) +INSTANTIATE_UNIFORM(std::int32_t) + +#define INSTANTIATE_UWR(F) \ + template ONEDAL_EXPORT sycl::event uniform_without_replacement(sycl::queue& queue, \ + std::int64_t count_, \ + F* dst, \ + F* buff, \ + device_engine& engine_, \ + F a, \ + F b, \ + const event_vector& deps); + +INSTANTIATE_UWR(float) +INSTANTIATE_UWR(double) +INSTANTIATE_UWR(std::int32_t) + +#define INSTANTIATE_SHUFFLE(F) \ + template ONEDAL_EXPORT sycl::event shuffle(sycl::queue& queue, \ + std::int64_t count_, \ + F* dst, \ + device_engine& engine_, \ + const event_vector& deps); + +INSTANTIATE_SHUFFLE(std::int32_t) + +#define INSTANTIATE_PARTIAL_SHUFFLE(F) \ + template ONEDAL_EXPORT sycl::event partial_fisher_yates_shuffle(sycl::queue& queue, \ + ndview& a, \ + std::int64_t top, \ + std::int64_t seed, \ + engine_type method, \ + const event_vector& deps); + +INSTANTIATE_PARTIAL_SHUFFLE(std::int32_t) +INSTANTIATE_PARTIAL_SHUFFLE(std::int64_t) + +} // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/rng/host_engine.hpp b/cpp/oneapi/dal/backend/primitives/rng/host_engine.hpp new file mode 100644 index 00000000000..7d6356196bd --- /dev/null +++ b/cpp/oneapi/dal/backend/primitives/rng/host_engine.hpp @@ -0,0 +1,194 @@ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +#pragma once + +#include +#include +#include +#include +#include + +#include "oneapi/dal/backend/primitives/ndarray.hpp" +#include "oneapi/dal/backend/primitives/rng/utils.hpp" +#include "oneapi/dal/backend/primitives/rng/rng_types.hpp" + +namespace oneapi::dal::backend::primitives { + +/// A class that provides an interface for random number generation on the host (CPU) only. +/// +/// This class serves as a wrapper for host-based random number generators (RNGs), supporting multiple engine +/// types for flexible and efficient random number generation on CPU. It abstracts the underlying engine +/// implementation and provides an interface to manage and retrieve the engine's state. +/// +/// @note The class only supports host-based RNG and does not require a SYCL queue or device context. +class host_engine { +public: + /// @param[in] seed The initial seed for the random number generator. Defaults to `777`. + /// @param[in] method The engine method. Defaults to `engine_type::mt2203`. + host_engine(std::int64_t seed = 777, engine_type method = engine_type::mt2203) { + switch (method) { + case engine_type::mt2203: + host_engine_ = daal::algorithms::engines::mt2203::Batch<>::create(seed); + break; + case engine_type::mcg59: + host_engine_ = daal::algorithms::engines::mcg59::Batch<>::create(seed); + break; + case engine_type::mrg32k3a: + host_engine_ = daal::algorithms::engines::mrg32k3a::Batch<>::create(seed); + break; + case engine_type::philox4x32x10: + host_engine_ = daal::algorithms::engines::philox4x32x10::Batch<>::create(seed); + break; + case engine_type::mt19937: + host_engine_ = daal::algorithms::engines::mt19937::Batch<>::create(seed); + break; + default: throw std::invalid_argument("Unsupported engine type 1"); + } + impl_ = + dynamic_cast(host_engine_.get()); + if (!impl_) { + throw std::domain_error("RNG engine is not supported"); + } + } + + explicit host_engine(const daal::algorithms::engines::EnginePtr& eng) : host_engine_(eng) { + impl_ = dynamic_cast(eng.get()); + if (!impl_) { + throw domain_error(dal::detail::error_messages::rng_engine_is_not_supported()); + } + } + + /// Assignment operator. + host_engine& operator=(const daal::algorithms::engines::EnginePtr& eng) { + host_engine_ = eng; + impl_ = dynamic_cast(eng.get()); + if (!impl_) { + throw domain_error(dal::detail::error_messages::rng_engine_is_not_supported()); + } + + return *this; + } + + // Copy constructor + host_engine(const host_engine&) = default; + + // Copy assignment operator + host_engine& operator=(const host_engine&) = default; + + // Move constructor + host_engine(host_engine&& other) noexcept = default; + + // Move assignment operator + host_engine& operator=(host_engine&& other) noexcept = default; + + /// Destructor. + ~host_engine() = default; + + /// Retrieves the state of the host rng engine(DAAL). + /// @return Pointer to the host engine state. + void* get_host_engine_state() const { + return impl_->getState(); + } + +private: + daal::algorithms::engines::EnginePtr host_engine_; + daal::algorithms::engines::internal::BatchBaseImpl* impl_; +}; + +/// Generates uniformly distributed random numbers on the CPU. +/// @tparam Type The data type of the generated numbers. +/// @param[in] count The number of random numbers to generate. +/// @param[out] dst Pointer to the output buffer. +/// @param[in] engine_ Reference to the device engine. +/// @param[in] a The lower bound of the uniform distribution. +/// @param[in] b The upper bound of the uniform distribution. +template +void uniform(std::int64_t count, Type* dst, host_engine& host_engine, Type a, Type b) { + auto state = host_engine.get_host_engine_state(); + uniform_dispatcher::uniform_by_cpu(count, dst, state, a, b); +} + +/// Generates a random permutation of elements without replacement on the CPU. +/// @tparam Type The data type of the elements. +/// @param[in] count The number of elements to generate. +/// @param[out] dst Pointer to the output buffer. +/// @param[out] buffer Temporary buffer used for computations. +/// @param[in] engine_ Reference to the device engine. +/// @param[in] a The lower bound of the range. +/// @param[in] b The upper bound of the range. +template +void uniform_without_replacement(std::int64_t count, + Type* dst, + Type* buffer, + host_engine host_engine, + Type a, + Type b) { + auto state = host_engine.get_host_engine_state(); + uniform_dispatcher::uniform_without_replacement_by_cpu(count, dst, buffer, state, a, b); +} + +/// Shuffles an array using random swaps on the CPU. +/// @tparam Type The data type of the array elements. +/// @param[in] count The number of elements to shuffle. +/// @param[in, out] dst Pointer to the array to be shuffled. +/// @param[in] engine_ Reference to the device engine. +template >> +void shuffle(std::int64_t count, Type* dst, host_engine host_engine) { + auto state = host_engine.get_host_engine_state(); + Type idx[2]; + for (std::int64_t i = 0; i < count; ++i) { + uniform_dispatcher::uniform_by_cpu(2, idx, state, 0, count); + std::swap(dst[idx[0]], dst[idx[1]]); + } +} + +/// Shuffles an array using random swaps on the CPU. +/// @tparam Type The data type of the array elements. +/// @param[in] count The number of elements to shuffle. +/// @param[in, out] dst Pointer to the array to be shuffled. +/// @param[in] engine_ Reference to the device engine. +template +void partial_fisher_yates_shuffle(ndview& result_array, + std::int64_t top, + std::int64_t seed, + engine_type method = engine_type::mt19937) { + host_engine eng_ = host_engine(seed, method); + const auto casted_top = dal::detail::integral_cast(top); + const std::int64_t count = result_array.get_count(); + const auto casted_count = dal::detail::integral_cast(count); + ONEDAL_ASSERT(casted_count < casted_top); + auto indices_ptr = result_array.get_mutable_data(); + + std::int64_t k = 0; + std::size_t value = 0; + auto state = eng_.get_host_engine_state(); + for (std::size_t i = 0; i < casted_count; i++) { + uniform_dispatcher::uniform_by_cpu(1, &value, state, i, casted_top); + for (std::size_t j = i; j > 0; j--) { + if (value == dal::detail::integral_cast(indices_ptr[j - 1])) { + value = j - 1; + } + } + if (value >= casted_top) + continue; + indices_ptr[i] = dal::detail::integral_cast(value); + k++; + } + ONEDAL_ASSERT(k == count); +} + +} // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/rng/rng_engine_collection.hpp b/cpp/oneapi/dal/backend/primitives/rng/host_engine_collection.hpp similarity index 65% rename from cpp/oneapi/dal/backend/primitives/rng/rng_engine_collection.hpp rename to cpp/oneapi/dal/backend/primitives/rng/host_engine_collection.hpp index 09a5a589141..9b9c241a1a8 100644 --- a/cpp/oneapi/dal/backend/primitives/rng/rng_engine_collection.hpp +++ b/cpp/oneapi/dal/backend/primitives/rng/host_engine_collection.hpp @@ -16,30 +16,31 @@ #pragma once -#include "oneapi/dal/backend/primitives/rng/rng_engine.hpp" +#include "oneapi/dal/backend/primitives/rng/host_engine.hpp" #include namespace oneapi::dal::backend::primitives { -template -class engine_collection { +class host_engine_collection { public: - explicit engine_collection(Size count, std::int64_t seed = 777) + explicit host_engine_collection(std::int64_t count, + std::int64_t seed = 777, + engine_type method = engine_type::mt2203) : count_(count), - engine_(daal::algorithms::engines::mt2203::Batch<>::create(seed)), + engine_(initialize_host_engine(seed, method)), params_(count), technique_(daal::algorithms::engines::internal::family), daal_engine_list_(count) {} template - std::vector operator()(Op&& op) { + std::vector operator()(Op&& op) { daal::services::Status status; - for (Size i = 0; i < count_; ++i) { + for (std::int64_t i = 0; i < count_; ++i) { op(i, params_.nSkip[i]); } select_parallelization_technique(technique_); - daal::algorithms::engines::internal::EnginesCollection engine_collection( + daal::algorithms::engines::internal::EnginesCollection host_engine_collection( engine_, technique_, params_, @@ -49,8 +50,8 @@ class engine_collection { dal::backend::interop::status_to_exception(status); } - std::vector engine_list(count_); - for (Size i = 0; i < count_; ++i) { + std::vector engine_list(count_); + for (std::int64_t i = 0; i < count_; ++i) { engine_list[i] = daal_engine_list_[i]; } @@ -59,6 +60,22 @@ class engine_collection { } private: + daal::algorithms::engines::EnginePtr initialize_host_engine(std::int64_t seed, + engine_type method) { + switch (method) { + case engine_type::mt2203: + return daal::algorithms::engines::mt2203::Batch<>::create(seed); + case engine_type::mcg59: return daal::algorithms::engines::mcg59::Batch<>::create(seed); + case engine_type::mrg32k3a: + return daal::algorithms::engines::mrg32k3a::Batch<>::create(seed); + case engine_type::philox4x32x10: + return daal::algorithms::engines::philox4x32x10::Batch<>::create(seed); + case engine_type::mt19937: + return daal::algorithms::engines::mt19937::Batch<>::create(seed); + default: throw std::invalid_argument("Unsupported engine type"); + } + } + void select_parallelization_technique( daal::algorithms::engines::internal::ParallelizationTechnique& technique) { auto daal_engine_impl = @@ -82,7 +99,7 @@ class engine_collection { } private: - Size count_; + std::int64_t count_; daal::algorithms::engines::EnginePtr engine_; daal::algorithms::engines::internal::Params params_; daal::algorithms::engines::internal::ParallelizationTechnique technique_; diff --git a/cpp/oneapi/dal/backend/primitives/rng/partial_shuffle.cpp b/cpp/oneapi/dal/backend/primitives/rng/partial_shuffle.cpp deleted file mode 100644 index 532fe9da069..00000000000 --- a/cpp/oneapi/dal/backend/primitives/rng/partial_shuffle.cpp +++ /dev/null @@ -1,67 +0,0 @@ -/******************************************************************************* -* Copyright 2021 Intel Corporation -* -* Licensed under the Apache License, Version 2.0 (the "License"); -* you may not use this file except in compliance with the License. -* You may obtain a copy of the License at -* -* http://www.apache.org/licenses/LICENSE-2.0 -* -* Unless required by applicable law or agreed to in writing, software -* distributed under the License is distributed on an "AS IS" BASIS, -* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -* See the License for the specific language governing permissions and -* limitations under the License. -*******************************************************************************/ - -#include "oneapi/dal/backend/primitives/rng/partial_shuffle.hpp" -#include "oneapi/dal/backend/primitives/rng/utils.hpp" - -#include - -#include - -namespace oneapi::dal::backend::primitives { - -void partial_fisher_yates_shuffle(ndview& result_array, - std::int64_t top, - std::int64_t seed) { - using msg = dal::detail::error_messages; - daal::algorithms::engines::EnginePtr engine = - daal::algorithms::engines::mt19937::Batch<>::create(seed); - if (engine.get() == nullptr) { - throw internal_error(msg::failed_to_generate_random_numbers()); - } - auto engine_impl = - dynamic_cast(&(*engine)); - if (engine_impl == nullptr) { - throw internal_error(msg::failed_to_generate_random_numbers()); - } - const auto casted_top = dal::detail::integral_cast(top); - const std::int64_t count = result_array.get_count(); - const auto casted_count = dal::detail::integral_cast(count); - ONEDAL_ASSERT(casted_count < casted_top); - auto indices_ptr = result_array.get_mutable_data(); - - std::int64_t k = 0; - std::size_t value = 0; - for (std::size_t i = 0; i < casted_count; i++) { - uniform_dispatcher::uniform_by_cpu(1, &value, engine_impl->getState(), i, casted_top); - for (std::size_t j = i; j > 0; j--) { - if (value == dal::detail::integral_cast(indices_ptr[j - 1])) { - value = j - 1; - } - } - if (value >= casted_top) - continue; - indices_ptr[i] = dal::detail::integral_cast(value); - k++; - } - ONEDAL_ASSERT(k == count); -} - -void partial_fisher_yates_shuffle(ndview& result_array, std::int64_t top) { - partial_fisher_yates_shuffle(result_array, top, 7777); -} - -} // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/rng/rng_engine.hpp b/cpp/oneapi/dal/backend/primitives/rng/rng_engine.hpp deleted file mode 100644 index c8ca3b13ce9..00000000000 --- a/cpp/oneapi/dal/backend/primitives/rng/rng_engine.hpp +++ /dev/null @@ -1,101 +0,0 @@ -/******************************************************************************* -* Copyright 2021 Intel Corporation -* -* Licensed under the Apache License, Version 2.0 (the "License"); -* you may not use this file except in compliance with the License. -* You may obtain a copy of the License at -* -* http://www.apache.org/licenses/LICENSE-2.0 -* -* Unless required by applicable law or agreed to in writing, software -* distributed under the License is distributed on an "AS IS" BASIS, -* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -* See the License for the specific language governing permissions and -* limitations under the License. -*******************************************************************************/ - -#pragma once - -#include - -#include "oneapi/dal/backend/primitives/rng/utils.hpp" - -namespace oneapi::dal::backend::primitives { - -template -class rng { -public: - rng() = default; - ~rng() = default; - - void uniform(Size count, Type* dst, void* state, Type a, Type b) { - uniform_dispatcher::uniform_by_cpu(count, dst, state, a, b); - } - - void uniform_without_replacement(Size count, - Type* dst, - Type* buffer, - void* state, - Type a, - Type b) { - uniform_dispatcher::uniform_without_replacement_by_cpu(count, - dst, - buffer, - state, - a, - b); - } - - template >> - void shuffle(Size count, Type* dst, void* state) { - Type idx[2]; - - for (Size i = 0; i < count; ++i) { - uniform_dispatcher::uniform_by_cpu(2, idx, state, 0, count); - std::swap(dst[idx[0]], dst[idx[1]]); - } - } - -private: - daal::internal::RNGsInst daal_rng_; -}; - -class engine { -public: - explicit engine(std::int64_t seed = 777) - : engine_(daal::algorithms::engines::mt2203::Batch<>::create(seed)) { - impl_ = dynamic_cast(engine_.get()); - if (!impl_) { - throw domain_error(dal::detail::error_messages::rng_engine_is_not_supported()); - } - } - - explicit engine(const daal::algorithms::engines::EnginePtr& eng) : engine_(eng) { - impl_ = dynamic_cast(eng.get()); - if (!impl_) { - throw domain_error(dal::detail::error_messages::rng_engine_is_not_supported()); - } - } - - virtual ~engine() = default; - - engine& operator=(const daal::algorithms::engines::EnginePtr& eng) { - engine_ = eng; - impl_ = dynamic_cast(eng.get()); - if (!impl_) { - throw domain_error(dal::detail::error_messages::rng_engine_is_not_supported()); - } - - return *this; - } - - void* get_state() const { - return impl_->getState(); - } - -private: - daal::algorithms::engines::EnginePtr engine_; - daal::algorithms::engines::internal::BatchBaseImpl* impl_; -}; - -} // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/rng/rng_types.hpp b/cpp/oneapi/dal/backend/primitives/rng/rng_types.hpp new file mode 100644 index 00000000000..2007a9bdf43 --- /dev/null +++ b/cpp/oneapi/dal/backend/primitives/rng/rng_types.hpp @@ -0,0 +1,34 @@ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +#pragma once + +namespace oneapi::dal::backend::primitives { + +/// Enum class representing different random number generation (RNG) engine methods. +/// +/// This enumeration defines the available RNG engines supported by the library. +/// Each engine method corresponds to a specific algorithm for generating random numbers +/// @enum engine_type +/// Enumeration of RNG engine methods: +/// - `mt2203`: Mersenne Twister engine with specific optimizations for parallel environments. +/// - `mcg59`: Multiplicative congruential generator with a modulus of \(2^{59}\). +/// - `mt19937`: Standard Mersenne Twister engine with a period of \(2^{19937} - 1\). +/// - `mrg32k3a`: Combined multiple recursive generator with a period of \(2^{191}\). +/// - `philox4x32x10`: Counter-based RNG engine optimized for parallel computations. +enum class engine_type { mt2203, mcg59, mt19937, mrg32k3a, philox4x32x10 }; + +} // namespace oneapi::dal::backend::primitives diff --git a/cpp/oneapi/dal/backend/primitives/rng/test/rng_dpc.cpp b/cpp/oneapi/dal/backend/primitives/rng/test/rng_dpc.cpp new file mode 100644 index 00000000000..49476e9bc65 --- /dev/null +++ b/cpp/oneapi/dal/backend/primitives/rng/test/rng_dpc.cpp @@ -0,0 +1,216 @@ +/******************************************************************************* +* Copyright contributors to the oneDAL project +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ + +#include "oneapi/dal/test/engine/common.hpp" +#include "oneapi/dal/test/engine/fixtures.hpp" +#include "oneapi/dal/test/engine/dataframe.hpp" + +#include "oneapi/dal/backend/primitives/rng/device_engine.hpp" +#include "oneapi/dal/backend/primitives/rng/host_engine.hpp" +#include "oneapi/dal/backend/primitives/rng/rng_types.hpp" +#include "oneapi/dal/backend/primitives/ndarray.hpp" +#include + +#include "oneapi/dal/backend/primitives/rng/utils.hpp" +#include "oneapi/dal/table/common.hpp" + +namespace oneapi::dal::backend::primitives::test { + +namespace te = dal::test::engine; + +class mt2203 {}; +class mcg59 {}; +class mrg32k3a {}; +class mt19937 {}; +class philox4x32x10 {}; + +template +struct engine_map {}; + +template <> +struct engine_map { + constexpr static auto value = engine_type::mt2203; +}; + +template <> +struct engine_map { + constexpr static auto value = engine_type::mcg59; +}; + +template <> +struct engine_map { + constexpr static auto value = engine_type::mrg32k3a; +}; + +template <> +struct engine_map { + constexpr static auto value = engine_type::philox4x32x10; +}; + +template <> +struct engine_map { + constexpr static auto value = engine_type::mt19937; +}; + +template +constexpr auto engine_v = engine_map::value; + +template +class rng_test : public te::float_algo_fixture> { +public: + using DataType = std::tuple_element_t<0, TestType>; + using EngineType = std::tuple_element_t<1, TestType>; + static constexpr auto engine_test_type = engine_v; + + auto get_host_engine(std::int64_t seed) { + auto rng_engine = host_engine(seed, engine_test_type); + return rng_engine; + } + + auto get_device_engine(std::int64_t seed) { + auto rng_engine = device_engine(this->get_queue(), seed, engine_test_type); + return rng_engine; + } + + auto allocate_array_host(std::int64_t elem_count) { + auto arr_host = ndarray::empty({ elem_count }); + return arr_host; + } + + auto allocate_array_device(std::int64_t elem_count) { + auto& q = this->get_queue(); + auto arr_gpu = ndarray::empty(q, { elem_count }, sycl::usm::alloc::device); + return arr_gpu; + } + + void check_results(const ndarray& arr_1, const ndarray& arr_2) { + const auto arr_1_host = arr_1.to_host(this->get_queue()); + const DataType* val_arr_1_host_ptr = arr_1_host.get_data(); + + const auto arr_2_host = arr_2.to_host(this->get_queue()); + const DataType* val_arr_2_host_ptr = arr_2_host.get_data(); + + for (std::int64_t el = 0; el < arr_2_host.get_count(); el++) { + // Due to MKL inside generates floats on GPU and doubles on CPU, it makes sense to add minor eps. + REQUIRE(abs(val_arr_1_host_ptr[el] - val_arr_2_host_ptr[el]) < 0.01); + } + } +}; + +using rng_types = COMBINE_TYPES((float, double), (mt2203, mt19937, mcg59, mrg32k3a, philox4x32x10)); + +TEMPLATE_LIST_TEST_M(rng_test, "rng cpu vs gpu", "[rng]", rng_types) { + SKIP_IF(this->get_policy().is_cpu()); + SKIP_IF(this->not_float64_friendly()); + using Float = std::tuple_element_t<0, TestType>; + + std::int64_t elem_count = GENERATE_COPY(10, 777, 10000, 50000); + std::int64_t seed = GENERATE_COPY(777, 999); + + auto arr_gpu = this->allocate_array_device(elem_count); + auto arr_host = this->allocate_array_host(elem_count); + auto arr_gpu_ptr = arr_gpu.get_mutable_data(); + auto arr_host_ptr = arr_host.get_mutable_data(); + + auto rng_engine = this->get_device_engine(seed); + auto rng_engine_ = this->get_device_engine(seed); + + uniform(elem_count, arr_host_ptr, rng_engine, 0, elem_count); + uniform(this->get_queue(), elem_count, arr_gpu_ptr, rng_engine_, 0, elem_count) + .wait_and_throw(); + + this->check_results(arr_gpu, arr_host); +} + +using rng_types_skip_ahead_support = COMBINE_TYPES((float, double), + (mt19937, mcg59, mrg32k3a, philox4x32x10)); + +TEMPLATE_LIST_TEST_M(rng_test, "mixed rng cpu skip", "[rng]", rng_types_skip_ahead_support) { + SKIP_IF(this->get_policy().is_cpu()); + SKIP_IF(this->not_float64_friendly()); + using Float = std::tuple_element_t<0, TestType>; + + std::int64_t elem_count = GENERATE_COPY(10, 777, 10000, 100000); + std::int64_t seed = GENERATE_COPY(777, 999); + + auto arr_host_init_1 = this->allocate_array_host(elem_count); + auto arr_host_init_2 = this->allocate_array_host(elem_count); + + auto arr_gpu = this->allocate_array_device(elem_count); + auto arr_host = this->allocate_array_host(elem_count); + + auto arr_host_init_1_ptr = arr_host_init_1.get_mutable_data(); + auto arr_host_init_2_ptr = arr_host_init_2.get_mutable_data(); + auto arr_gpu_ptr = arr_gpu.get_mutable_data(); + auto arr_host_ptr = arr_host.get_mutable_data(); + + auto rng_engine = this->get_device_engine(seed); + auto rng_engine_2 = this->get_device_engine(seed); + + uniform(elem_count, arr_host_init_1_ptr, rng_engine, 0, elem_count); + uniform(elem_count, arr_host_init_2_ptr, rng_engine_2, 0, elem_count); + + uniform(this->get_queue(), elem_count, arr_gpu_ptr, rng_engine, 0, elem_count) + .wait_and_throw(); + uniform(elem_count, arr_host_ptr, rng_engine_2, 0, elem_count); + + this->check_results(arr_host_init_1, arr_host_init_2); + this->check_results(arr_gpu, arr_host); +} + +TEMPLATE_LIST_TEST_M(rng_test, "mixed rng gpu skip", "[rng]", rng_types_skip_ahead_support) { + SKIP_IF(this->get_policy().is_cpu()); + SKIP_IF(this->not_float64_friendly()); + using Float = std::tuple_element_t<0, TestType>; + + std::int64_t elem_count = GENERATE_COPY(10, 100, 777, 10000); + std::int64_t seed = GENERATE_COPY(1, 777, 999); + + auto arr_device_init_1 = this->allocate_array_device(elem_count); + auto arr_device_init_2 = this->allocate_array_device(elem_count); + + auto arr_gpu = this->allocate_array_device(elem_count); + auto arr_host = this->allocate_array_host(elem_count); + + auto arr_device_init_1_ptr = arr_device_init_1.get_mutable_data(); + auto arr_device_init_2_ptr = arr_device_init_2.get_mutable_data(); + auto arr_gpu_ptr = arr_gpu.get_mutable_data(); + auto arr_host_ptr = arr_host.get_mutable_data(); + + auto rng_engine = this->get_device_engine(seed); + auto rng_engine_2 = this->get_device_engine(seed); + + uniform(this->get_queue(), elem_count, arr_device_init_1_ptr, rng_engine, 0, elem_count) + .wait_and_throw(); + uniform(this->get_queue(), + elem_count, + arr_device_init_2_ptr, + rng_engine_2, + 0, + elem_count) + .wait_and_throw(); + + uniform(this->get_queue(), elem_count, arr_gpu_ptr, rng_engine, 0, elem_count) + .wait_and_throw(); + uniform(elem_count, arr_host_ptr, rng_engine_2, 0, elem_count); + + this->check_results(arr_device_init_1, arr_device_init_2); + this->check_results(arr_gpu, arr_host); +} + +//TODO: add engine collection test + separate host_engine tests + +} // namespace oneapi::dal::backend::primitives::test diff --git a/docs/source/daal/algorithms/engines/index.rst b/docs/source/daal/algorithms/engines/index.rst index 9def0af4c81..e4435c6547c 100644 --- a/docs/source/daal/algorithms/engines/index.rst +++ b/docs/source/daal/algorithms/engines/index.rst @@ -109,6 +109,8 @@ These methods are represented with member functions of classes that represent fu .. toctree:: :maxdepth: 1 - mt19937.rst mcg59.rst + mt19937.rst mt2203.rst + mrg32k3a.rst + philox4x32x10.rst diff --git a/docs/source/daal/algorithms/engines/mrg32k3a.rst b/docs/source/daal/algorithms/engines/mrg32k3a.rst new file mode 100644 index 00000000000..2b00c990588 --- /dev/null +++ b/docs/source/daal/algorithms/engines/mrg32k3a.rst @@ -0,0 +1,62 @@ +.. Copyright contributors to the oneDAL project +.. +.. Licensed under the Apache License, Version 2.0 (the "License"); +.. you may not use this file except in compliance with the License. +.. You may obtain a copy of the License at +.. +.. http://www.apache.org/licenses/LICENSE-2.0 +.. +.. Unless required by applicable law or agreed to in writing, software +.. distributed under the License is distributed on an "AS IS" BASIS, +.. WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +.. See the License for the specific language governing permissions and +.. limitations under the License. + +mrg32k3a +======== + +The engine is based on a 32-bit combined multiple recursive generator +with two components of order 3, optimized for batch processing. + +.. rubric:: Subsequence selection methods support + +skipAhead (nskip) + Supported +leapfrog (threadIdx, nThreads) + Supported + +Batch Processing +**************** + +MRG32k3a engine needs the initial condition (``seed``) for state initialization. +The seed can be either an integer scalar or a vector of :math:`p` integer elements, the inputs to the respective engine constructors. + +.. rubric:: Algorithm Parameters + +MRG32k3a engine has the following parameters: + +.. tabularcolumns:: |\Y{0.2}|\Y{0.2}|\Y{0.6}| + +.. list-table:: Algorithm Parameters for mcg58 engine (Batch Processing) + :header-rows: 1 + :widths: 10 20 30 + :align: left + :class: longtable + + * - Parameter + - Default Value + - Description + * - ``algorithmFPType`` + - ``float`` + - The floating-point type that the algorithm uses for intermediate computations. Can be ``float`` or ``double``. + * - ``method`` + - ``defaultDense`` + - Performance-oriented computation method; the only method supported by the algorithm. + * - ``seed`` + - + - :math:`777` for a scalar seed + - NA for a vector seed + - Initial condition for state initialization, scalar or vector: + + - Scalar, value of ``size_t`` type + - Vector, pointer to ``HomogenNumericTable`` of size :math:`1 \times p` diff --git a/docs/source/daal/algorithms/engines/philox4x32x10.rst b/docs/source/daal/algorithms/engines/philox4x32x10.rst new file mode 100644 index 00000000000..ee315261604 --- /dev/null +++ b/docs/source/daal/algorithms/engines/philox4x32x10.rst @@ -0,0 +1,62 @@ +.. Copyright contributors to the oneDAL project +.. +.. Licensed under the Apache License, Version 2.0 (the "License"); +.. you may not use this file except in compliance with the License. +.. You may obtain a copy of the License at +.. +.. http://www.apache.org/licenses/LICENSE-2.0 +.. +.. Unless required by applicable law or agreed to in writing, software +.. distributed under the License is distributed on an "AS IS" BASIS, +.. WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +.. See the License for the specific language governing permissions and +.. limitations under the License. + +philox4x32x10 +============= + +Implementation of the Philox4x32-10 engine: a counter-based pseudorandom number generator (PRNG) +that uses 4x32-bit keys and performs 10 rounds of mixing to produce high-quality randomness. + +.. rubric:: Subsequence selection methods support + +skipAhead (nskip) + Supported +leapfrog (threadIdx, nThreads) + Supported + +Batch Processing +**************** + +Philox4x32x10 engine needs the initial condition (``seed``) for state initialization. +The seed can be either an integer scalar or a vector of :math:`p` integer elements, the inputs to the respective engine constructors. + +.. rubric:: Algorithm Parameters + +Philox4x32x10 engine has the following parameters: + +.. tabularcolumns:: |\Y{0.2}|\Y{0.2}|\Y{0.6}| + +.. list-table:: Algorithm Parameters for mcg58 engine (Batch Processing) + :header-rows: 1 + :widths: 10 20 30 + :align: left + :class: longtable + + * - Parameter + - Default Value + - Description + * - ``algorithmFPType`` + - ``float`` + - The floating-point type that the algorithm uses for intermediate computations. Can be ``float`` or ``double``. + * - ``method`` + - ``defaultDense`` + - Performance-oriented computation method; the only method supported by the algorithm. + * - ``seed`` + - + - :math:`777` for a scalar seed + - NA for a vector seed + - Initial condition for state initialization, scalar or vector: + + - Scalar, value of ``size_t`` type + - Vector, pointer to ``HomogenNumericTable`` of size :math:`1 \times p` diff --git a/makefile.lst b/makefile.lst index 92dc52ff521..b042ede80a7 100755 --- a/makefile.lst +++ b/makefile.lst @@ -65,7 +65,7 @@ multiclassclassifier += classifier k_nearest_neighbors += engines classifier logistic_regression += classifier optimization_solver objective_function engines implicit_als += engines distributions -engines += engines/mt19937 engines/mcg59 engines/mt2203 +engines += engines/mt19937 engines/mcg59 engines/mrg32k3a engines/philox4x32x10 engines/mt2203 distributions += distributions/bernoulli distributions/normal distributions/uniform tsne += @@ -95,6 +95,8 @@ CORE.ALGORITHMS.FULL := \ elastic_net \ engines \ engines/mcg59 \ + engines/mrg32k3a \ + engines/philox4x32x10 \ engines/mt19937 \ engines/mt2203 \ em \ @@ -309,6 +311,8 @@ JJ.ALGORITHMS := adaboost elastic_net/prediction \ engines \ engines/mcg59 \ + engines/mrg32k3a \ + engines/philox4x32x10 \ engines/mt19937 \ engines/mt2203 \ em_gmm \