Skip to content

Commit

Permalink
Split local and global scope constants
Browse files Browse the repository at this point in the history 
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.
  • Loading branch information
@ddemidov
ddemidov committed Sep 20, 2013
1 parent 69af029 commit 31ccbdd
Show file tree
Hide file tree
Showing 2 changed files with 132 additions and 79 deletions.
46 changes: 40 additions & 6 deletions README.md
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down Expand Up @@ -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] ) );
}
}
~~~
Expand All @@ -203,6 +204,39 @@ Please do not hesitate to open an issue in this case.
Z = sqrt(2 * X) + pow(cos(Y), 2.0);
~~~

### <a name="constants"></a>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()`).

### <a name="element-indices"></a>Element indices

Function `vex::element_index(size_t offset = 0)` allows to use an index of each
Expand Down
165 changes: 92 additions & 73 deletions vexcl/constants.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -133,12 +133,28 @@ struct kernel_arg_setter< user_constant<Impl> >

/// \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<constant_##name>, \
vex::vector_domain>::type \
operator()() const { \
return boost::proto::as_expr<vex::vector_domain>( \
vex::user_constant<constant_##name>()); \
} \
} 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; \
Expand All @@ -147,82 +163,85 @@ namespace constants { }
} \
}; \
inline typename boost::proto::result_of::as_expr< \
user_constant<constant_##name>, vector_domain>::type name() { \
user_constant<constant_##name>, vector_domain>::type \
name() { \
return boost::proto::as_expr<vector_domain>( \
user_constant<constant_##name>()); \
} \
}

VEX_CONSTANT( pi, boost::math::constants::pi<double>() )
VEX_CONSTANT( root_pi, boost::math::constants::root_pi<double>() )
VEX_CONSTANT( root_half_pi, boost::math::constants::root_half_pi<double>() )
VEX_CONSTANT( root_two_pi, boost::math::constants::root_two_pi<double>() )
VEX_CONSTANT( root_ln_four, boost::math::constants::root_ln_four<double>() )
VEX_CONSTANT( e, boost::math::constants::e<double>() )
VEX_CONSTANT( half, boost::math::constants::half<double>() )
VEX_CONSTANT( euler, boost::math::constants::euler<double>() )
VEX_CONSTANT( root_two, boost::math::constants::root_two<double>() )
VEX_CONSTANT( ln_two, boost::math::constants::ln_two<double>() )
VEX_CONSTANT( ln_ln_two, boost::math::constants::ln_ln_two<double>() )
VEX_CONSTANT( third, boost::math::constants::third<double>() )
VEX_CONSTANT( twothirds, boost::math::constants::twothirds<double>() )
VEX_CONSTANT( pi_minus_three, boost::math::constants::pi_minus_three<double>() )
VEX_CONSTANT( four_minus_pi, boost::math::constants::four_minus_pi<double>() )
VEX_CONSTANT( two_pi, boost::math::constants::two_pi<double>() )
VEX_CONSTANT( one_div_two_pi, boost::math::constants::one_div_two_pi<double>() )
VEX_CONSTANT( half_root_two, boost::math::constants::half_root_two<double>() )
VEX_CONSTANT( pow23_four_minus_pi, boost::math::constants::pow23_four_minus_pi<double>() )
VEX_CONSTANT( exp_minus_half, boost::math::constants::exp_minus_half<double>() )
VEX_GLOBAL_CONSTANT( pi, boost::math::constants::pi<double>() );
VEX_GLOBAL_CONSTANT( root_pi, boost::math::constants::root_pi<double>() );
VEX_GLOBAL_CONSTANT( root_half_pi, boost::math::constants::root_half_pi<double>() );
VEX_GLOBAL_CONSTANT( root_two_pi, boost::math::constants::root_two_pi<double>() );
VEX_GLOBAL_CONSTANT( root_ln_four, boost::math::constants::root_ln_four<double>() );
VEX_GLOBAL_CONSTANT( e, boost::math::constants::e<double>() );
VEX_GLOBAL_CONSTANT( half, boost::math::constants::half<double>() );
VEX_GLOBAL_CONSTANT( euler, boost::math::constants::euler<double>() );
VEX_GLOBAL_CONSTANT( root_two, boost::math::constants::root_two<double>() );
VEX_GLOBAL_CONSTANT( ln_two, boost::math::constants::ln_two<double>() );
VEX_GLOBAL_CONSTANT( ln_ln_two, boost::math::constants::ln_ln_two<double>() );
VEX_GLOBAL_CONSTANT( third, boost::math::constants::third<double>() );
VEX_GLOBAL_CONSTANT( twothirds, boost::math::constants::twothirds<double>() );
VEX_GLOBAL_CONSTANT( pi_minus_three, boost::math::constants::pi_minus_three<double>() );
VEX_GLOBAL_CONSTANT( four_minus_pi, boost::math::constants::four_minus_pi<double>() );
VEX_GLOBAL_CONSTANT( two_pi, boost::math::constants::two_pi<double>() );
VEX_GLOBAL_CONSTANT( one_div_two_pi, boost::math::constants::one_div_two_pi<double>() );
VEX_GLOBAL_CONSTANT( half_root_two, boost::math::constants::half_root_two<double>() );
VEX_GLOBAL_CONSTANT( pow23_four_minus_pi, boost::math::constants::pow23_four_minus_pi<double>() );
VEX_GLOBAL_CONSTANT( exp_minus_half, boost::math::constants::exp_minus_half<double>() );

