From 31ccbdd67cebeee1e44f65442d51165711b2fa70 Mon Sep 17 00:00:00 2001 From: Denis Demidov Date: Fri, 20 Sep 2013 22:58:52 +0400 Subject: [PATCH] Split local and global scope constants Made VEX_CONSTANT return a functor. This makes it suitable for definition in local scope. Predefined constants are defined (with VEX_GLOBAL_CONSTANT)as inline functions which makes them safe for inclusion into several object modules. --- README.md | 46 ++++++++++-- vexcl/constants.hpp | 165 ++++++++++++++++++++++++-------------------- 2 files changed, 132 insertions(+), 79 deletions(-) diff --git a/README.md b/README.md index 9b1a4a4cb..8ff48ffc4 100644 --- a/README.md +++ b/README.md @@ -32,6 +32,7 @@ performance of several GPGPU libraries, including VexCL. * [Copies between host and devices](#copies-between-host-and-devices) * [Vector expressions](#vector-expressions) * [Builtin operations](#builtin-operations) + * [Constants](#constants) * [Element indices](#element-indices) * [User-defined functions](#user-defined-functions) * [Tagged terminals](#tagged-terminals) @@ -176,16 +177,16 @@ X = 2 * Y - sin(Z); ~~~ will lead to the launch of the following OpenCL kernel: ~~~{.c} -kernel void minus_multiplies_term_term_sin_term_( +kernel void vexcl_vector_kernel( ulong n, - global double *res, - int prm_1, - global double *prm_2, - global double *prm_3 + global double * prm_1, + int prm_2, + global double * prm_3, + global double * prm_4 ) { for(size_t idx = get_global_id(0); idx < n; idx += get_global_size(0)) { - res[idx] = ( ( prm_1 * prm_2[idx] ) - sin( prm_3[idx] ) ); + prm_1[idx] = ( ( prm_2 * prm_3[idx] ) - sin( prm_4[idx] ) ); } } ~~~ @@ -203,6 +204,39 @@ Please do not hesitate to open an issue in this case. Z = sqrt(2 * X) + pow(cos(Y), 2.0); ~~~ +### Constants + +As you have seen above, `2` in the expression `2 * Y - sin(Z)` is passed to the +generated OpenCL kernel as an `int` parameter (`prm_1`). Sometimes this is +desired behaviour, because same kernel will be reused for expressions `42 * +Z - sin(Y)` or `a * Y - sin(Y)` (where `a` is an integer variable). But this +may lead to a slight overhead if an expression involves true constant that will +always have same value. Macro `VEX_CONSTANT` allows to define such constants +for use in vector expressions. Compare the generated kernel for the following +example with the kernel above: +~~~{.cpp} +VEX_CONSTANT(two, 2); + +X = two() * Y - sin(Z); +~~~ + +~~~{.c} +kernel void vexcl_vector_kernel( + ulong n, + global double * prm_1, + global double * prm_3, + global double * prm_4 +) +{ + for(size_t idx = get_global_id(0); idx < n; idx += get_global_size(0)) { + prm_1[idx] = ( ( ( 2 ) * prm_3[idx] ) - sin( prm_4[idx] ) ); + } +} +~~~ + +VexCL provides some predefined constants in `vex::constants` namespace that +correspond to boost::math::constants (e.g. `vex::constants::pi()`). + ### Element indices Function `vex::element_index(size_t offset = 0)` allows to use an index of each diff --git a/vexcl/constants.hpp b/vexcl/constants.hpp index 12e7efb90..6c3530e6a 100644 --- a/vexcl/constants.hpp +++ b/vexcl/constants.hpp @@ -133,12 +133,28 @@ struct kernel_arg_setter< user_constant > /// \endcond -/// Mathematical constants imported from Boost. -namespace constants { } - -/// Create constant for use in VexCL expressions +/// Create user-defined constant for use in VexCL expressions #define VEX_CONSTANT(name, value) \ - namespace constants { \ + struct constant_##name { \ + static std::string get() { \ + std::ostringstream s; \ + s << "( " << std::scientific << std::setprecision(16) << value << " )"; \ + return s.str(); \ + } \ + typename boost::proto::result_of::as_expr< \ + vex::user_constant, \ + vex::vector_domain>::type \ + operator()() const { \ + return boost::proto::as_expr( \ + vex::user_constant()); \ + } \ + } name + +/// Mathematical constants. +namespace constants { + +// Can not implement these as functors, could end up in multiple object files. +#define VEX_GLOBAL_CONSTANT(name, value) \ struct constant_##name { \ static std::string get() { \ std::ostringstream s; \ @@ -147,82 +163,85 @@ namespace constants { } } \ }; \ inline typename boost::proto::result_of::as_expr< \ - user_constant, vector_domain>::type name() { \ + user_constant, vector_domain>::type \ + name() { \ return boost::proto::as_expr( \ user_constant()); \ - } \ } -VEX_CONSTANT( pi, boost::math::constants::pi() ) -VEX_CONSTANT( root_pi, boost::math::constants::root_pi() ) -VEX_CONSTANT( root_half_pi, boost::math::constants::root_half_pi() ) -VEX_CONSTANT( root_two_pi, boost::math::constants::root_two_pi() ) -VEX_CONSTANT( root_ln_four, boost::math::constants::root_ln_four() ) -VEX_CONSTANT( e, boost::math::constants::e() ) -VEX_CONSTANT( half, boost::math::constants::half() ) -VEX_CONSTANT( euler, boost::math::constants::euler() ) -VEX_CONSTANT( root_two, boost::math::constants::root_two() ) -VEX_CONSTANT( ln_two, boost::math::constants::ln_two() ) -VEX_CONSTANT( ln_ln_two, boost::math::constants::ln_ln_two() ) -VEX_CONSTANT( third, boost::math::constants::third() ) -VEX_CONSTANT( twothirds, boost::math::constants::twothirds() ) -VEX_CONSTANT( pi_minus_three, boost::math::constants::pi_minus_three() ) -VEX_CONSTANT( four_minus_pi, boost::math::constants::four_minus_pi() ) -VEX_CONSTANT( two_pi, boost::math::constants::two_pi() ) -VEX_CONSTANT( one_div_two_pi, boost::math::constants::one_div_two_pi() ) -VEX_CONSTANT( half_root_two, boost::math::constants::half_root_two() ) -VEX_CONSTANT( pow23_four_minus_pi, boost::math::constants::pow23_four_minus_pi() ) -VEX_CONSTANT( exp_minus_half, boost::math::constants::exp_minus_half() ) +VEX_GLOBAL_CONSTANT( pi, boost::math::constants::pi() ); +VEX_GLOBAL_CONSTANT( root_pi, boost::math::constants::root_pi() ); +VEX_GLOBAL_CONSTANT( root_half_pi, boost::math::constants::root_half_pi() ); +VEX_GLOBAL_CONSTANT( root_two_pi, boost::math::constants::root_two_pi() ); +VEX_GLOBAL_CONSTANT( root_ln_four, boost::math::constants::root_ln_four() ); +VEX_GLOBAL_CONSTANT( e, boost::math::constants::e() ); +VEX_GLOBAL_CONSTANT( half, boost::math::constants::half() ); +VEX_GLOBAL_CONSTANT( euler, boost::math::constants::euler() ); +VEX_GLOBAL_CONSTANT( root_two, boost::math::constants::root_two() ); +VEX_GLOBAL_CONSTANT( ln_two, boost::math::constants::ln_two() ); +VEX_GLOBAL_CONSTANT( ln_ln_two, boost::math::constants::ln_ln_two() ); +VEX_GLOBAL_CONSTANT( third, boost::math::constants::third() ); +VEX_GLOBAL_CONSTANT( twothirds, boost::math::constants::twothirds() ); +VEX_GLOBAL_CONSTANT( pi_minus_three, boost::math::constants::pi_minus_three() ); +VEX_GLOBAL_CONSTANT( four_minus_pi, boost::math::constants::four_minus_pi() ); +VEX_GLOBAL_CONSTANT( two_pi, boost::math::constants::two_pi() ); +VEX_GLOBAL_CONSTANT( one_div_two_pi, boost::math::constants::one_div_two_pi() ); +VEX_GLOBAL_CONSTANT( half_root_two, boost::math::constants::half_root_two() ); +VEX_GLOBAL_CONSTANT( pow23_four_minus_pi, boost::math::constants::pow23_four_minus_pi() ); +VEX_GLOBAL_CONSTANT( exp_minus_half, boost::math::constants::exp_minus_half() ); #if (BOOST_VERSION >= 105000) || defined(DOXYGEN) -VEX_CONSTANT( catalan, boost::math::constants::catalan() ) -VEX_CONSTANT( cbrt_pi, boost::math::constants::cbrt_pi() ) -VEX_CONSTANT( cosh_one, boost::math::constants::cosh_one() ) -VEX_CONSTANT( cos_one, boost::math::constants::cos_one() ) -VEX_CONSTANT( degree, boost::math::constants::degree() ) -VEX_CONSTANT( e_pow_pi, boost::math::constants::e_pow_pi() ) -VEX_CONSTANT( euler_sqr, boost::math::constants::euler_sqr() ) -VEX_CONSTANT( extreme_value_skewness, boost::math::constants::extreme_value_skewness() ) -VEX_CONSTANT( four_thirds_pi, boost::math::constants::four_thirds_pi() ) -VEX_CONSTANT( glaisher, boost::math::constants::glaisher() ) -VEX_CONSTANT( half_pi, boost::math::constants::half_pi() ) -VEX_CONSTANT( khinchin, boost::math::constants::khinchin() ) -VEX_CONSTANT( ln_phi, boost::math::constants::ln_phi() ) -VEX_CONSTANT( ln_ten, boost::math::constants::ln_ten() ) -VEX_CONSTANT( log10_e, boost::math::constants::log10_e() ) -VEX_CONSTANT( one_div_cbrt_pi, boost::math::constants::one_div_cbrt_pi() ) -VEX_CONSTANT( one_div_euler, boost::math::constants::one_div_euler() ) -VEX_CONSTANT( one_div_ln_phi, boost::math::constants::one_div_ln_phi() ) -VEX_CONSTANT( one_div_log10_e, boost::math::constants::one_div_log10_e() ) -VEX_CONSTANT( one_div_root_pi, boost::math::constants::one_div_root_pi() ) -VEX_CONSTANT( one_div_root_two, boost::math::constants::one_div_root_two() ) -VEX_CONSTANT( one_div_root_two_pi, boost::math::constants::one_div_root_two_pi() ) -VEX_CONSTANT( phi, boost::math::constants::phi() ) -VEX_CONSTANT( pi_cubed, boost::math::constants::pi_cubed() ) -VEX_CONSTANT( pi_pow_e, boost::math::constants::pi_pow_e() ) -VEX_CONSTANT( pi_sqr, boost::math::constants::pi_sqr() ) -VEX_CONSTANT( pi_sqr_div_six, boost::math::constants::pi_sqr_div_six() ) -VEX_CONSTANT( radian, boost::math::constants::radian() ) -VEX_CONSTANT( rayleigh_kurtosis, boost::math::constants::rayleigh_kurtosis() ) -VEX_CONSTANT( rayleigh_kurtosis_excess, boost::math::constants::rayleigh_kurtosis_excess() ) -VEX_CONSTANT( rayleigh_skewness, boost::math::constants::rayleigh_skewness() ) -VEX_CONSTANT( root_e, boost::math::constants::root_e() ) -VEX_CONSTANT( root_one_div_pi, boost::math::constants::root_one_div_pi() ) -VEX_CONSTANT( root_three, boost::math::constants::root_three() ) -VEX_CONSTANT( root_two_div_pi, boost::math::constants::root_two_div_pi() ) -VEX_CONSTANT( sinh_one, boost::math::constants::sinh_one() ) -VEX_CONSTANT( sin_one, boost::math::constants::sin_one() ) -VEX_CONSTANT( sixth_pi, boost::math::constants::sixth_pi() ) -VEX_CONSTANT( third_pi, boost::math::constants::third_pi() ) -VEX_CONSTANT( three_quarters, boost::math::constants::three_quarters() ) -VEX_CONSTANT( three_quarters_pi, boost::math::constants::three_quarters_pi() ) -VEX_CONSTANT( two_div_pi, boost::math::constants::two_div_pi() ) -VEX_CONSTANT( two_thirds, boost::math::constants::two_thirds() ) -VEX_CONSTANT( two_thirds_pi, boost::math::constants::two_thirds_pi() ) -VEX_CONSTANT( zeta_three, boost::math::constants::zeta_three() ) -VEX_CONSTANT( zeta_two, boost::math::constants::zeta_two() ) +VEX_GLOBAL_CONSTANT( catalan, boost::math::constants::catalan() ); +VEX_GLOBAL_CONSTANT( cbrt_pi, boost::math::constants::cbrt_pi() ); +VEX_GLOBAL_CONSTANT( cosh_one, boost::math::constants::cosh_one() ); +VEX_GLOBAL_CONSTANT( cos_one, boost::math::constants::cos_one() ); +VEX_GLOBAL_CONSTANT( degree, boost::math::constants::degree() ); +VEX_GLOBAL_CONSTANT( e_pow_pi, boost::math::constants::e_pow_pi() ); +VEX_GLOBAL_CONSTANT( euler_sqr, boost::math::constants::euler_sqr() ); +VEX_GLOBAL_CONSTANT( extreme_value_skewness, boost::math::constants::extreme_value_skewness() ); +VEX_GLOBAL_CONSTANT( four_thirds_pi, boost::math::constants::four_thirds_pi() ); +VEX_GLOBAL_CONSTANT( glaisher, boost::math::constants::glaisher() ); +VEX_GLOBAL_CONSTANT( half_pi, boost::math::constants::half_pi() ); +VEX_GLOBAL_CONSTANT( khinchin, boost::math::constants::khinchin() ); +VEX_GLOBAL_CONSTANT( ln_phi, boost::math::constants::ln_phi() ); +VEX_GLOBAL_CONSTANT( ln_ten, boost::math::constants::ln_ten() ); +VEX_GLOBAL_CONSTANT( log10_e, boost::math::constants::log10_e() ); +VEX_GLOBAL_CONSTANT( one_div_cbrt_pi, boost::math::constants::one_div_cbrt_pi() ); +VEX_GLOBAL_CONSTANT( one_div_euler, boost::math::constants::one_div_euler() ); +VEX_GLOBAL_CONSTANT( one_div_ln_phi, boost::math::constants::one_div_ln_phi() ); +VEX_GLOBAL_CONSTANT( one_div_log10_e, boost::math::constants::one_div_log10_e() ); +VEX_GLOBAL_CONSTANT( one_div_root_pi, boost::math::constants::one_div_root_pi() ); +VEX_GLOBAL_CONSTANT( one_div_root_two, boost::math::constants::one_div_root_two() ); +VEX_GLOBAL_CONSTANT( one_div_root_two_pi, boost::math::constants::one_div_root_two_pi() ); +VEX_GLOBAL_CONSTANT( phi, boost::math::constants::phi() ); +VEX_GLOBAL_CONSTANT( pi_cubed, boost::math::constants::pi_cubed() ); +VEX_GLOBAL_CONSTANT( pi_pow_e, boost::math::constants::pi_pow_e() ); +VEX_GLOBAL_CONSTANT( pi_sqr, boost::math::constants::pi_sqr() ); +VEX_GLOBAL_CONSTANT( pi_sqr_div_six, boost::math::constants::pi_sqr_div_six() ); +VEX_GLOBAL_CONSTANT( radian, boost::math::constants::radian() ); +VEX_GLOBAL_CONSTANT( rayleigh_kurtosis, boost::math::constants::rayleigh_kurtosis() ); +VEX_GLOBAL_CONSTANT( rayleigh_kurtosis_excess, boost::math::constants::rayleigh_kurtosis_excess() ); +VEX_GLOBAL_CONSTANT( rayleigh_skewness, boost::math::constants::rayleigh_skewness() ); +VEX_GLOBAL_CONSTANT( root_e, boost::math::constants::root_e() ); +VEX_GLOBAL_CONSTANT( root_one_div_pi, boost::math::constants::root_one_div_pi() ); +VEX_GLOBAL_CONSTANT( root_three, boost::math::constants::root_three() ); +VEX_GLOBAL_CONSTANT( root_two_div_pi, boost::math::constants::root_two_div_pi() ); +VEX_GLOBAL_CONSTANT( sinh_one, boost::math::constants::sinh_one() ); +VEX_GLOBAL_CONSTANT( sin_one, boost::math::constants::sin_one() ); +VEX_GLOBAL_CONSTANT( sixth_pi, boost::math::constants::sixth_pi() ); +VEX_GLOBAL_CONSTANT( third_pi, boost::math::constants::third_pi() ); +VEX_GLOBAL_CONSTANT( three_quarters, boost::math::constants::three_quarters() ); +VEX_GLOBAL_CONSTANT( three_quarters_pi, boost::math::constants::three_quarters_pi() ); +VEX_GLOBAL_CONSTANT( two_div_pi, boost::math::constants::two_div_pi() ); +VEX_GLOBAL_CONSTANT( two_thirds, boost::math::constants::two_thirds() ); +VEX_GLOBAL_CONSTANT( two_thirds_pi, boost::math::constants::two_thirds_pi() ); +VEX_GLOBAL_CONSTANT( zeta_three, boost::math::constants::zeta_three() ); +VEX_GLOBAL_CONSTANT( zeta_two, boost::math::constants::zeta_two() ); #endif + +} //namespace constants + } // namespace vex #endif