ViennaCL - The Vienna Computing Library  1.5.2
compressed_compressed_matrix.hpp
Go to the documentation of this file.
1 #ifndef VIENNACL_LINALG_OPENCL_KERNELS_COMPRESSED_COMPRESSED_MATRIX_HPP
2 #define VIENNACL_LINALG_OPENCL_KERNELS_COMPRESSED_COMPRESSED_MATRIX_HPP
3 
7 #include "viennacl/ocl/utils.hpp"
8 
11 namespace viennacl
12 {
13  namespace linalg
14  {
15  namespace opencl
16  {
17  namespace kernels
18  {
19 
21 
22  template <typename StringType>
23  void generate_vec_mul(StringType & source, std::string const & numeric_string)
24  {
25  source.append("__kernel void vec_mul( \n");
26  source.append(" __global const unsigned int * row_jumper, \n");
27  source.append(" __global const unsigned int * row_indices, \n");
28  source.append(" __global const unsigned int * column_indices, \n");
29  source.append(" __global const "); source.append(numeric_string); source.append(" * elements, \n");
30  source.append(" uint nonzero_rows, \n");
31  source.append(" __global const "); source.append(numeric_string); source.append(" * x, \n");
32  source.append(" uint4 layout_x, \n");
33  source.append(" __global "); source.append(numeric_string); source.append(" * result, \n");
34  source.append(" uint4 layout_result) \n");
35  source.append("{ \n");
36  source.append(" for (unsigned int i = get_global_id(0); i < nonzero_rows; i += get_global_size(0)) \n");
37  source.append(" { \n");
38  source.append(" "); source.append(numeric_string); source.append(" dot_prod = 0; \n");
39  source.append(" unsigned int row_end = row_jumper[i+1]; \n");
40  source.append(" for (unsigned int j = row_jumper[i]; j < row_end; ++j) \n");
41  source.append(" dot_prod += elements[j] * x[column_indices[j] * layout_x.y + layout_x.x]; \n");
42  source.append(" result[row_indices[i] * layout_result.y + layout_result.x] = dot_prod; \n");
43  source.append(" } \n");
44  source.append(" } \n");
45  }
46 
48 
50  template <typename NumericT>
52  {
53  static std::string program_name()
54  {
55  return viennacl::ocl::type_to_string<NumericT>::apply() + "_compressed_compressed_matrix";
56  }
57 
58  static void init(viennacl::ocl::context & ctx)
59  {
61  std::string numeric_string = viennacl::ocl::type_to_string<NumericT>::apply();
62 
63  static std::map<cl_context, bool> init_done;
64  if (!init_done[ctx.handle().get()])
65  {
66  std::string source;
67  source.reserve(8192);
68 
69  viennacl::ocl::append_double_precision_pragma<NumericT>(ctx, source);
70 
71  // fully parametrized kernels:
72  generate_vec_mul(source, numeric_string);
73 
74  std::string prog_name = program_name();
75  #ifdef VIENNACL_BUILD_INFO
76  std::cout << "Creating program " << prog_name << std::endl;
77  #endif
78  ctx.add_program(source, prog_name);
79  init_done[ctx.handle().get()] = true;
80  } //if
81  } //init
82  };
83 
84  } // namespace kernels
85  } // namespace opencl
86  } // namespace linalg
87 } // namespace viennacl
88 #endif
89 
Implements a OpenCL platform within ViennaCL.
Various little tools used here and there in ViennaCL.
static void init(viennacl::ocl::context &ctx)
Definition: compressed_compressed_matrix.hpp:58
Manages an OpenCL context and provides the respective convenience functions for creating buffers...
Definition: context.hpp:51
Provides OpenCL-related utilities.
const OCL_TYPE & get() const
Definition: handle.hpp:189
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
void generate_vec_mul(StringType &source, std::string const &numeric_string)
Definition: compressed_compressed_matrix.hpp:23
static void apply(viennacl::ocl::context const &)
Definition: utils.hpp:40
Representation of an OpenCL kernel in ViennaCL.
Helper class for converting a type to its string representation.
Definition: utils.hpp:57
static std::string program_name()
Definition: compressed_compressed_matrix.hpp:53
Main kernel class for generating OpenCL kernels for compressed_compressed_matrix. ...
Definition: compressed_compressed_matrix.hpp:51