1 #ifndef VIENNACL_LINALG_CUDA_SCALAR_OPERATIONS_HPP_
2 #define VIENNACL_LINALG_CUDA_SCALAR_OPERATIONS_HPP_
35 #include <cuda_runtime.h>
53 __global__
void as_kernel(T * s1,
const T * fac2,
unsigned int options2,
const T * s2)
56 if (options2 & (1 << 0))
58 if (options2 & (1 << 1))
59 alpha = ((T)(1)) / alpha;
65 __global__
void as_kernel(T * s1, T fac2,
unsigned int options2,
const T * s2)
68 if (options2 & (1 << 0))
70 if (options2 & (1 << 1))
71 alpha = ((T)(1)) / alpha;
76 template <
typename S1,
77 typename S2,
typename ScalarType1>
83 S2
const & s2, ScalarType1
const & alpha,
vcl_size_t len_alpha,
bool reciprocal_alpha,
bool flip_sign_alpha)
89 value_type temporary_alpha = 0;
91 temporary_alpha = alpha;
93 as_kernel<<<1, 1>>>(detail::cuda_arg<value_type>(s1),
96 detail::cuda_arg<value_type>(s2));
103 template <
typename T>
105 const T * fac2,
unsigned int options2,
const T * s2,
106 const T * fac3,
unsigned int options3,
const T * s3)
109 if (options2 & (1 << 0))
111 if (options2 & (1 << 1))
112 alpha = ((T)(1)) / alpha;
115 if (options3 & (1 << 0))
117 if (options3 & (1 << 1))
118 beta = ((T)(1)) / beta;
120 *s1 = *s2 * alpha + *s3 * beta;
124 template <
typename T>
126 T fac2,
unsigned int options2,
const T * s2,
127 const T * fac3,
unsigned int options3,
const T * s3)
130 if (options2 & (1 << 0))
132 if (options2 & (1 << 1))
133 alpha = ((T)(1)) / alpha;
136 if (options3 & (1 << 0))
138 if (options3 & (1 << 1))
139 beta = ((T)(1)) / beta;
141 *s1 = *s2 * alpha + *s3 * beta;
145 template <
typename T>
147 const T * fac2,
unsigned int options2,
const T * s2,
148 T fac3,
unsigned int options3,
const T * s3)
151 if (options2 & (1 << 0))
153 if (options2 & (1 << 1))
154 alpha = ((T)(1)) / alpha;
157 if (options3 & (1 << 0))
159 if (options3 & (1 << 1))
160 beta = ((T)(1)) / beta;
162 *s1 = *s2 * alpha + *s3 * beta;
166 template <
typename T>
168 T fac2,
unsigned int options2,
const T * s2,
169 T fac3,
unsigned int options3,
const T * s3)
172 if (options2 & (1 << 0))
174 if (options2 & (1 << 1))
175 alpha = ((T)(1)) / alpha;
178 if (options3 & (1 << 0))
180 if (options3 & (1 << 1))
181 beta = ((T)(1)) / beta;
183 *s1 = *s2 * alpha + *s3 * beta;
187 template <
typename S1,
188 typename S2,
typename ScalarType1,
189 typename S3,
typename ScalarType2>
197 S2
const & s2, ScalarType1
const & alpha,
vcl_size_t len_alpha,
bool reciprocal_alpha,
bool flip_sign_alpha,
198 S3
const & s3, ScalarType2
const & beta,
vcl_size_t len_beta,
bool reciprocal_beta,
bool flip_sign_beta)
202 unsigned int options_alpha =
detail::make_options(len_alpha, reciprocal_alpha, flip_sign_alpha);
205 value_type temporary_alpha = 0;
207 temporary_alpha = alpha;
209 value_type temporary_beta = 0;
211 temporary_beta = beta;
213 asbs_kernel<<<1, 1>>>(detail::cuda_arg<value_type>(s1),
216 detail::cuda_arg<value_type>(s2),
219 detail::cuda_arg<value_type>(s3) );
226 template <
typename T>
228 const T * fac2,
unsigned int options2,
const T * s2,
229 const T * fac3,
unsigned int options3,
const T * s3)
232 if (options2 & (1 << 0))
234 if (options2 & (1 << 1))
235 alpha = ((T)(1)) / alpha;
238 if (options3 & (1 << 0))
240 if (options3 & (1 << 1))
241 beta = ((T)(1)) / beta;
243 *s1 += *s2 * alpha + *s3 * beta;
247 template <
typename T>
249 T fac2,
unsigned int options2,
const T * s2,
250 const T * fac3,
unsigned int options3,
const T * s3)
253 if (options2 & (1 << 0))
255 if (options2 & (1 << 1))
256 alpha = ((T)(1)) / alpha;
259 if (options3 & (1 << 0))
261 if (options3 & (1 << 1))
262 beta = ((T)(1)) / beta;
264 *s1 += *s2 * alpha + *s3 * beta;
268 template <
typename T>
270 const T * fac2,
unsigned int options2,
const T * s2,
271 T fac3,
unsigned int options3,
const T * s3)
274 if (options2 & (1 << 0))
276 if (options2 & (1 << 1))
277 alpha = ((T)(1)) / alpha;
280 if (options3 & (1 << 0))
282 if (options3 & (1 << 1))
283 beta = ((T)(1)) / beta;
285 *s1 += *s2 * alpha + *s3 * beta;
289 template <
typename T>
291 T fac2,
unsigned int options2,
const T * s2,
292 T fac3,
unsigned int options3,
const T * s3)
295 if (options2 & (1 << 0))
297 if (options2 & (1 << 1))
298 alpha = ((T)(1)) / alpha;
301 if (options3 & (1 << 0))
303 if (options3 & (1 << 1))
304 beta = ((T)(1)) / beta;
306 *s1 += *s2 * alpha + *s3 * beta;
310 template <
typename S1,
311 typename S2,
typename ScalarType1,
312 typename S3,
typename ScalarType2>
320 S2
const & s2, ScalarType1
const & alpha,
vcl_size_t len_alpha,
bool reciprocal_alpha,
bool flip_sign_alpha,
321 S3
const & s3, ScalarType2
const & beta,
vcl_size_t len_beta,
bool reciprocal_beta,
bool flip_sign_beta)
325 unsigned int options_alpha =
detail::make_options(len_alpha, reciprocal_alpha, flip_sign_alpha);
328 value_type temporary_alpha = 0;
330 temporary_alpha = alpha;
332 value_type temporary_beta = 0;
334 temporary_beta = beta;
336 std::cout <<
"Launching asbs_s_kernel..." << std::endl;
337 asbs_s_kernel<<<1, 1>>>(detail::cuda_arg<value_type>(s1),
340 detail::cuda_arg<value_type>(s2),
343 detail::cuda_arg<value_type>(s3) );
349 template <
typename T>
362 template <
typename S1,
typename S2>
370 scalar_swap_kernel<<<1, 1>>>(detail::cuda_arg<value_type>(s1),detail::cuda_arg<value_type>(s2));
__global__ void as_kernel(T *s1, const T *fac2, unsigned int options2, const T *s2)
Definition: scalar_operations.hpp:53
Simple enable-if variant that uses the SFINAE pattern.
Definition: enable_if.hpp:29
std::size_t vcl_size_t
Definition: forwards.h:58
unsigned int make_options(vcl_size_t length, bool reciprocal, bool flip_sign)
Definition: common.hpp:37
Generic size and resize functionality for different vector and matrix types.
Common routines for CUDA execution.
viennacl::backend::mem_handle::cuda_handle_type & arg_reference(viennacl::scalar< T > &s, U)
Definition: common.hpp:127
Extracts the underlying OpenCL start index handle from a vector, a matrix, an expression etc...
This file provides the forward declarations for the main types used within ViennaCL.
Determines row and column increments for matrices and matrix proxies.
#define VIENNACL_CUDA_LAST_ERROR_CHECK(message)
Definition: common.hpp:27
__global__ void asbs_s_kernel(T *s1, const T *fac2, unsigned int options2, const T *s2, const T *fac3, unsigned int options3, const T *s3)
Definition: scalar_operations.hpp:227
__global__ void scalar_swap_kernel(T *s1, T *s2)
Definition: scalar_operations.hpp:350
Helper struct for checking whether a type is a host scalar type (e.g. float, double) ...
Definition: forwards.h:363
viennacl::enable_if< viennacl::is_scalar< S1 >::value &&viennacl::is_scalar< S2 >::value &&viennacl::is_scalar< S3 >::value &&viennacl::is_any_scalar< ScalarType1 >::value &&viennacl::is_any_scalar< ScalarType2 >::value >::type asbs(S1 &s1, S2 const &s2, ScalarType1 const &alpha, vcl_size_t len_alpha, bool reciprocal_alpha, bool flip_sign_alpha, S3 const &s3, ScalarType2 const &beta, vcl_size_t len_beta, bool reciprocal_beta, bool flip_sign_beta)
Definition: scalar_operations.hpp:196
Helper struct for checking whether the provided type represents a scalar (either host, from ViennaCL, or a flip-sign proxy)
Definition: forwards.h:384
viennacl::enable_if< viennacl::is_scalar< S1 >::value &&viennacl::is_scalar< S2 >::value &&viennacl::is_any_scalar< ScalarType1 >::value >::type as(S1 &s1, S2 const &s2, ScalarType1 const &alpha, vcl_size_t len_alpha, bool reciprocal_alpha, bool flip_sign_alpha)
Definition: scalar_operations.hpp:82
__global__ void asbs_kernel(T *s1, const T *fac2, unsigned int options2, const T *s2, const T *fac3, unsigned int options3, const T *s3)
Definition: scalar_operations.hpp:104
All the predicates used within ViennaCL. Checks for expressions to be vectors, etc.
Helper struct for checking whether a type is a viennacl::scalar<>
Definition: forwards.h:370
T::ERROR_CANNOT_DEDUCE_CPU_SCALAR_TYPE_FOR_T type
Definition: result_of.hpp:276
viennacl::enable_if< viennacl::is_scalar< S1 >::value &&viennacl::is_scalar< S2 >::value >::type swap(S1 &s1, S2 &s2)
Swaps the contents of two scalars, data is copied.
Definition: scalar_operations.hpp:366
viennacl::enable_if< viennacl::is_scalar< S1 >::value &&viennacl::is_scalar< S2 >::value &&viennacl::is_scalar< S3 >::value &&viennacl::is_any_scalar< ScalarType1 >::value &&viennacl::is_any_scalar< ScalarType2 >::value >::type asbs_s(S1 &s1, S2 const &s2, ScalarType1 const &alpha, vcl_size_t len_alpha, bool reciprocal_alpha, bool flip_sign_alpha, S3 const &s3, ScalarType2 const &beta, vcl_size_t len_beta, bool reciprocal_beta, bool flip_sign_beta)
Definition: scalar_operations.hpp:319
Simple enable-if variant that uses the SFINAE pattern.