ViennaCL - The Vienna Computing Library  1.5.2
ilu.hpp
Go to the documentation of this file.
1 #ifndef VIENNACL_LINALG_OPENCL_KERNELS_ILU_HPP
2 #define VIENNACL_LINALG_OPENCL_KERNELS_ILU_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  template <typename StringType>
20  void generate_ilu_level_scheduling_substitute(StringType & source, std::string const & numeric_string)
21  {
22  source.append("__kernel void level_scheduling_substitute( \n");
23  source.append(" __global const unsigned int * row_index_array, \n");
24  source.append(" __global const unsigned int * row_indices, \n");
25  source.append(" __global const unsigned int * column_indices, \n");
26  source.append(" __global const "); source.append(numeric_string); source.append(" * elements, \n");
27  source.append(" __global "); source.append(numeric_string); source.append(" * vec, \n");
28  source.append(" unsigned int size) \n");
29  source.append("{ \n");
30  source.append(" for (unsigned int row = get_global_id(0); \n");
31  source.append(" row < size; \n");
32  source.append(" row += get_global_size(0)) \n");
33  source.append(" { \n");
34  source.append(" unsigned int eq_row = row_index_array[row]; \n");
35  source.append(" "); source.append(numeric_string); source.append(" vec_entry = vec[eq_row]; \n");
36  source.append(" unsigned int row_end = row_indices[row+1]; \n");
37 
38  source.append(" for (unsigned int j = row_indices[row]; j < row_end; ++j) \n");
39  source.append(" vec_entry -= vec[column_indices[j]] * elements[j]; \n");
40 
41  source.append(" vec[eq_row] = vec_entry; \n");
42  source.append(" } \n");
43  source.append("} \n");
44  }
45 
46  // main kernel class
48  template <class NumericT>
49  struct ilu
50  {
51  static std::string program_name()
52  {
54  }
55 
56  static void init(viennacl::ocl::context & ctx)
57  {
59  std::string numeric_string = viennacl::ocl::type_to_string<NumericT>::apply();
60 
61  static std::map<cl_context, bool> init_done;
62  if (!init_done[ctx.handle().get()])
63  {
64  std::string source;
65  source.reserve(1024);
66 
67  viennacl::ocl::append_double_precision_pragma<NumericT>(ctx, source);
68 
69  // only generate for floating points (forces error for integers)
70  if (numeric_string == "float" || numeric_string == "double")
71  {
72  generate_ilu_level_scheduling_substitute(source, numeric_string);
73  }
74 
75  std::string prog_name = program_name();
76  #ifdef VIENNACL_BUILD_INFO
77  std::cout << "Creating program " << prog_name << std::endl;
78  #endif
79  ctx.add_program(source, prog_name);
80  init_done[ctx.handle().get()] = true;
81  } //if
82  } //init
83  };
84 
85  } // namespace kernels
86  } // namespace opencl
87  } // namespace linalg
88 } // namespace viennacl
89 #endif
90 
Implements a OpenCL platform within ViennaCL.
Various little tools used here and there in ViennaCL.
Manages an OpenCL context and provides the respective convenience functions for creating buffers...
Definition: context.hpp:51
Provides OpenCL-related utilities.
void generate_ilu_level_scheduling_substitute(StringType &source, std::string const &numeric_string)
Definition: ilu.hpp:20
const viennacl::ocl::handle< cl_context > & handle() const
Returns the context handle.
Definition: context.hpp:476
Main kernel class for generating OpenCL kernels for incomplete LU factorization preconditioners.
Definition: ilu.hpp:49
const OCL_TYPE & get() const
Definition: handle.hpp:189
static void init(viennacl::ocl::context &ctx)
Definition: ilu.hpp:56
static void apply(viennacl::ocl::context const &)
Definition: utils.hpp:40
Representation of an OpenCL kernel in ViennaCL.
static std::string program_name()
Definition: ilu.hpp:51
Helper class for converting a type to its string representation.
Definition: utils.hpp:57