1 #ifndef VIENNACL_LINALG_OPENCL_KERNELS_SCALAR_HPP
2 #define VIENNACL_LINALG_OPENCL_KERNELS_SCALAR_HPP
42 template <
typename StringType>
45 source.append(
" *s1 "); source.append(cfg.
assign_op); source.append(1, sign_a); source.append(
" *s2 ");
47 source.append(
"* alpha ");
49 source.append(
"/ alpha ");
52 source.append(1, sign_b); source.append(
" *s3 ");
54 source.append(
"* beta");
56 source.append(
"/ beta");
58 source.append(
"; \n");
61 template <
typename StringType>
64 source.append(
" if (options2 & (1 << 1)) { \n");
67 source.append(
" if (options3 & (1 << 1)) \n");
69 source.append(
" else \n");
74 source.append(
" } else { \n");
77 source.append(
" if (options3 & (1 << 1)) \n");
79 source.append(
" else \n");
84 source.append(
" } \n");
88 template <
typename StringType>
91 source.append(
"__kernel void as");
98 source.append(
"_cpu");
100 source.append(
"_gpu");
103 source.append(
"_cpu");
105 source.append(
"_gpu");
106 source.append(
"( \n");
107 source.append(
" __global "); source.append(numeric_string); source.append(
" * s1, \n");
108 source.append(
" \n");
111 source.append(
" "); source.append(numeric_string); source.append(
" fac2, \n");
115 source.append(
" __global "); source.append(numeric_string); source.append(
" * fac2, \n");
117 source.append(
" unsigned int options2, \n");
118 source.append(
" __global const "); source.append(numeric_string); source.append(
" * s2");
122 source.append(
", \n\n");
125 source.append(
" "); source.append(numeric_string); source.append(
" fac3, \n");
129 source.append(
" __global "); source.append(numeric_string); source.append(
" * fac3, \n");
131 source.append(
" unsigned int options3, \n");
132 source.append(
" __global const "); source.append(numeric_string); source.append(
" * s3");
134 source.append(
") \n{ \n");
138 source.append(
" "); source.append(numeric_string); source.append(
" alpha = fac2; \n");
142 source.append(
" "); source.append(numeric_string); source.append(
" alpha = fac2[0]; \n");
144 source.append(
" \n");
148 source.append(
" "); source.append(numeric_string); source.append(
" beta = fac3; \n");
152 source.append(
" "); source.append(numeric_string); source.append(
" beta = fac3[0]; \n");
155 source.append(
" if (options2 & (1 << 0)) { \n");
158 source.append(
" if (options3 & (1 << 0)) { \n");
160 source.append(
" } else { \n");
162 source.append(
" } \n");
166 source.append(
" } else { \n");
169 source.append(
" if (options3 & (1 << 0)) { \n");
171 source.append(
" } else { \n");
173 source.append(
" } \n");
178 source.append(
" } \n");
179 source.append(
"} \n");
182 template <
typename StringType>
208 template <
typename StringType>
211 source.append(
"__kernel void swap( \n");
212 source.append(
" __global "); source.append(numeric_string); source.append(
" * s1, \n");
213 source.append(
" __global "); source.append(numeric_string); source.append(
" * s2) \n");
214 source.append(
"{ \n");
215 source.append(
" "); source.append(numeric_string); source.append(
" tmp = *s2; \n");
216 source.append(
" *s2 = *s1; \n");
217 source.append(
" *s1 = tmp; \n");
218 source.append(
"} \n");
225 template <
class TYPE>
238 static std::map<cl_context, bool> init_done;
242 source.reserve(8192);
244 viennacl::ocl::append_double_precision_pragma<TYPE>(ctx, source);
252 #ifdef VIENNACL_BUILD_INFO
253 std::cout <<
"Creating program " << prog_name << std::endl;
255 ctx.add_program(source, prog_name);
256 init_done[ctx.handle().get()] =
true;
Manages an OpenCL context and provides the respective convenience functions for creating buffers...
Definition: context.hpp:51
Provides OpenCL-related utilities.
void generate_asbs_impl2(StringType &source, char sign_a, char sign_b, asbs_config const &cfg)
Definition: scalar.hpp:62
bool with_stride_and_range
Definition: scalar.hpp:35
const OCL_TYPE & get() const
Definition: handle.hpp:189
asbs_config()
Definition: scalar.hpp:33
Definition: scalar.hpp:25
Definition: scalar.hpp:27
const viennacl::ocl::handle< cl_context > & handle() const
Returns the context handle.
Definition: context.hpp:476
Main namespace in ViennaCL. Holds all the basic types such as vector, matrix, etc. and defines operations upon them.
Definition: cpu_ram.hpp:29
asbs_scalar_type
Enumeration for the scalar type in avbv-like operations.
Definition: scalar.hpp:23
void generate_asbs_impl3(StringType &source, char sign_a, char sign_b, asbs_config const &cfg, bool mult_alpha, bool mult_beta)
Definition: scalar.hpp:43
Main kernel class for generating OpenCL kernels for operations involving viennacl::scalar<>, but not viennacl::vector<> or viennacl::matrix<>.
Definition: scalar.hpp:226
Definition: scalar.hpp:26
std::string assign_op
Definition: scalar.hpp:36
void generate_scalar_swap(StringType &source, std::string const &numeric_string)
Definition: scalar.hpp:209
void generate_asbs_impl(StringType &source, std::string const &numeric_string, asbs_config const &cfg)
Definition: scalar.hpp:89
Configuration struct for generating OpenCL kernels for linear combinations of viennacl::scalar<> obje...
Definition: scalar.hpp:31
static void apply(viennacl::ocl::context const &)
Definition: utils.hpp:40
asbs_scalar_type b
Definition: scalar.hpp:38
static void init(viennacl::ocl::context &ctx)
Definition: scalar.hpp:233
Representation of an OpenCL kernel in ViennaCL.
asbs_scalar_type a
Definition: scalar.hpp:37
Helper class for converting a type to its string representation.
Definition: utils.hpp:57
static std::string program_name()
Definition: scalar.hpp:228
void generate_asbs(StringType &source, std::string const &numeric_string)
Definition: scalar.hpp:183