ViennaCL - The Vienna Computing Library  1.5.2
enqueue.hpp
Go to the documentation of this file.
1 #ifndef VIENNACL_OCL_ENQUEUE_HPP_
2 #define VIENNACL_OCL_ENQUEUE_HPP_
3 
4 /* =========================================================================
5  Copyright (c) 2010-2014, Institute for Microelectronics,
6  Institute for Analysis and Scientific Computing,
7  TU Wien.
8  Portions of this software are copyright by UChicago Argonne, LLC.
9 
10  -----------------
11  ViennaCL - The Vienna Computing Library
12  -----------------
13 
14  Project Head: Karl Rupp rupp@iue.tuwien.ac.at
15 
16  (A list of authors and contributors can be found in the PDF manual)
17 
18  License: MIT (X11), see file LICENSE in the base directory
19 ============================================================================= */
20 
25 #ifdef __APPLE__
26 #include <OpenCL/cl.h>
27 #else
28 #include <CL/cl.h>
29 #endif
30 
31 #include "viennacl/ocl/backend.hpp"
32 #include "viennacl/ocl/kernel.hpp"
34 #include "viennacl/ocl/context.hpp"
35 
36 namespace viennacl
37 {
38  namespace generator{
39  class custom_operation;
40  void enqueue_custom_op(viennacl::generator::custom_operation & op, viennacl::ocl::command_queue const & queue);
41  }
42 
43  namespace ocl
44  {
45 
47  template <typename KernelType>
48  void enqueue(KernelType & k, viennacl::ocl::command_queue const & queue)
49  {
50  // 1D kernel:
51  if (k.local_work_size(1) == 0)
52  {
53  #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
54  std::cout << "ViennaCL: Starting 1D-kernel '" << k.name() << "'..." << std::endl;
55  std::cout << "ViennaCL: Global work size: '" << k.global_work_size() << "'..." << std::endl;
56  std::cout << "ViennaCL: Local work size: '" << k.local_work_size() << "'..." << std::endl;
57  #endif
58 
59  vcl_size_t tmp_global = k.global_work_size();
60  vcl_size_t tmp_local = k.local_work_size();
61 
62  cl_int err;
63  if (tmp_global == 1 && tmp_local == 1)
64  err = clEnqueueTask(queue.handle().get(), k.handle().get(), 0, NULL, NULL);
65  else
66  err = clEnqueueNDRangeKernel(queue.handle().get(), k.handle().get(), 1, NULL, &tmp_global, &tmp_local, 0, NULL, NULL);
67 
68  if (err != CL_SUCCESS)
69  {
70  std::cerr << "ViennaCL: FATAL ERROR: Kernel start failed for '" << k.name() << "'." << std::endl;
71  std::cerr << "ViennaCL: Smaller work sizes could not solve the problem. " << std::endl;
72  VIENNACL_ERR_CHECK(err);
73  }
74  }
75  else //2D or 3D kernel
76  {
77  #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
78  std::cout << "ViennaCL: Starting 2D/3D-kernel '" << k.name() << "'..." << std::endl;
79  std::cout << "ViennaCL: Global work size: '" << k.global_work_size(0) << ", " << k.global_work_size(1) << ", " << k.global_work_size(2) << "'..." << std::endl;
80  std::cout << "ViennaCL: Local work size: '" << k.local_work_size(0) << ", " << k.local_work_size(1) << ", " << k.local_work_size(2) << "'..." << std::endl;
81  #endif
82 
83  vcl_size_t tmp_global[3];
84  tmp_global[0] = k.global_work_size(0);
85  tmp_global[1] = k.global_work_size(1);
86  tmp_global[2] = k.global_work_size(2);
87 
88  vcl_size_t tmp_local[3];
89  tmp_local[0] = k.local_work_size(0);
90  tmp_local[1] = k.local_work_size(1);
91  tmp_local[2] = k.local_work_size(2);
92 
93  cl_int err = clEnqueueNDRangeKernel(queue.handle().get(), k.handle().get(), (tmp_global[2] == 0) ? 2 : 3, NULL, tmp_global, tmp_local, 0, NULL, NULL);
94 
95  if (err != CL_SUCCESS)
96  {
97  //could not start kernel with any parameters
98  std::cerr << "ViennaCL: FATAL ERROR: Kernel start failed for '" << k.name() << "'." << std::endl;
99  VIENNACL_ERR_CHECK(err);
100  }
101  }
102 
103  #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
104  queue.finish();
105  std::cout << "ViennaCL: Kernel " << k.name() << " finished!" << std::endl;
106  #endif
107  } //enqueue()
108 
109 
111  template <typename KernelType>
112  void enqueue(KernelType & k)
113  {
114  enqueue(k, k.context().get_queue());
115  }
116 
117  inline void enqueue(viennacl::generator::custom_operation & op, viennacl::ocl::command_queue const & queue)
118  {
120  }
121 
122  inline void enqueue(viennacl::generator::custom_operation & op)
123  {
125  }
126 
127  } // namespace ocl
128 } // namespace viennacl
129 #endif
std::size_t vcl_size_t
Definition: forwards.h:58
void enqueue_custom_op(viennacl::generator::custom_operation &op, viennacl::ocl::command_queue const &queue)
A class representing a command queue.
Definition: command_queue.hpp:45
viennacl::ocl::handle< cl_command_queue > const & handle() const
Definition: command_queue.hpp:81
void finish() const
Waits until all kernels in the queue have finished their execution.
Definition: command_queue.hpp:70
Implementations of command queue representations.
#define VIENNACL_ERR_CHECK(err)
Definition: error.hpp:655
void enqueue(KernelType &k, viennacl::ocl::command_queue const &queue)
Enqueues a kernel in the provided queue.
Definition: enqueue.hpp:48
const OCL_TYPE & get() const
Definition: handle.hpp:189
viennacl::ocl::command_queue & get_queue()
Convenience function for getting the default queue for the currently active device in the active cont...
Definition: backend.hpp:299
viennacl::ocl::context & current_context()
Convenience function for returning the current context.
Definition: backend.hpp:192
Represents an OpenCL context within ViennaCL.
Implementations of the OpenCL backend, where all contexts are stored in.
Representation of an OpenCL kernel in ViennaCL.