#if (BOOST_VERSION >= 105000) || defined(DOXYGEN)
VEX_CONSTANT( catalan, boost::math::constants::catalan<double>() )
VEX_CONSTANT( cbrt_pi, boost::math::constants::cbrt_pi<double>() )
VEX_CONSTANT( cosh_one, boost::math::constants::cosh_one<double>() )
VEX_CONSTANT( cos_one, boost::math::constants::cos_one<double>() )
VEX_CONSTANT( degree, boost::math::constants::degree<double>() )
VEX_CONSTANT( e_pow_pi, boost::math::constants::e_pow_pi<double>() )
VEX_CONSTANT( euler_sqr, boost::math::constants::euler_sqr<double>() )
VEX_CONSTANT( extreme_value_skewness, boost::math::constants::extreme_value_skewness<double>() )
VEX_CONSTANT( four_thirds_pi, boost::math::constants::four_thirds_pi<double>() )
VEX_CONSTANT( glaisher, boost::math::constants::glaisher<double>() )
VEX_CONSTANT( half_pi, boost::math::constants::half_pi<double>() )
VEX_CONSTANT( khinchin, boost::math::constants::khinchin<double>() )
VEX_CONSTANT( ln_phi, boost::math::constants::ln_phi<double>() )
VEX_CONSTANT( ln_ten, boost::math::constants::ln_ten<double>() )
VEX_CONSTANT( log10_e, boost::math::constants::log10_e<double>() )
VEX_CONSTANT( one_div_cbrt_pi, boost::math::constants::one_div_cbrt_pi<double>() )
VEX_CONSTANT( one_div_euler, boost::math::constants::one_div_euler<double>() )
VEX_CONSTANT( one_div_ln_phi, boost::math::constants::one_div_ln_phi<double>() )
VEX_CONSTANT( one_div_log10_e, boost::math::constants::one_div_log10_e<double>() )
VEX_CONSTANT( one_div_root_pi, boost::math::constants::one_div_root_pi<double>() )
VEX_CONSTANT( one_div_root_two, boost::math::constants::one_div_root_two<double>() )
VEX_CONSTANT( one_div_root_two_pi, boost::math::constants::one_div_root_two_pi<double>() )
VEX_CONSTANT( phi, boost::math::constants::phi<double>() )
VEX_CONSTANT( pi_cubed, boost::math::constants::pi_cubed<double>() )
VEX_CONSTANT( pi_pow_e, boost::math::constants::pi_pow_e<double>() )
VEX_CONSTANT( pi_sqr, boost::math::constants::pi_sqr<double>() )
VEX_CONSTANT( pi_sqr_div_six, boost::math::constants::pi_sqr_div_six<double>() )
VEX_CONSTANT( radian, boost::math::constants::radian<double>() )
VEX_CONSTANT( rayleigh_kurtosis, boost::math::constants::rayleigh_kurtosis<double>() )
VEX_CONSTANT( rayleigh_kurtosis_excess, boost::math::constants::rayleigh_kurtosis_excess<double>() )
VEX_CONSTANT( rayleigh_skewness, boost::math::constants::rayleigh_skewness<double>() )
VEX_CONSTANT( root_e, boost::math::constants::root_e<double>() )
VEX_CONSTANT( root_one_div_pi, boost::math::constants::root_one_div_pi<double>() )
VEX_CONSTANT( root_three, boost::math::constants::root_three<double>() )
VEX_CONSTANT( root_two_div_pi, boost::math::constants::root_two_div_pi<double>() )
VEX_CONSTANT( sinh_one, boost::math::constants::sinh_one<double>() )
VEX_CONSTANT( sin_one, boost::math::constants::sin_one<double>() )
VEX_CONSTANT( sixth_pi, boost::math::constants::sixth_pi<double>() )
VEX_CONSTANT( third_pi, boost::math::constants::third_pi<double>() )
VEX_CONSTANT( three_quarters, boost::math::constants::three_quarters<double>() )
VEX_CONSTANT( three_quarters_pi, boost::math::constants::three_quarters_pi<double>() )
VEX_CONSTANT( two_div_pi, boost::math::constants::two_div_pi<double>() )
VEX_CONSTANT( two_thirds, boost::math::constants::two_thirds<double>() )
VEX_CONSTANT( two_thirds_pi, boost::math::constants::two_thirds_pi<double>() )
VEX_CONSTANT( zeta_three, boost::math::constants::zeta_three<double>() )
VEX_CONSTANT( zeta_two, boost::math::constants::zeta_two<double>() )
VEX_GLOBAL_CONSTANT( catalan, boost::math::constants::catalan<double>() );
VEX_GLOBAL_CONSTANT( cbrt_pi, boost::math::constants::cbrt_pi<double>() );
VEX_GLOBAL_CONSTANT( cosh_one, boost::math::constants::cosh_one<double>() );
VEX_GLOBAL_CONSTANT( cos_one, boost::math::constants::cos_one<double>() );
VEX_GLOBAL_CONSTANT( degree, boost::math::constants::degree<double>() );
VEX_GLOBAL_CONSTANT( e_pow_pi, boost::math::constants::e_pow_pi<double>() );
VEX_GLOBAL_CONSTANT( euler_sqr, boost::math::constants::euler_sqr<double>() );
VEX_GLOBAL_CONSTANT( extreme_value_skewness, boost::math::constants::extreme_value_skewness<double>() );
VEX_GLOBAL_CONSTANT( four_thirds_pi, boost::math::constants::four_thirds_pi<double>() );
VEX_GLOBAL_CONSTANT( glaisher, boost::math::constants::glaisher<double>() );
VEX_GLOBAL_CONSTANT( half_pi, boost::math::constants::half_pi<double>() );
VEX_GLOBAL_CONSTANT( khinchin, boost::math::constants::khinchin<double>() );
VEX_GLOBAL_CONSTANT( ln_phi, boost::math::constants::ln_phi<double>() );
VEX_GLOBAL_CONSTANT( ln_ten, boost::math::constants::ln_ten<double>() );
VEX_GLOBAL_CONSTANT( log10_e, boost::math::constants::log10_e<double>() );
VEX_GLOBAL_CONSTANT( one_div_cbrt_pi, boost::math::constants::one_div_cbrt_pi<double>() );
VEX_GLOBAL_CONSTANT( one_div_euler, boost::math::constants::one_div_euler<double>() );
VEX_GLOBAL_CONSTANT( one_div_ln_phi, boost::math::constants::one_div_ln_phi<double>() );
VEX_GLOBAL_CONSTANT( one_div_log10_e, boost::math::constants::one_div_log10_e<double>() );
VEX_GLOBAL_CONSTANT( one_div_root_pi, boost::math::constants::one_div_root_pi<double>() );
VEX_GLOBAL_CONSTANT( one_div_root_two, boost::math::constants::one_div_root_two<double>() );
VEX_GLOBAL_CONSTANT( one_div_root_two_pi, boost::math::constants::one_div_root_two_pi<double>() );
VEX_GLOBAL_CONSTANT( phi, boost::math::constants::phi<double>() );
VEX_GLOBAL_CONSTANT( pi_cubed, boost::math::constants::pi_cubed<double>() );
VEX_GLOBAL_CONSTANT( pi_pow_e, boost::math::constants::pi_pow_e<double>() );
VEX_GLOBAL_CONSTANT( pi_sqr, boost::math::constants::pi_sqr<double>() );
VEX_GLOBAL_CONSTANT( pi_sqr_div_six, boost::math::constants::pi_sqr_div_six<double>() );
VEX_GLOBAL_CONSTANT( radian, boost::math::constants::radian<double>() );
VEX_GLOBAL_CONSTANT( rayleigh_kurtosis, boost::math::constants::rayleigh_kurtosis<double>() );
VEX_GLOBAL_CONSTANT( rayleigh_kurtosis_excess, boost::math::constants::rayleigh_kurtosis_excess<double>() );
VEX_GLOBAL_CONSTANT( rayleigh_skewness, boost::math::constants::rayleigh_skewness<double>() );
VEX_GLOBAL_CONSTANT( root_e, boost::math::constants::root_e<double>() );
VEX_GLOBAL_CONSTANT( root_one_div_pi, boost::math::constants::root_one_div_pi<double>() );
VEX_GLOBAL_CONSTANT( root_three, boost::math::constants::root_three<double>() );
VEX_GLOBAL_CONSTANT( root_two_div_pi, boost::math::constants::root_two_div_pi<double>() );
VEX_GLOBAL_CONSTANT( sinh_one, boost::math::constants::sinh_one<double>() );
VEX_GLOBAL_CONSTANT( sin_one, boost::math::constants::sin_one<double>() );
VEX_GLOBAL_CONSTANT( sixth_pi, boost::math::constants::sixth_pi<double>() );
VEX_GLOBAL_CONSTANT( third_pi, boost::math::constants::third_pi<double>() );
VEX_GLOBAL_CONSTANT( three_quarters, boost::math::constants::three_quarters<double>() );
VEX_GLOBAL_CONSTANT( three_quarters_pi, boost::math::constants::three_quarters_pi<double>() );
VEX_GLOBAL_CONSTANT( two_div_pi, boost::math::constants::two_div_pi<double>() );
VEX_GLOBAL_CONSTANT( two_thirds, boost::math::constants::two_thirds<double>() );
VEX_GLOBAL_CONSTANT( two_thirds_pi, boost::math::constants::two_thirds_pi<double>() );
VEX_GLOBAL_CONSTANT( zeta_three, boost::math::constants::zeta_three<double>() );
VEX_GLOBAL_CONSTANT( zeta_two, boost::math::constants::zeta_two<double>() );
#endif


} //namespace constants

} // namespace vex

#endif

0 comments on commit 31ccbdd

Please sign in to comment.