diff --git a/dpnp/backend/include/dpnp_iface_random.hpp b/dpnp/backend/include/dpnp_iface_random.hpp index aea8fb83730e..e6396b9e0aa2 100644 --- a/dpnp/backend/include/dpnp_iface_random.hpp +++ b/dpnp/backend/include/dpnp_iface_random.hpp @@ -335,8 +335,8 @@ INP_DLLEXPORT void dpnp_rng_lognormal_c(void* result, const _DataType mean, cons * @param [in] q_ref Reference to SYCL queue. * @param [out] result Output array. * @param [in] ntrial Number of independent trials. - * @param [in] p_vector Probability vector of possible outcomes (k length). - * @param [in] p_vector_size Length of `p_vector`. + * @param [in] p_in Probability of possible outcomes (k length). + * @param [in] p_size Length of `p_in`. * @param [in] size Number of elements in `result` arrays. * @param [in] dep_event_vec_ref Reference to vector of SYCL events. */ @@ -344,14 +344,14 @@ template INP_DLLEXPORT DPCTLSyclEventRef dpnp_rng_multinomial_c(DPCTLSyclQueueRef q_ref, void* result, const int ntrial, - const double* p_vector, - const size_t p_vector_size, + const double* p_in, + const size_t p_size, const size_t size, const DPCTLEventVectorRef dep_event_vec_ref); template INP_DLLEXPORT void dpnp_rng_multinomial_c( - void* result, const int ntrial, const double* p_vector, const size_t p_vector_size, const size_t size); + void* result, const int ntrial, const double* p_in, const size_t p_size, const size_t size); /** * @ingroup BACKEND_RANDOM_API @@ -360,10 +360,10 @@ INP_DLLEXPORT void dpnp_rng_multinomial_c( * @param [in] q_ref Reference to SYCL queue. * @param [out] result Output array. * @param [in] dimen Dimension of output random vectors. - * @param [in] mean_vector Mean vector a of dimension. - * @param [in] mean_vector_size Length of `mean_vector`. - * @param [in] cov_vector Variance-covariance matrix. - * @param [in] cov_vector_size Length of `cov_vector`. + * @param [in] mean_in Mean arry of dimension. + * @param [in] mean_size Length of `mean_in`. + * @param [in] cov Variance-covariance matrix. + * @param [in] cov_size Length of `cov_in`. * @param [in] size Number of elements in `result` arrays. * @param [in] dep_event_vec_ref Reference to vector of SYCL events. */ @@ -371,20 +371,20 @@ template INP_DLLEXPORT DPCTLSyclEventRef dpnp_rng_multivariate_normal_c(DPCTLSyclQueueRef q_ref, void* result, const int dimen, - const double* mean_vector, - const size_t mean_vector_size, - const double* cov_vector, - const size_t cov_vector_size, + const double* mean_in, + const size_t mean_size, + const double* cov_in, + const size_t cov_size, const size_t size, const DPCTLEventVectorRef dep_event_vec_ref); template INP_DLLEXPORT void dpnp_rng_multivariate_normal_c(void* result, const int dimen, - const double* mean_vector, - const size_t mean_vector_size, - const double* cov_vector, - const size_t cov_vector_size, + const double* mean_in, + const size_t mean_size, + const double* cov_in, + const size_t cov_size, const size_t size); /** diff --git a/dpnp/backend/kernels/dpnp_krnl_random.cpp b/dpnp/backend/kernels/dpnp_krnl_random.cpp index 67d34fcff1cf..afc5df8187d3 100644 --- a/dpnp/backend/kernels/dpnp_krnl_random.cpp +++ b/dpnp/backend/kernels/dpnp_krnl_random.cpp @@ -101,9 +101,10 @@ DPCTLSyclEventRef dpnp_rng_beta_c(DPCTLSyclQueueRef q_ref, mkl_rng::beta<_DataType> distribution(a, b, displacement, scalefactor); // perform generation auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); - return event_ref; + event_ref = reinterpret_cast(&event_out); + + return DPCTLEvent_Copy(event_ref); } template @@ -157,23 +158,25 @@ DPCTLSyclEventRef dpnp_rng_binomial_c(DPCTLSyclQueueRef q_ref, if (ntrial == 0 || p == 0) { - dpnp_zeros_c<_DataType>(result, size); + event_ref = dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); } else if (p == 1) { _DataType* fill_value = reinterpret_cast<_DataType*>(sycl::malloc_shared(sizeof(_DataType), q)); fill_value[0] = static_cast<_DataType>(ntrial); - dpnp_initval_c<_DataType>(result, fill_value, size); - sycl::free(fill_value, q); + + event_ref = dpnp_initval_c<_DataType>(q_ref, result, fill_value, size, dep_event_vec_ref); + DPCTLEvent_Wait(event_ref); + dpnp_memory_free_c(q_ref, fill_value); } else { _DataType* result1 = reinterpret_cast<_DataType*>(result); mkl_rng::binomial<_DataType> distribution(ntrial, p); auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_ref = reinterpret_cast(&event_out); } - return event_ref; + return DPCTLEvent_Copy(event_ref); } @@ -225,9 +228,10 @@ DPCTLSyclEventRef dpnp_rng_chisquare_c(DPCTLSyclQueueRef q_ref, mkl_rng::chi_square<_DataType> distribution(df); auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); - return event_ref; + event_ref = reinterpret_cast(&event_out); + + return DPCTLEvent_Copy(event_ref); } template @@ -280,9 +284,10 @@ DPCTLSyclEventRef dpnp_rng_exponential_c(DPCTLSyclQueueRef q_ref, mkl_rng::exponential<_DataType> distribution(a, beta); // perform generation auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); - return event_ref; + event_ref = reinterpret_cast(&event_out); + + return DPCTLEvent_Copy(event_ref); } template @@ -398,6 +403,7 @@ DPCTLSyclEventRef dpnp_rng_gamma_c(DPCTLSyclQueueRef q_ref, (void)dep_event_vec_ref; DPCTLSyclEventRef event_ref = nullptr; + sycl::event event_out; if (!size || result == nullptr) { @@ -408,7 +414,7 @@ DPCTLSyclEventRef dpnp_rng_gamma_c(DPCTLSyclQueueRef q_ref, if (shape == 0.0 || scale == 0.0) { - dpnp_zeros_c<_DataType>(result, size); + event_ref = dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); } else { @@ -416,11 +422,11 @@ DPCTLSyclEventRef dpnp_rng_gamma_c(DPCTLSyclQueueRef q_ref, const _DataType a = (_DataType(0.0)); mkl_rng::gamma<_DataType> distribution(shape, a, scale); - auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); + event_ref = reinterpret_cast(&event_out); } - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -473,9 +479,9 @@ DPCTLSyclEventRef dpnp_rng_gaussian_c(DPCTLSyclQueueRef q_ref, mkl_rng::gaussian<_DataType> distribution(mean, stddev); // perform generation auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_ref = reinterpret_cast(&event_out); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -517,6 +523,7 @@ DPCTLSyclEventRef dpnp_rng_geometric_c(DPCTLSyclQueueRef q_ref, (void)dep_event_vec_ref; DPCTLSyclEventRef event_ref = nullptr; + sycl::event event_out; if (!size || !result) { @@ -527,17 +534,17 @@ DPCTLSyclEventRef dpnp_rng_geometric_c(DPCTLSyclQueueRef q_ref, if (p == 1.0) { - dpnp_ones_c<_DataType>(result, size); + event_ref = dpnp_ones_c<_DataType>(q_ref, result, size, dep_event_vec_ref); } else { _DataType* result1 = reinterpret_cast<_DataType*>(result); mkl_rng::geometric<_DataType> distribution(p); // perform generation - auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); + event_ref = reinterpret_cast(&event_out); } - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -575,6 +582,7 @@ DPCTLSyclEventRef dpnp_rng_gumbel_c(DPCTLSyclQueueRef q_ref, (void)dep_event_vec_ref; DPCTLSyclEventRef event_ref = nullptr; + sycl::event event_out; if (!size || !result) { @@ -587,8 +595,10 @@ DPCTLSyclEventRef dpnp_rng_gumbel_c(DPCTLSyclQueueRef q_ref, { _DataType* fill_value = reinterpret_cast<_DataType*>(sycl::malloc_shared(sizeof(_DataType), q)); fill_value[0] = static_cast<_DataType>(loc); - dpnp_initval_c<_DataType>(result, fill_value, size); - sycl::free(fill_value, q); + + event_ref = dpnp_initval_c<_DataType>(q_ref, result, fill_value, size, dep_event_vec_ref); + DPCTLEvent_Wait(event_ref); + dpnp_memory_free_c(q_ref, fill_value); } else { @@ -600,11 +610,10 @@ DPCTLSyclEventRef dpnp_rng_gumbel_c(DPCTLSyclQueueRef q_ref, mkl_rng::gumbel<_DataType> distribution(negloc, scale); auto event_distribution = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - sycl::event prod_event; - prod_event = mkl_blas::scal(q, size, alpha, result1, incx, {event_distribution}); - prod_event.wait(); + event_out = mkl_blas::scal(q, size, alpha, result1, incx, {event_distribution}); + event_ref = reinterpret_cast(&event_out); } - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -645,6 +654,7 @@ DPCTLSyclEventRef dpnp_rng_hypergeometric_c(DPCTLSyclQueueRef q_ref, (void)dep_event_vec_ref; DPCTLSyclEventRef event_ref = nullptr; + sycl::event event_out; if (!size || !result) { @@ -655,23 +665,26 @@ DPCTLSyclEventRef dpnp_rng_hypergeometric_c(DPCTLSyclQueueRef q_ref, if (m == 0) { - dpnp_zeros_c<_DataType>(result, size); + event_ref = dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); } else if (l == m) { _DataType* fill_value = reinterpret_cast<_DataType*>(sycl::malloc_shared(sizeof(_DataType), q)); fill_value[0] = static_cast<_DataType>(s); - dpnp_initval_c<_DataType>(result, fill_value, size); - sycl::free(fill_value, q); + + event_ref = dpnp_initval_c<_DataType>(q_ref, result, fill_value, size, dep_event_vec_ref); + DPCTLEvent_Wait(event_ref); + dpnp_memory_free_c(q_ref, fill_value); } else { _DataType* result1 = reinterpret_cast<_DataType*>(result); mkl_rng::hypergeometric<_DataType> distribution(l, s, m); - auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); + event_ref = reinterpret_cast(&event_out); + } - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -717,6 +730,7 @@ DPCTLSyclEventRef dpnp_rng_laplace_c(DPCTLSyclQueueRef q_ref, (void)dep_event_vec_ref; DPCTLSyclEventRef event_ref = nullptr; + sycl::event event_out; if (!size || !result) { @@ -727,17 +741,17 @@ DPCTLSyclEventRef dpnp_rng_laplace_c(DPCTLSyclQueueRef q_ref, if (scale == 0.0) { - dpnp_zeros_c<_DataType>(result, size); + event_ref = dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); } else { _DataType* result1 = reinterpret_cast<_DataType*>(result); mkl_rng::laplace<_DataType> distribution(loc, scale); // perform generation - auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); + event_ref = reinterpret_cast(&event_out); } - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -808,9 +822,9 @@ DPCTLSyclEventRef dpnp_rng_logistic_c(DPCTLSyclQueueRef q_ref, cgh.parallel_for>(gws, kernel_parallel_for_func); }; auto event = q.submit(kernel_func); - event.wait(); + event_ref = reinterpret_cast(&event); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -850,6 +864,7 @@ DPCTLSyclEventRef dpnp_rng_lognormal_c(DPCTLSyclQueueRef q_ref, (void)dep_event_vec_ref; DPCTLSyclEventRef event_ref = nullptr; + sycl::event event_out; if (!size || !result) { @@ -864,8 +879,10 @@ DPCTLSyclEventRef dpnp_rng_lognormal_c(DPCTLSyclQueueRef q_ref, { _DataType* fill_value = reinterpret_cast<_DataType*>(sycl::malloc_shared(sizeof(_DataType), q)); fill_value[0] = static_cast<_DataType>(std::exp(mean + (stddev * stddev) / 2)); - dpnp_initval_c<_DataType>(result, fill_value, size); - sycl::free(fill_value, q); + + event_ref = dpnp_initval_c<_DataType>(q_ref, result, fill_value, size, dep_event_vec_ref); + DPCTLEvent_Wait(event_ref); + dpnp_memory_free_c(q_ref, fill_value); } else { @@ -873,10 +890,10 @@ DPCTLSyclEventRef dpnp_rng_lognormal_c(DPCTLSyclQueueRef q_ref, const _DataType scalefactor = _DataType(1.0); mkl_rng::lognormal<_DataType> distribution(mean, stddev, displacement, scalefactor); - auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); + event_ref = reinterpret_cast(&event_out); } - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -911,8 +928,8 @@ template DPCTLSyclEventRef dpnp_rng_multinomial_c(DPCTLSyclQueueRef q_ref, void* result, const int ntrial, - const double* p_vector, - const size_t p_vector_size, + void* p_in, + const size_t p_size, const size_t size, const DPCTLEventVectorRef dep_event_vec_ref) { @@ -920,6 +937,7 @@ DPCTLSyclEventRef dpnp_rng_multinomial_c(DPCTLSyclQueueRef q_ref, (void)dep_event_vec_ref; DPCTLSyclEventRef event_ref = nullptr; + sycl::event event_out; if (!size || !result || (ntrial < 0)) { @@ -930,56 +948,58 @@ DPCTLSyclEventRef dpnp_rng_multinomial_c(DPCTLSyclQueueRef q_ref, if (ntrial == 0) { - dpnp_zeros_c<_DataType>(result, size); + event_ref = dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); } else { - std::vector p(p_vector, p_vector + p_vector_size); + DPNPC_ptr_adapter p_ptr(q_ref, p_in, p_size, true); + const double* p = p_ptr.get_ptr(); + std::vector p_vec(p, p + p_size); // size = size // `result` is a array for random numbers // `size` is a `result`'s len. `size = n * p.size()` // `n` is a number of random values to be generated. - size_t n = size / p.size(); + size_t n = size / p_vec.size(); size_t is_cpu_queue = dpnp_queue_is_cpu_c(); // math library supports the distribution generation on GPU device with input parameters // which follow the condition - if (is_cpu_queue || (!is_cpu_queue && (p_vector_size >= ((size_t)ntrial * 16)) && (ntrial <= 16))) + if (is_cpu_queue || (!is_cpu_queue && (p_size >= ((size_t)ntrial * 16)) && (ntrial <= 16))) { DPNPC_ptr_adapter result_ptr(q_ref, result, size, false, true); std::int32_t* result1 = result_ptr.get_ptr(); - mkl_rng::multinomial distribution(ntrial, p); + mkl_rng::multinomial distribution(ntrial, p_vec); // perform generation - auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, n, result1); - event_out.wait(); + event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, n, result1); + event_ref = reinterpret_cast(&event_out); } else { DPNPC_ptr_adapter result_ptr(q_ref, result, size, true, true); std::int32_t* result1 = result_ptr.get_ptr(); int errcode = viRngMultinomial( - VSL_RNG_METHOD_MULTINOMIAL_MULTPOISSON, get_rng_stream(), n, result1, ntrial, p_vector_size, p_vector); + VSL_RNG_METHOD_MULTINOMIAL_MULTPOISSON, get_rng_stream(), n, result1, ntrial, p_size, p); if (errcode != VSL_STATUS_OK) { throw std::runtime_error("DPNP RNG Error: dpnp_rng_multinomial_c() failed."); } } } - return event_ref; + return DPCTLEvent_Copy(event_ref); } template void dpnp_rng_multinomial_c( - void* result, const int ntrial, const double* p_vector, const size_t p_vector_size, const size_t size) + void* result, const int ntrial, void* p_in, const size_t p_size, const size_t size) { DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); DPCTLEventVectorRef dep_event_vec_ref = nullptr; DPCTLSyclEventRef event_ref = dpnp_rng_multinomial_c<_DataType>(q_ref, result, ntrial, - p_vector, - p_vector_size, + p_in, + p_size, size, dep_event_vec_ref); DPCTLEvent_WaitAndThrow(event_ref); @@ -988,7 +1008,7 @@ void dpnp_rng_multinomial_c( template void (*dpnp_rng_multinomial_default_c)(void*, const int, - const double*, + void*, const size_t, const size_t) = dpnp_rng_multinomial_c<_DataType>; @@ -996,7 +1016,7 @@ template DPCTLSyclEventRef (*dpnp_rng_multinomial_ext_c)(DPCTLSyclQueueRef, void*, const int, - const double*, + void*, const size_t, const size_t, const DPCTLEventVectorRef) = dpnp_rng_multinomial_c<_DataType>; @@ -1005,10 +1025,10 @@ template DPCTLSyclEventRef dpnp_rng_multivariate_normal_c(DPCTLSyclQueueRef q_ref, void* result, const int dimen, - const double* mean_vector, - const size_t mean_vector_size, - const double* cov_vector, - const size_t cov_vector_size, + void* mean_in, + const size_t mean_size, + void* cov_in, + const size_t cov_size, const size_t size, const DPCTLEventVectorRef dep_event_vec_ref) { @@ -1024,30 +1044,35 @@ DPCTLSyclEventRef dpnp_rng_multivariate_normal_c(DPCTLSyclQueueRef q_ref, sycl::queue q = *(reinterpret_cast(q_ref)); + DPNPC_ptr_adapter mean_ptr(q_ref, mean_in, mean_size, true); + const double* mean = mean_ptr.get_ptr(); + DPNPC_ptr_adapter cov_ptr(q_ref, cov_in, cov_size, true); + const double* cov = cov_ptr.get_ptr(); + _DataType* result1 = reinterpret_cast<_DataType*>(result); - std::vector mean(mean_vector, mean_vector + mean_vector_size); - std::vector cov(cov_vector, cov_vector + cov_vector_size); + std::vector mean_vec(mean, mean + mean_size); + std::vector cov_vec(cov, cov + cov_size); // `result` is a array for random numbers // `size` is a `result`'s len. // `size1` is a number of random values to be generated for each dimension. size_t size1 = size / dimen; - mkl_rng::gaussian_mv<_DataType> distribution(dimen, mean, cov); + mkl_rng::gaussian_mv<_DataType> distribution(dimen, mean_vec, cov_vec); auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size1, result1); - event_out.wait(); + event_ref = reinterpret_cast(&event_out); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template void dpnp_rng_multivariate_normal_c(void* result, const int dimen, - const double* mean_vector, - const size_t mean_vector_size, - const double* cov_vector, - const size_t cov_vector_size, + void* mean_in, + const size_t mean_size, + void* cov_in, + const size_t cov_size, const size_t size) { DPCTLSyclQueueRef q_ref = reinterpret_cast(&DPNP_QUEUE); @@ -1055,10 +1080,10 @@ void dpnp_rng_multivariate_normal_c(void* result, DPCTLSyclEventRef event_ref = dpnp_rng_multivariate_normal_c<_DataType>(q_ref, result, dimen, - mean_vector, - mean_vector_size, - cov_vector, - cov_vector_size, + mean_in, + mean_size, + cov_in, + cov_size, size, dep_event_vec_ref); DPCTLEvent_WaitAndThrow(event_ref); @@ -1067,23 +1092,22 @@ void dpnp_rng_multivariate_normal_c(void* result, template void (*dpnp_rng_multivariate_normal_default_c)(void*, const int, - const double*, + void*, const size_t, - const double*, + void*, const size_t, const size_t) = dpnp_rng_multivariate_normal_c<_DataType>; template -DPCTLSyclEventRef (*dpnp_rng_multivariate_normal_ext_c)( - DPCTLSyclQueueRef, - void*, - const int, - const double*, - const size_t, - const double*, - const size_t, - const size_t, - const DPCTLEventVectorRef) = dpnp_rng_multivariate_normal_c<_DataType>; +DPCTLSyclEventRef (*dpnp_rng_multivariate_normal_ext_c)(DPCTLSyclQueueRef, + void*, + const int, + void*, + const size_t, + void*, + const size_t, + const size_t, + const DPCTLEventVectorRef) = dpnp_rng_multivariate_normal_c<_DataType>; template DPCTLSyclEventRef dpnp_rng_negative_binomial_c(DPCTLSyclQueueRef q_ref, @@ -1108,9 +1132,9 @@ DPCTLSyclEventRef dpnp_rng_negative_binomial_c(DPCTLSyclQueueRef q_ref, _DataType* result1 = reinterpret_cast<_DataType*>(result); mkl_rng::negative_binomial<_DataType> distribution(a, p); auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_ref = reinterpret_cast(&event_out); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -1159,6 +1183,7 @@ DPCTLSyclEventRef dpnp_rng_noncentral_chisquare_c(DPCTLSyclQueueRef q_ref, (void)dep_event_vec_ref; DPCTLSyclEventRef event_ref = nullptr; + sycl::event event_out; if (!size || !result) { @@ -1209,7 +1234,7 @@ DPCTLSyclEventRef dpnp_rng_noncentral_chisquare_c(DPCTLSyclQueueRef q_ref, lambda = 0.5 * nonc; mkl_rng::poisson poisson_distribution(lambda); - auto event_out = mkl_rng::generate(poisson_distribution, DPNP_RNG_ENGINE, size, pvec); + event_out = mkl_rng::generate(poisson_distribution, DPNP_RNG_ENGINE, size, pvec); event_out.wait(); shape = 0.5 * df; @@ -1287,10 +1312,10 @@ DPCTLSyclEventRef dpnp_rng_noncentral_chisquare_c(DPCTLSyclQueueRef q_ref, loc = sqrt(nonc); mkl_rng::gaussian<_DataType> gaussian_distribution(loc, d_one); auto event_gaussian_distr = mkl_rng::generate(gaussian_distribution, DPNP_RNG_ENGINE, size, result1); - auto event_out = mkl_vm::sqr(q, size, result1, result1, {event_gaussian_distr}, mkl_vm::mode::ha); - event_out.wait(); + event_out = mkl_vm::sqr(q, size, result1, result1, {event_gaussian_distr}, mkl_vm::mode::ha); + event_ref = reinterpret_cast(&event_out); } - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -1347,9 +1372,9 @@ DPCTLSyclEventRef dpnp_rng_normal_c(DPCTLSyclQueueRef q_ref, mkl_rng::gaussian<_DataType> distribution(mean, stddev); // perform generation auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_ref = reinterpret_cast(&event_out); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -1411,9 +1436,9 @@ DPCTLSyclEventRef dpnp_rng_pareto_c(DPCTLSyclQueueRef q_ref, event_out.wait(); event_out = mkl_vm::powx(q, size, result1, neg_rec_alp, result1, no_deps, mkl_vm::mode::ha); - event_out.wait(); + event_ref = reinterpret_cast(&event_out); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -1463,9 +1488,9 @@ DPCTLSyclEventRef dpnp_rng_poisson_c(DPCTLSyclQueueRef q_ref, mkl_rng::poisson<_DataType> distribution(lambda); // perform generation auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_ref = reinterpret_cast(&event_out); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -1522,9 +1547,9 @@ DPCTLSyclEventRef dpnp_rng_power_c(DPCTLSyclQueueRef q_ref, event_out.wait(); event_out = mkl_vm::powx(q, size, result1, neg_rec_alp, result1, no_deps, mkl_vm::mode::ha); - event_out.wait(); + event_ref = reinterpret_cast(&event_out); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -1581,9 +1606,9 @@ DPCTLSyclEventRef dpnp_rng_rayleigh_c(DPCTLSyclQueueRef q_ref, auto exponential_rng_event = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); auto sqrt_event = mkl_vm::sqrt(q, size, result1, result1, {exponential_rng_event}, mkl_vm::mode::ha); auto scal_event = mkl_blas::scal(q, size, scale, result1, 1, {sqrt_event}); - scal_event.wait(); + event_ref = reinterpret_cast(&scal_event); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -1760,9 +1785,9 @@ DPCTLSyclEventRef dpnp_rng_standard_cauchy_c(DPCTLSyclQueueRef q_ref, mkl_rng::cauchy<_DataType> distribution(displacement, scalefactor); // perform generation auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_ref = reinterpret_cast(&event_out); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -1807,9 +1832,9 @@ DPCTLSyclEventRef dpnp_rng_standard_exponential_c(DPCTLSyclQueueRef q_ref, // set displacement a const _DataType beta = (_DataType(1.0)); - dpnp_rng_exponential_c(result, beta, size); + event_ref = dpnp_rng_exponential_c(q_ref, result, beta, size, dep_event_vec_ref); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -1855,9 +1880,9 @@ DPCTLSyclEventRef dpnp_rng_standard_gamma_c(DPCTLSyclQueueRef q_ref, const _DataType scale = _DataType(1.0); - dpnp_rng_gamma_c(result, shape, scale, size); + event_ref = dpnp_rng_gamma_c(q_ref, result, shape, scale, size, dep_event_vec_ref); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -1904,9 +1929,9 @@ DPCTLSyclEventRef dpnp_rng_standard_normal_c(DPCTLSyclQueueRef q_ref, const _DataType mean = _DataType(0.0); const _DataType stddev = _DataType(1.0); - dpnp_rng_normal_c(result, mean, stddev, size); + event_ref = dpnp_rng_normal_c(q_ref, result, mean, stddev, size, dep_event_vec_ref); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -2076,9 +2101,9 @@ DPCTLSyclEventRef dpnp_rng_triangular_c(DPCTLSyclQueueRef q_ref, kernel_parallel_for_func); }; auto event_ration_acceptance = q.submit(kernel_func); - event_ration_acceptance.wait(); + event_ref = reinterpret_cast(&event_ration_acceptance); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -2143,9 +2168,9 @@ DPCTLSyclEventRef dpnp_rng_uniform_c(DPCTLSyclQueueRef q_ref, mkl_rng::uniform<_DataType> distribution(a, b); // perform generation auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_ref = reinterpret_cast(&event_out); - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -2614,6 +2639,7 @@ DPCTLSyclEventRef dpnp_rng_weibull_c(DPCTLSyclQueueRef q_ref, (void)dep_event_vec_ref; DPCTLSyclEventRef event_ref = nullptr; + sycl::event event_out; if (!size) { @@ -2624,7 +2650,7 @@ DPCTLSyclEventRef dpnp_rng_weibull_c(DPCTLSyclQueueRef q_ref, if (alpha == 0) { - dpnp_zeros_c<_DataType>(result, size); + event_ref = dpnp_zeros_c<_DataType>(q_ref, result, size, dep_event_vec_ref); } else { @@ -2633,10 +2659,10 @@ DPCTLSyclEventRef dpnp_rng_weibull_c(DPCTLSyclQueueRef q_ref, const _DataType beta = (_DataType(1.0)); mkl_rng::weibull<_DataType> distribution(alpha, a, beta); - auto event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); - event_out.wait(); + event_out = mkl_rng::generate(distribution, DPNP_RNG_ENGINE, size, result1); + event_ref = reinterpret_cast(&event_out); } - return event_ref; + return DPCTLEvent_Copy(event_ref); } template @@ -2760,96 +2786,188 @@ void func_map_init_random(func_map_t& fmap) { fmap[DPNPFuncName::DPNP_FN_RNG_BETA][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_beta_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_BETA_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_beta_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_BINOMIAL][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_rng_binomial_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_BINOMIAL_EXT][eft_INT][eft_INT] = {eft_INT, + (void*)dpnp_rng_binomial_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_CHISQUARE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_chisquare_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_CHISQUARE_EXT][eft_DBL][eft_DBL] = {eft_DBL, + (void*)dpnp_rng_chisquare_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_EXPONENTIAL][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_exponential_default_c}; fmap[DPNPFuncName::DPNP_FN_RNG_EXPONENTIAL][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_rng_exponential_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_EXPONENTIAL_EXT][eft_DBL][eft_DBL] = {eft_DBL, + (void*)dpnp_rng_exponential_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_EXPONENTIAL_EXT][eft_FLT][eft_FLT] = {eft_FLT, + (void*)dpnp_rng_exponential_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_F][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_f_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_F_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_f_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_GAMMA][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_gamma_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_GAMMA_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_gamma_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_GAUSSIAN][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_gaussian_default_c}; fmap[DPNPFuncName::DPNP_FN_RNG_GAUSSIAN][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_rng_gaussian_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_GAUSSIAN_EXT][eft_DBL][eft_DBL] = {eft_DBL, + (void*)dpnp_rng_gaussian_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_GAUSSIAN_EXT][eft_FLT][eft_FLT] = {eft_FLT, + (void*)dpnp_rng_gaussian_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_GEOMETRIC][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_rng_geometric_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_GEOMETRIC_EXT][eft_INT][eft_INT] = {eft_INT, + (void*)dpnp_rng_geometric_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_GUMBEL][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_gumbel_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_GUMBEL_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_gumbel_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_HYPERGEOMETRIC][eft_INT][eft_INT] = { eft_INT, (void*)dpnp_rng_hypergeometric_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_HYPERGEOMETRIC_EXT][eft_INT][eft_INT] = { + eft_INT, (void*)dpnp_rng_hypergeometric_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_LAPLACE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_laplace_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_LAPLACE_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_laplace_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_LOGISTIC][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_logistic_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_LOGISTIC_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_logistic_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_LOGNORMAL][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_lognormal_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_LOGNORMAL_EXT][eft_DBL][eft_DBL] = {eft_DBL, + (void*)dpnp_rng_lognormal_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_MULTINOMIAL][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_rng_multinomial_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_MULTINOMIAL_EXT][eft_INT][eft_INT] = {eft_INT, + (void*)dpnp_rng_multinomial_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_MULTIVARIATE_NORMAL][eft_DBL][eft_DBL] = { eft_DBL, (void*)dpnp_rng_multivariate_normal_default_c}; fmap[DPNPFuncName::DPNP_FN_RNG_NEGATIVE_BINOMIAL][eft_INT][eft_INT] = { eft_INT, (void*)dpnp_rng_negative_binomial_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_NEGATIVE_BINOMIAL_EXT][eft_INT][eft_INT] = { + eft_INT, (void*)dpnp_rng_negative_binomial_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_NONCENTRAL_CHISQUARE][eft_DBL][eft_DBL] = { eft_DBL, (void*)dpnp_rng_noncentral_chisquare_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_NONCENTRAL_CHISQUARE_EXT][eft_DBL][eft_DBL] = { + eft_DBL, (void*)dpnp_rng_noncentral_chisquare_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_NORMAL][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_normal_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_NORMAL_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_normal_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_PARETO][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_pareto_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_PARETO_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_pareto_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_POISSON][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_rng_poisson_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_POISSON_EXT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_rng_poisson_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_POWER][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_power_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_POWER_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_power_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_RAYLEIGH][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_rayleigh_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_RAYLEIGH_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_rayleigh_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_shuffle_default_c}; fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_rng_shuffle_default_c}; fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_rng_shuffle_default_c}; fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_rng_shuffle_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_shuffle_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE_EXT][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_rng_shuffle_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE_EXT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_rng_shuffle_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_SHUFFLE_EXT][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_rng_shuffle_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_SRAND][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_srand_c}; fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_CAUCHY][eft_DBL][eft_DBL] = { eft_DBL, (void*)dpnp_rng_standard_cauchy_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_CAUCHY_EXT][eft_DBL][eft_DBL] = { + eft_DBL, (void*)dpnp_rng_standard_cauchy_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_EXPONENTIAL][eft_DBL][eft_DBL] = { eft_DBL, (void*)dpnp_rng_standard_exponential_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_EXPONENTIAL_EXT][eft_DBL][eft_DBL] = { + eft_DBL, (void*)dpnp_rng_standard_exponential_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_GAMMA][eft_DBL][eft_DBL] = { eft_DBL, (void*)dpnp_rng_standard_gamma_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_GAMMA_EXT][eft_DBL][eft_DBL] = { + eft_DBL, (void*)dpnp_rng_standard_gamma_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_NORMAL][eft_DBL][eft_DBL] = { eft_DBL, (void*)dpnp_rng_standard_normal_default_c}; + + fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_NORMAL_EXT][eft_DBL][eft_DBL] = { + eft_DBL, (void*)dpnp_rng_standard_normal_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_T][eft_DBL][eft_DBL] = { eft_DBL, (void*)dpnp_rng_standard_t_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_STANDARD_T_EXT][eft_DBL][eft_DBL] = { + eft_DBL, (void*)dpnp_rng_standard_t_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_TRIANGULAR][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_triangular_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_TRIANGULAR_EXT][eft_DBL][eft_DBL] = {eft_DBL, + (void*)dpnp_rng_triangular_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_uniform_default_c}; fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_rng_uniform_default_c}; fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_rng_uniform_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_uniform_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM_EXT][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_rng_uniform_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_UNIFORM_EXT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_rng_uniform_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_VONMISES][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_vonmises_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_VONMISES_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_vonmises_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_WALD][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_wald_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_WALD_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_wald_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_WEIBULL][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_weibull_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_WEIBULL_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_weibull_ext_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_ZIPF][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_zipf_default_c}; + fmap[DPNPFuncName::DPNP_FN_RNG_ZIPF_EXT][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_rng_zipf_ext_c}; + return; } diff --git a/dpnp/dpnp_algo/dpnp_algo.pxd b/dpnp/dpnp_algo/dpnp_algo.pxd index 1cc650bdaaa8..10679c035afe 100644 --- a/dpnp/dpnp_algo/dpnp_algo.pxd +++ b/dpnp/dpnp_algo/dpnp_algo.pxd @@ -24,6 +24,7 @@ # ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF # THE POSSIBILITY OF SUCH DAMAGE. # ***************************************************************************** +cimport dpctl as c_dpctl from libcpp cimport bool as cpp_bool @@ -134,41 +135,77 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na DPNP_FN_REPEAT DPNP_FN_RIGHT_SHIFT DPNP_FN_RNG_BETA + DPNP_FN_RNG_BETA_EXT DPNP_FN_RNG_BINOMIAL + DPNP_FN_RNG_BINOMIAL_EXT DPNP_FN_RNG_CHISQUARE + DPNP_FN_RNG_CHISQUARE_EXT DPNP_FN_RNG_EXPONENTIAL + DPNP_FN_RNG_EXPONENTIAL_EXT DPNP_FN_RNG_F + DPNP_FN_RNG_F_EXT DPNP_FN_RNG_GAMMA + DPNP_FN_RNG_GAMMA_EXT DPNP_FN_RNG_GAUSSIAN + DPNP_FN_RNG_GAUSSIAN_EXT DPNP_FN_RNG_GEOMETRIC + DPNP_FN_RNG_GEOMETRIC_EXT DPNP_FN_RNG_GUMBEL + DPNP_FN_RNG_GUMBEL_EXT DPNP_FN_RNG_HYPERGEOMETRIC + DPNP_FN_RNG_HYPERGEOMETRIC_EXT DPNP_FN_RNG_LAPLACE + DPNP_FN_RNG_LAPLACE_EXT DPNP_FN_RNG_LOGISTIC + DPNP_FN_RNG_LOGISTIC_EXT DPNP_FN_RNG_LOGNORMAL + DPNP_FN_RNG_LOGNORMAL_EXT DPNP_FN_RNG_MULTINOMIAL + DPNP_FN_RNG_MULTINOMIAL_EXT DPNP_FN_RNG_MULTIVARIATE_NORMAL + DPNP_FN_RNG_MULTIVARIATE_NORMAL_EXT DPNP_FN_RNG_NEGATIVE_BINOMIAL + DPNP_FN_RNG_NEGATIVE_BINOMIAL_EXT DPNP_FN_RNG_NONCENTRAL_CHISQUARE + DPNP_FN_RNG_NONCENTRAL_CHISQUARE_EXT DPNP_FN_RNG_NORMAL + DPNP_FN_RNG_NORMAL_EXT DPNP_FN_RNG_PARETO + DPNP_FN_RNG_PARETO_EXT DPNP_FN_RNG_POISSON + DPNP_FN_RNG_POISSON_EXT DPNP_FN_RNG_POWER + DPNP_FN_RNG_POWER_EXT DPNP_FN_PUT_ALONG_AXIS + DPNP_FN_PUT_ALONG_AXIS_EXT DPNP_FN_RNG_RAYLEIGH + DPNP_FN_RNG_RAYLEIGH_EXT DPNP_FN_RNG_SHUFFLE + DPNP_FN_RNG_SHUFFLE_EXT DPNP_FN_RNG_SRAND + DPNP_FN_RNG_SRAND_EXT DPNP_FN_RNG_STANDARD_CAUCHY + DPNP_FN_RNG_STANDARD_CAUCHY_EXT DPNP_FN_RNG_STANDARD_EXPONENTIAL + DPNP_FN_RNG_STANDARD_EXPONENTIAL_EXT DPNP_FN_RNG_STANDARD_GAMMA + DPNP_FN_RNG_STANDARD_GAMMA_EXT DPNP_FN_RNG_STANDARD_NORMAL + DPNP_FN_RNG_STANDARD_NORMAL_EXT DPNP_FN_RNG_STANDARD_T + DPNP_FN_RNG_STANDARD_T_EXT DPNP_FN_RNG_TRIANGULAR + DPNP_FN_RNG_TRIANGULAR_EXT DPNP_FN_RNG_UNIFORM + DPNP_FN_RNG_UNIFORM_EXT DPNP_FN_RNG_VONMISES + DPNP_FN_RNG_VONMISES_EXT DPNP_FN_RNG_WALD + DPNP_FN_RNG_WALD_EXT DPNP_FN_RNG_WEIBULL + DPNP_FN_RNG_WEIBULL_EXT DPNP_FN_RNG_ZIPF + DPNP_FN_RNG_ZIPF_EXT DPNP_FN_SEARCHSORTED DPNP_FN_SIGN DPNP_FN_SIN diff --git a/dpnp/random/dpnp_algo_random.pyx b/dpnp/random/dpnp_algo_random.pyx index f4376f96a44a..3b01b3d0bf07 100644 --- a/dpnp/random/dpnp_algo_random.pyx +++ b/dpnp/random/dpnp_algo_random.pyx @@ -82,59 +82,193 @@ __all__ = [ ] -ctypedef void(*fptr_dpnp_rng_beta_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_binomial_c_1out_t)(void * , const int, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_chisquare_c_1out_t)(void * , const int, const size_t) except + -ctypedef void(*fptr_dpnp_rng_exponential_c_1out_t)(void * , const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_f_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_gamma_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_geometric_c_1out_t)(void * , const float, const size_t) except + -ctypedef void(*fptr_dpnp_rng_gaussian_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_gumbel_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_hypergeometric_c_1out_t)(void * , const int, const int, const int, const size_t) except + -ctypedef void(*fptr_dpnp_rng_laplace_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_logistic_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_lognormal_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_multinomial_c_1out_t)(void * result, - const int, - const double * , - const size_t, - const size_t) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_beta_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_binomial_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const int, const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_chisquare_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const int, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_exponential_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_f_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_gamma_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_geometric_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const float, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_gaussian_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_gumbel_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_hypergeometric_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const int, + const int, + const int, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_laplace_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_logistic_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_lognormal_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_multinomial_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * result, + const int, + void * , + const size_t, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + ctypedef void(*fptr_dpnp_rng_multivariate_normal_c_1out_t)(void * , const int, - const double * , + void * , const size_t, - const double * , + void * , const size_t, const size_t) except + -ctypedef void(*fptr_dpnp_rng_negative_binomial_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_noncentral_chisquare_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_normal_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_pareto_c_1out_t)(void * , const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_poisson_c_1out_t)(void * , const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_power_c_1out_t)(void * , const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_rayleigh_c_1out_t)(void * , const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_shuffle_c_1out_t)(void * , - const size_t, - const size_t, - const size_t, - const size_t) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_negative_binomial_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_noncentral_chisquare_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_normal_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_pareto_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_poisson_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_power_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_rayleigh_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_shuffle_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, void * , + const size_t, + const size_t, + const size_t, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + ctypedef void(*fptr_dpnp_rng_srand_c_1out_t)(const size_t) except + -ctypedef void(*fptr_dpnp_rng_standard_cauchy_c_1out_t)(void * , const size_t) except + -ctypedef void(*fptr_dpnp_rng_standard_exponential_c_1out_t)(void * , const size_t) except + -ctypedef void(*fptr_dpnp_rng_standard_gamma_c_1out_t)(void * , const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_standard_normal_c_1out_t)(void * , const size_t) except + -ctypedef void(*fptr_dpnp_rng_standard_t_c_1out_t)(void * , const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_triangular_c_1out_t)(void * , - const double, - const double, - const double, - const size_t) except + -ctypedef void(*fptr_dpnp_rng_uniform_c_1out_t)(void * , const long, const long, const size_t) except + -ctypedef void(*fptr_dpnp_rng_vonmises_c_1out_t)(void * , const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_wald_c_1out_t)(void *, const double, const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_weibull_c_1out_t)(void * , const double, const size_t) except + -ctypedef void(*fptr_dpnp_rng_zipf_c_1out_t)(void * , const double, const size_t) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_standard_cauchy_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_standard_exponential_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_standard_gamma_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_standard_normal_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_standard_t_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_triangular_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, void * , + const double, + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_uniform_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const long, + const long, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_vonmises_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_wald_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void *, + const double, + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_weibull_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + +ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_dpnp_rng_zipf_c_1out_t)(c_dpctl.DPCTLSyclQueueRef, + void * , + const double, + const size_t, + const c_dpctl.DPCTLEventVectorRef) except + cpdef utils.dpnp_descriptor dpnp_rng_beta(double a, double b, size): @@ -149,15 +283,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_beta(double a, double b, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_BETA, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_BETA_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_beta_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), a, b, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), a, b, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -180,16 +322,24 @@ cpdef utils.dpnp_descriptor dpnp_rng_binomial(int ntrial, double p, size): param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_BINOMIAL, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_BINOMIAL_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), ntrial, p, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), ntrial, p, result.size, NULL) + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) + return result @@ -205,15 +355,24 @@ cpdef utils.dpnp_descriptor dpnp_rng_chisquare(int df, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_CHISQUARE, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_CHISQUARE_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + + cdef fptr_dpnp_rng_chisquare_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), df, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), df, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -231,15 +390,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_exponential(double beta, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_EXPONENTIAL, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_EXPONENTIAL_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_exponential_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), beta, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), beta, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -256,15 +423,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_f(double df_num, double df_den, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_F, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_F_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_f_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), df_num, df_den, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), df_num, df_den, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -284,15 +459,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_gamma(double shape, double scale, size): param1_type = dpnp_dtype_to_DPNPFuncType(dtype) - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_GAMMA, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_GAMMA_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), shape, scale, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), shape, scale, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -315,15 +498,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_geometric(float p, size): param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_GEOMETRIC, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_GEOMETRIC_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), p, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), p, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -342,15 +533,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_gumbel(double loc, double scale, size): cdef fptr_dpnp_rng_gumbel_c_1out_t func param1_type = dpnp_dtype_to_DPNPFuncType(dtype) - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_GUMBEL, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_GUMBEL_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), loc, scale, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), loc, scale, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -372,15 +571,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_hypergeometric(int l, int s, int m, size): param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_HYPERGEOMETRIC, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_HYPERGEOMETRIC_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), l, s, m, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), l, s, m, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -402,15 +609,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_laplace(double loc, double scale, size): param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_LAPLACE, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_LAPLACE_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), loc, scale, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), loc, scale, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -427,15 +642,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_logistic(double loc, double scale, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_LOGISTIC, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_LOGISTIC_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_logistic_c_1out_t func = < fptr_dpnp_rng_logistic_c_1out_t > kernel_data.ptr # call FPTR function - func(result.get_data(), loc, scale, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), loc, scale, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -457,20 +680,28 @@ cpdef utils.dpnp_descriptor dpnp_rng_lognormal(double mean, double stddev, size) param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_LOGNORMAL, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_LOGNORMAL_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), mean, stddev, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), mean, stddev, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result -cpdef utils.dpnp_descriptor dpnp_rng_multinomial(int ntrial, p, size): +cpdef utils.dpnp_descriptor dpnp_rng_multinomial(int ntrial, utils.dpnp_descriptor p, size): """ Returns an array populated with samples from multinomial distribution. @@ -484,29 +715,43 @@ cpdef utils.dpnp_descriptor dpnp_rng_multinomial(int ntrial, p, size): cdef DPNPFuncType param1_type cdef DPNPFuncData kernel_data cdef fptr_dpnp_rng_multinomial_c_1out_t func - p = numpy.asarray(p, dtype=numpy.float64) - cdef double * p_vector = numpy.PyArray_DATA(p) - cdef size_t p_vector_size = len(p) + cdef size_t p_size = p.size # convert string type names (array.dtype) to C enum DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_MULTINOMIAL, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_MULTINOMIAL_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) - cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + + p_obj = p.get_array() + + cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, + kernel_data.return_type, + None, + device=p_obj.device, + usm_type=p_obj.usm_type, + sycl_queue=p_obj.sycl_queue) + + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() func = kernel_data.ptr # call FPTR function - func(result.get_data(), ntrial, p_vector, p_vector_size, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), ntrial, p.get_data(), p_size, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result -cpdef utils.dpnp_descriptor dpnp_rng_multivariate_normal(numpy.ndarray mean, numpy.ndarray cov, size): +cpdef utils.dpnp_descriptor dpnp_rng_multivariate_normal(utils.dpnp_descriptor mean, utils.dpnp_descriptor cov, size): """ Returns an array populated with samples from multivariate normal distribution. `dpnp_rng_multivariate_normal` generates a matrix filled with random floats sampled from a @@ -516,21 +761,15 @@ cpdef utils.dpnp_descriptor dpnp_rng_multivariate_normal(numpy.ndarray mean, num dtype = numpy.float64 cdef int dimen - cdef double * mean_vector - cdef double * cov_vector - cdef size_t mean_vector_size - cdef size_t cov_vector_size + cdef size_t mean_size + cdef size_t cov_size cdef DPNPFuncType param1_type cdef DPNPFuncData kernel_data cdef fptr_dpnp_rng_multivariate_normal_c_1out_t func - # mean and cov expected numpy.ndarray in C order (row major) - mean_vector = numpy.PyArray_DATA(mean) - cov_vector = numpy.PyArray_DATA(cov) - - mean_vector_size = mean.size - cov_vector_size = cov.size + mean_size = mean.size + cov_size = cov.size # convert string type names (array.dtype) to C enum DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype) @@ -540,17 +779,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_multivariate_normal(numpy.ndarray mean, num # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) - cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_device, result_usm_type, result_sycl_queue = utils.get_common_usm_allocation(mean, cov) + cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, + kernel_data.return_type, + None, + device=result_sycl_device, + usm_type=result_usm_type, + sycl_queue=result_sycl_queue) - dimen = len(mean) + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() func = kernel_data.ptr # call FPTR function - func(result.get_data(), dimen, mean_vector, mean_vector_size, cov_vector, cov_vector_size, result.size) + func(result.get_data(), mean_size, mean.get_data(), mean_size, cov.get_data(), cov_size, result.size) return result - cpdef utils.dpnp_descriptor dpnp_rng_negative_binomial(double a, double p, size): """ Returns an array populated with samples from negative binomial distribution. @@ -567,6 +812,9 @@ cpdef utils.dpnp_descriptor dpnp_rng_negative_binomial(double a, double p, size) cdef DPNPFuncData kernel_data cdef fptr_dpnp_rng_negative_binomial_c_1out_t func cdef shape_type_c result_shape + cdef c_dpctl.SyclQueue q + cdef c_dpctl.DPCTLSyclQueueRef q_ref + cdef c_dpctl.DPCTLSyclEventRef event_ref if p == 0.0: filled_val = numpy.iinfo(dtype).min @@ -578,15 +826,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_negative_binomial(double a, double p, size) param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_NEGATIVE_BINOMIAL, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_NEGATIVE_BINOMIAL_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data result_shape = utils._object_to_tuple(size) result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + q = result_sycl_queue + q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), a, p, result.size) + event_ref = func(q_ref, result.get_data(), a, p, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -603,15 +859,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_noncentral_chisquare(double df, double nonc cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_NONCENTRAL_CHISQUARE, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_NONCENTRAL_CHISQUARE_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_noncentral_chisquare_c_1out_t func = < fptr_dpnp_rng_noncentral_chisquare_c_1out_t > kernel_data.ptr # call FPTR function - func(result.get_data(), df, nonc, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), df, nonc, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -630,6 +894,9 @@ cpdef utils.dpnp_descriptor dpnp_rng_normal(double loc, double scale, size): cdef DPNPFuncType param1_type cdef DPNPFuncData kernel_data cdef fptr_dpnp_rng_normal_c_1out_t func + cdef c_dpctl.SyclQueue q + cdef c_dpctl.DPCTLSyclQueueRef q_ref + cdef c_dpctl.DPCTLSyclEventRef event_ref if scale == 0.0: return dpnp_full(size, loc, dtype) @@ -638,15 +905,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_normal(double loc, double scale, size): param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_NORMAL, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_NORMAL_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data result_shape = utils._object_to_tuple(size) result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + q = result_sycl_queue + q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), loc, scale, result.size) + event_ref = func(q_ref, result.get_data(), loc, scale, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -664,15 +939,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_pareto(double alpha, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_PARETO, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_PARETO_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_pareto_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), alpha, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), alpha, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -692,6 +975,9 @@ cpdef utils.dpnp_descriptor dpnp_rng_poisson(double lam, size): cdef DPNPFuncType param1_type cdef DPNPFuncData kernel_data cdef fptr_dpnp_rng_poisson_c_1out_t func + cdef c_dpctl.SyclQueue q + cdef c_dpctl.DPCTLSyclQueueRef q_ref + cdef c_dpctl.DPCTLSyclEventRef event_ref if lam == 0: return dpnp_full(size, 0, dtype) @@ -700,15 +986,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_poisson(double lam, size): param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_POISSON, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_POISSON_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data result_shape = utils._object_to_tuple(size) result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + q = result_sycl_queue + q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), lam, result.size) + event_ref = func(q_ref, result.get_data(), lam, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -725,15 +1019,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_power(double alpha, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_POWER, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_POWER_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_power_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), alpha, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), alpha, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -752,15 +1054,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_randn(dims): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_GAUSSIAN, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_GAUSSIAN_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(dims) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_gaussian_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), mean, stddev, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), mean, stddev, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -778,15 +1088,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_random(dims): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_UNIFORM, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_UNIFORM_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(dims) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_uniform_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), low, high, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), low, high, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -805,6 +1123,9 @@ cpdef utils.dpnp_descriptor dpnp_rng_rayleigh(double scale, size): cdef DPNPFuncType param1_type cdef DPNPFuncData kernel_data cdef fptr_dpnp_rng_rayleigh_c_1out_t func + cdef c_dpctl.SyclQueue q + cdef c_dpctl.DPCTLSyclQueueRef q_ref + cdef c_dpctl.DPCTLSyclEventRef event_ref if scale == 0.0: return dpnp_full(size, 0.0, dtype) @@ -813,15 +1134,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_rayleigh(double scale, size): param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_RAYLEIGH, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_RAYLEIGH_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data result_shape = utils._object_to_tuple(size) result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + q = result_sycl_queue + q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), scale, result.size) + event_ref = func(q_ref, result.get_data(), scale, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -839,11 +1168,19 @@ cpdef utils.dpnp_descriptor dpnp_rng_shuffle(utils.dpnp_descriptor x1): cdef size_t high_dim_size = x1.get_pyobj().size # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_SHUFFLE, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_SHUFFLE_EXT, param1_type, param1_type) + + x1_sycl_queue = x1.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = x1_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() cdef fptr_dpnp_rng_shuffle_c_1out_t func = < fptr_dpnp_rng_shuffle_c_1out_t > kernel_data.ptr # call FPTR function - func(x1.get_data(), itemsize, ndim, high_dim_size, x1.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, x1.get_data(), itemsize, ndim, high_dim_size, x1.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return x1 @@ -877,15 +1214,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_standard_cauchy(size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_STANDARD_CAUCHY, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_STANDARD_CAUCHY_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_standard_cauchy_c_1out_t func = < fptr_dpnp_rng_standard_cauchy_c_1out_t > kernel_data.ptr # call FPTR function - func(result.get_data(), result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -904,15 +1249,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_standard_exponential(size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_STANDARD_EXPONENTIAL, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_STANDARD_EXPONENTIAL_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + func = < fptr_dpnp_rng_standard_exponential_c_1out_t > kernel_data.ptr # call FPTR function - func(result.get_data(), result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -931,6 +1284,9 @@ cpdef utils.dpnp_descriptor dpnp_rng_standard_gamma(double shape, size): cdef DPNPFuncType param1_type cdef DPNPFuncData kernel_data cdef fptr_dpnp_rng_standard_gamma_c_1out_t func + cdef c_dpctl.SyclQueue q + cdef c_dpctl.DPCTLSyclQueueRef q_ref + cdef c_dpctl.DPCTLSyclEventRef event_ref if shape == 0.0: return dpnp_full(size, 0.0, dtype) @@ -939,15 +1295,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_standard_gamma(double shape, size): param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_STANDARD_GAMMA, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_STANDARD_GAMMA_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data result_shape = utils._object_to_tuple(size) result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + q = result_sycl_queue + q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), shape, result.size) + event_ref = func(q_ref, result.get_data(), shape, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -964,15 +1328,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_standard_normal(size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_STANDARD_NORMAL, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_STANDARD_NORMAL_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_standard_normal_c_1out_t func = < fptr_dpnp_rng_standard_normal_c_1out_t > kernel_data.ptr # call FPTR function - func(result.get_data(), result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -988,15 +1360,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_standard_t(double df, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_STANDARD_T, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_STANDARD_T_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_standard_t_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), df, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), df, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -1014,15 +1394,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_triangular(double left, double mode, double cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_TRIANGULAR, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_TRIANGULAR_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_triangular_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), left, mode, right, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), left, mode, right, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -1041,6 +1429,9 @@ cpdef utils.dpnp_descriptor dpnp_rng_uniform(long low, long high, size, dtype): cdef DPNPFuncType param1_type cdef DPNPFuncData kernel_data cdef fptr_dpnp_rng_uniform_c_1out_t func + cdef c_dpctl.SyclQueue q + cdef c_dpctl.DPCTLSyclQueueRef q_ref + cdef c_dpctl.DPCTLSyclEventRef event_ref if low == high: return dpnp_full(size, low, dtype) @@ -1049,15 +1440,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_uniform(long low, long high, size, dtype): param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_UNIFORM, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_UNIFORM_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data result_shape = utils._object_to_tuple(size) result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + q = result_sycl_queue + q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), low, high, result.size) + event_ref = func(q_ref, result.get_data(), low, high, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -1074,15 +1473,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_vonmises(double mu, double kappa, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_VONMISES, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_VONMISES_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_vonmises_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), mu, kappa, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), mu, kappa, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -1100,15 +1507,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_wald(double mean, double scale, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_WALD, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_WALD_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_wald_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), mean, scale, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), mean, scale, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -1130,15 +1545,23 @@ cpdef utils.dpnp_descriptor dpnp_rng_weibull(double a, size): param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_WEIBULL, param1_type, param1_type) + kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_WEIBULL_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + func = kernel_data.ptr # call FPTR function - func(result.get_data(), a, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), a, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result @@ -1155,14 +1578,22 @@ cpdef utils.dpnp_descriptor dpnp_rng_zipf(double a, size): cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(numpy.float64) # get the FPTR data structure - cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_ZIPF, param1_type, param1_type) + cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_RNG_ZIPF_EXT, param1_type, param1_type) # ceate result array with type given by FPTR data cdef shape_type_c result_shape = utils._object_to_tuple(size) cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None) + result_sycl_queue = result.get_array().sycl_queue + + cdef c_dpctl.SyclQueue q = result_sycl_queue + cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref() + cdef fptr_dpnp_rng_zipf_c_1out_t func = kernel_data.ptr # call FPTR function - func(result.get_data(), a, result.size) + cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, result.get_data(), a, result.size, NULL) + + with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref) + c_dpctl.DPCTLEvent_Delete(event_ref) return result diff --git a/dpnp/random/dpnp_iface_random.py b/dpnp/random/dpnp_iface_random.py index 2b7b2a70e6c3..533ba6949955 100644 --- a/dpnp/random/dpnp_iface_random.py +++ b/dpnp/random/dpnp_iface_random.py @@ -656,6 +656,7 @@ def multinomial(n, pvals, size=None): if not use_origin_backend(n): pvals_sum = sum(pvals) + pvals_desc = dpnp.get_dpnp_descriptor(dpnp.array(pvals)) d = len(pvals) if n < 0: pass @@ -674,7 +675,7 @@ def multinomial(n, pvals, size=None): except: shape = tuple(size) + (d,) - return dpnp_rng_multinomial(int(n), pvals, shape).get_pyobj() + return dpnp_rng_multinomial(int(n), pvals_desc, shape).get_pyobj() return call_origin(numpy.random.multinomial, n, pvals, size) @@ -703,8 +704,8 @@ def multivariate_normal(mean, cov, size=None, check_valid='warn', tol=1e-8): """ if not use_origin_backend(mean): - mean_ = numpy.array(mean, dtype=numpy.float64, order='C') - cov_ = numpy.array(cov, dtype=numpy.float64, order='C') + mean_ = dpnp.get_dpnp_descriptor(dpnp.array(mean, dtype=numpy.float64)) + cov_ = dpnp.get_dpnp_descriptor(dpnp.array(cov, dtype=numpy.float64)) if size is None: shape = [] elif isinstance(size, (int, numpy.integer)):