All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
scalar_operations.hpp
Go to the documentation of this file.
#ifndef VIENNACL_LINALG_CUDA_SCALAR_OPERATIONS_HPP_
#define VIENNACL_LINALG_CUDA_SCALAR_OPERATIONS_HPP_
/* =========================================================================
Copyright (c) 2010-2016, Institute for Microelectronics,
Institute for Analysis and Scientific Computing,
TU Wien.
Portions of this software are copyright by UChicago Argonne, LLC.
-----------------
ViennaCL - The Vienna Computing Library
-----------------
Project Head: Karl Rupp rupp@iue.tuwien.ac.at
(A list of authors and contributors can be found in the manual)
License: MIT (X11), see file LICENSE in the base directory
============================================================================= */
#include "viennacl/forwards.h"
#include "viennacl/tools/tools.hpp"
#include "viennacl/meta/predicate.hpp"
#include "viennacl/meta/enable_if.hpp"
#include "viennacl/traits/size.hpp"
#include "viennacl/traits/start.hpp"
#include "viennacl/traits/stride.hpp"
#include "viennacl/linalg/cuda/common.hpp"
// includes CUDA
#include <cuda_runtime.h>
namespace viennacl
{
namespace linalg
{
namespace cuda
{
template<typename NumericT>
__global__ void as_kernel(NumericT * s1, const NumericT * fac2, unsigned int options2, const NumericT * s2)
{
NumericT alpha = *fac2;
if (options2 & (1 << 0))
alpha = -alpha;
if (options2 & (1 << 1))
alpha = NumericT(1) / alpha;
*s1 = *s2 * alpha;
}
template<typename NumericT>
__global__ void as_kernel(NumericT * s1, NumericT fac2, unsigned int options2, const NumericT * s2)
{
NumericT alpha = fac2;
if (options2 & (1 << 0))
alpha = -alpha;
if (options2 & (1 << 1))
alpha = NumericT(1) / alpha;
*s1 = *s2 * alpha;
}
template<typename ScalarT1,
typename ScalarT2, typename NumericT>
typename viennacl::enable_if< viennacl::is_scalar<ScalarT1>::value
&& viennacl::is_scalar<ScalarT2>::value
&& viennacl::is_any_scalar<NumericT>::value
>::type
as(ScalarT1 & s1,
ScalarT2 const & s2, NumericT const & alpha, vcl_size_t len_alpha, bool reciprocal_alpha, bool flip_sign_alpha)
{
typedef typename viennacl::result_of::cpu_value_type<ScalarT1>::type value_type;
unsigned int options_alpha = detail::make_options(len_alpha, reciprocal_alpha, flip_sign_alpha);
value_type temporary_alpha = 0;
if (viennacl::is_cpu_scalar<NumericT>::value)
temporary_alpha = alpha;
as_kernel<<<1, 1>>>(viennacl::cuda_arg(s1),
viennacl::cuda_arg<value_type>(detail::arg_reference(alpha, temporary_alpha)),
options_alpha,
viennacl::cuda_arg(s2));
VIENNACL_CUDA_LAST_ERROR_CHECK("as_kernel");
}
// alpha and beta on GPU
template<typename NumericT>
__global__ void asbs_kernel(NumericT * s1,
const NumericT * fac2, unsigned int options2, const NumericT * s2,
const NumericT * fac3, unsigned int options3, const NumericT * s3)
{
NumericT alpha = *fac2;
if (options2 & (1 << 0))
alpha = -alpha;
if (options2 & (1 << 1))
alpha = NumericT(1) / alpha;
NumericT beta = *fac3;
if (options3 & (1 << 0))
beta = -beta;
if (options3 & (1 << 1))
beta = NumericT(1) / beta;
*s1 = *s2 * alpha + *s3 * beta;
}
// alpha on CPU, beta on GPU
template<typename NumericT>
__global__ void asbs_kernel(NumericT * s1,
NumericT fac2, unsigned int options2, const NumericT * s2,
NumericT const * fac3, unsigned int options3, const NumericT * s3)
{
NumericT alpha = fac2;
if (options2 & (1 << 0))
alpha = -alpha;
if (options2 & (1 << 1))
alpha = NumericT(1) / alpha;
NumericT beta = *fac3;
if (options3 & (1 << 0))
beta = -beta;
if (options3 & (1 << 1))
beta = NumericT(1) / beta;
*s1 = *s2 * alpha + *s3 * beta;
}
// alpha on GPU, beta on CPU
template<typename NumericT>
__global__ void asbs_kernel(NumericT * s1,
NumericT const * fac2, unsigned int options2, const NumericT * s2,
NumericT fac3, unsigned int options3, const NumericT * s3)
{
NumericT alpha = *fac2;
if (options2 & (1 << 0))
alpha = -alpha;
if (options2 & (1 << 1))
alpha = NumericT(1) / alpha;
NumericT beta = fac3;
if (options3 & (1 << 0))
beta = -beta;
if (options3 & (1 << 1))
beta = NumericT(1) / beta;
*s1 = *s2 * alpha + *s3 * beta;
}
// alpha and beta on CPU
template<typename NumericT>
__global__ void asbs_kernel(NumericT * s1,
NumericT fac2, unsigned int options2, const NumericT * s2,
NumericT fac3, unsigned int options3, const NumericT * s3)
{
NumericT alpha = fac2;
if (options2 & (1 << 0))
alpha = -alpha;
if (options2 & (1 << 1))
alpha = NumericT(1) / alpha;
NumericT beta = fac3;
if (options3 & (1 << 0))
beta = -beta;
if (options3 & (1 << 1))
beta = NumericT(1) / beta;
*s1 = *s2 * alpha + *s3 * beta;
}
template<typename ScalarT1,
typename ScalarT2, typename NumericT1,
typename ScalarT3, typename NumericT2>
typename viennacl::enable_if< viennacl::is_scalar<ScalarT1>::value
&& viennacl::is_scalar<ScalarT2>::value
&& viennacl::is_scalar<ScalarT3>::value
&& viennacl::is_any_scalar<NumericT1>::value
&& viennacl::is_any_scalar<NumericT2>::value
>::type
asbs(ScalarT1 & s1,
ScalarT2 const & s2, NumericT1 const & alpha, vcl_size_t len_alpha, bool reciprocal_alpha, bool flip_sign_alpha,
ScalarT3 const & s3, NumericT2 const & beta, vcl_size_t len_beta, bool reciprocal_beta, bool flip_sign_beta)
{
typedef typename viennacl::result_of::cpu_value_type<ScalarT1>::type value_type;
unsigned int options_alpha = detail::make_options(len_alpha, reciprocal_alpha, flip_sign_alpha);
unsigned int options_beta = detail::make_options(len_beta, reciprocal_beta, flip_sign_beta);
value_type temporary_alpha = 0;
if (viennacl::is_cpu_scalar<NumericT1>::value)
temporary_alpha = alpha;
value_type temporary_beta = 0;
if (viennacl::is_cpu_scalar<NumericT2>::value)
temporary_beta = beta;
asbs_kernel<<<1, 1>>>(viennacl::cuda_arg(s1),
viennacl::cuda_arg<value_type>(detail::arg_reference(alpha, temporary_alpha)),
options_alpha,
viennacl::cuda_arg(s2),
viennacl::cuda_arg<value_type>(detail::arg_reference(beta, temporary_beta)),
options_beta,
viennacl::cuda_arg(s3) );
VIENNACL_CUDA_LAST_ERROR_CHECK("asbs_kernel");
}
// alpha and beta on GPU
template<typename NumericT>
__global__ void asbs_s_kernel(NumericT * s1,
const NumericT * fac2, unsigned int options2, const NumericT * s2,
const NumericT * fac3, unsigned int options3, const NumericT * s3)
{
NumericT alpha = *fac2;
if (options2 & (1 << 0))
alpha = -alpha;
if (options2 & (1 << 1))
alpha = NumericT(1) / alpha;
NumericT beta = *fac3;
if (options3 & (1 << 0))
beta = -beta;
if (options3 & (1 << 1))
beta = NumericT(1) / beta;
*s1 += *s2 * alpha + *s3 * beta;
}
// alpha on CPU, beta on GPU
template<typename NumericT>
__global__ void asbs_s_kernel(NumericT * s1,
NumericT fac2, unsigned int options2, const NumericT * s2,
NumericT const * fac3, unsigned int options3, const NumericT * s3)
{
NumericT alpha = fac2;
if (options2 & (1 << 0))
alpha = -alpha;
if (options2 & (1 << 1))
alpha = NumericT(1) / alpha;
NumericT beta = *fac3;
if (options3 & (1 << 0))
beta = -beta;
if (options3 & (1 << 1))
beta = NumericT(1) / beta;
*s1 += *s2 * alpha + *s3 * beta;
}
// alpha on GPU, beta on CPU
template<typename NumericT>
__global__ void asbs_s_kernel(NumericT * s1,
NumericT const * fac2, unsigned int options2, const NumericT * s2,
NumericT fac3, unsigned int options3, const NumericT * s3)
{
NumericT alpha = *fac2;
if (options2 & (1 << 0))
alpha = -alpha;
if (options2 & (1 << 1))
alpha = NumericT(1) / alpha;
NumericT beta = fac3;
if (options3 & (1 << 0))
beta = -beta;
if (options3 & (1 << 1))
beta = NumericT(1) / beta;
*s1 += *s2 * alpha + *s3 * beta;
}
// alpha and beta on CPU
template<typename NumericT>
__global__ void asbs_s_kernel(NumericT * s1,
NumericT fac2, unsigned int options2, const NumericT * s2,
NumericT fac3, unsigned int options3, const NumericT * s3)
{
NumericT alpha = fac2;
if (options2 & (1 << 0))
alpha = -alpha;
if (options2 & (1 << 1))
alpha = NumericT(1) / alpha;
NumericT beta = fac3;
if (options3 & (1 << 0))
beta = -beta;
if (options3 & (1 << 1))
beta = NumericT(1) / beta;
*s1 += *s2 * alpha + *s3 * beta;
}
template<typename ScalarT1,
typename ScalarT2, typename NumericT1,
typename ScalarT3, typename NumericT2>
typename viennacl::enable_if< viennacl::is_scalar<ScalarT1>::value
&& viennacl::is_scalar<ScalarT2>::value
&& viennacl::is_scalar<ScalarT3>::value
&& viennacl::is_any_scalar<NumericT1>::value
&& viennacl::is_any_scalar<NumericT2>::value
>::type
asbs_s(ScalarT1 & s1,
ScalarT2 const & s2, NumericT1 const & alpha, vcl_size_t len_alpha, bool reciprocal_alpha, bool flip_sign_alpha,
ScalarT3 const & s3, NumericT2 const & beta, vcl_size_t len_beta, bool reciprocal_beta, bool flip_sign_beta)
{
typedef typename viennacl::result_of::cpu_value_type<ScalarT1>::type value_type;
unsigned int options_alpha = detail::make_options(len_alpha, reciprocal_alpha, flip_sign_alpha);
unsigned int options_beta = detail::make_options(len_beta, reciprocal_beta, flip_sign_beta);
value_type temporary_alpha = 0;
if (viennacl::is_cpu_scalar<NumericT1>::value)
temporary_alpha = alpha;
value_type temporary_beta = 0;
if (viennacl::is_cpu_scalar<NumericT2>::value)
temporary_beta = beta;
std::cout << "Launching asbs_s_kernel..." << std::endl;
asbs_s_kernel<<<1, 1>>>(viennacl::cuda_arg(s1),
viennacl::cuda_arg<value_type>(detail::arg_reference(alpha, temporary_alpha)),
options_alpha,
viennacl::cuda_arg(s2),
viennacl::cuda_arg<value_type>(detail::arg_reference(beta, temporary_beta)),
options_beta,
viennacl::cuda_arg(s3) );
VIENNACL_CUDA_LAST_ERROR_CHECK("asbs_s_kernel");
}
template<typename NumericT>
__global__ void scalar_swap_kernel(NumericT * s1, NumericT * s2)
{
NumericT tmp = *s2;
*s2 = *s1;
*s1 = tmp;
}
template<typename ScalarT1, typename ScalarT2>
typename viennacl::enable_if< viennacl::is_scalar<ScalarT1>::value
&& viennacl::is_scalar<ScalarT2>::value
>::type
swap(ScalarT1 & s1, ScalarT2 & s2)
{
typedef typename viennacl::result_of::cpu_value_type<ScalarT1>::type value_type;
scalar_swap_kernel<<<1, 1>>>(viennacl::cuda_arg(s1), viennacl::cuda_arg(s2));
}
} //namespace single_threaded
} //namespace linalg
} //namespace viennacl
#endif
viennacl::enable_if
Simple enable-if variant that uses the SFINAE pattern.
Definition: enable_if.hpp:30
viennacl::linalg::cuda::detail::make_options
unsigned int make_options(vcl_size_t length, bool reciprocal, bool flip_sign)
Definition: common.hpp:160
size.hpp
Generic size and resize functionality for different vector and matrix types.
start.hpp
Extracts the underlying OpenCL start index handle from a vector, a matrix, an expression etc...
tools.hpp
Various little tools used here and there in ViennaCL.
forwards.h
This file provides the forward declarations for the main types used within ViennaCL.
stride.hpp
Determines row and column increments for matrices and matrix proxies.
s2
viennacl::scalar< int > s2
Definition: global_variables.cpp:58
s1
viennacl::scalar< float > s1
Definition: global_variables.cpp:57
viennacl::linalg::cuda::asbs
viennacl::enable_if< viennacl::is_scalar< ScalarT1 >::value &&viennacl::is_scalar< ScalarT2 >::value &&viennacl::is_scalar< ScalarT3 >::value &&viennacl::is_any_scalar< NumericT1 >::value &&viennacl::is_any_scalar< NumericT2 >::value >::type asbs(ScalarT1 &s1, ScalarT2 const &s2, NumericT1 const &alpha, vcl_size_t len_alpha, bool reciprocal_alpha, bool flip_sign_alpha, ScalarT3 const &s3, NumericT2 const &beta, vcl_size_t len_beta, bool reciprocal_beta, bool flip_sign_beta)
Definition: scalar_operations.hpp:191
NumericT
float NumericT
Definition: bisect.cpp:40
viennacl::linalg::cuda::as_kernel
__global__ void as_kernel(NumericT *s1, const NumericT *fac2, unsigned int options2, const NumericT *s2)
Definition: scalar_operations.hpp:48
viennacl::is_cpu_scalar
Helper struct for checking whether a type is a host scalar type (e.g. float, double) ...
Definition: forwards.h:448
viennacl::is_any_scalar
Helper struct for checking whether the provided type represents a scalar (either host, from ViennaCL, or a flip-sign proxy)
Definition: forwards.h:469
viennacl::vcl_size_t
std::size_t vcl_size_t
Definition: forwards.h:75
viennacl::result_of::cpu_value_type::type
T::ERROR_CANNOT_DEDUCE_CPU_SCALAR_TYPE_FOR_T type
Definition: result_of.hpp:271
viennacl::linalg::cuda::swap
viennacl::enable_if< viennacl::is_scalar< ScalarT1 >::value &&viennacl::is_scalar< ScalarT2 >::value >::type swap(ScalarT1 &s1, ScalarT2 &s2)
Swaps the contents of two scalars, data is copied.
Definition: scalar_operations.hpp:361
viennacl::linalg::cuda::asbs_s
viennacl::enable_if< viennacl::is_scalar< ScalarT1 >::value &&viennacl::is_scalar< ScalarT2 >::value &&viennacl::is_scalar< ScalarT3 >::value &&viennacl::is_any_scalar< NumericT1 >::value &&viennacl::is_any_scalar< NumericT2 >::value >::type asbs_s(ScalarT1 &s1, ScalarT2 const &s2, NumericT1 const &alpha, vcl_size_t len_alpha, bool reciprocal_alpha, bool flip_sign_alpha, ScalarT3 const &s3, NumericT2 const &beta, vcl_size_t len_beta, bool reciprocal_beta, bool flip_sign_beta)
Definition: scalar_operations.hpp:314
predicate.hpp
All the predicates used within ViennaCL. Checks for expressions to be vectors, etc.
viennacl::is_scalar
Helper struct for checking whether a type is a viennacl::scalar<>
Definition: forwards.h:455
common.hpp
Common routines for CUDA execution.
viennacl::linalg::cuda::as
viennacl::enable_if< viennacl::is_scalar< ScalarT1 >::value &&viennacl::is_scalar< ScalarT2 >::value &&viennacl::is_any_scalar< NumericT >::value >::type as(ScalarT1 &s1, ScalarT2 const &s2, NumericT const &alpha, vcl_size_t len_alpha, bool reciprocal_alpha, bool flip_sign_alpha)
Definition: scalar_operations.hpp:77
viennacl::linalg::cuda::scalar_swap_kernel
__global__ void scalar_swap_kernel(NumericT *s1, NumericT *s2)
Definition: scalar_operations.hpp:345
viennacl::linalg::cuda::asbs_kernel
__global__ void asbs_kernel(NumericT *s1, const NumericT *fac2, unsigned int options2, const NumericT *s2, const NumericT *fac3, unsigned int options3, const NumericT *s3)
Definition: scalar_operations.hpp:99
VIENNACL_CUDA_LAST_ERROR_CHECK
#define VIENNACL_CUDA_LAST_ERROR_CHECK(message)
Definition: common.hpp:30
viennacl::cuda_arg
NumericT * cuda_arg(scalar< NumericT > &obj)
Convenience helper function for extracting the CUDA handle from a ViennaCL scalar. Non-const version.
Definition: common.hpp:39
viennacl::linalg::cuda::asbs_s_kernel
__global__ void asbs_s_kernel(NumericT *s1, const NumericT *fac2, unsigned int options2, const NumericT *s2, const NumericT *fac3, unsigned int options3, const NumericT *s3)
Definition: scalar_operations.hpp:222
viennacl::linalg::cuda::detail::arg_reference
viennacl::backend::mem_handle::cuda_handle_type & arg_reference(viennacl::scalar< NumericT > &s, OtherT)
Definition: common.hpp:188
enable_if.hpp
Simple enable-if variant that uses the SFINAE pattern.