ViennaCL - The Vienna Computing Library  1.7.0
Free open-source GPU-accelerated linear algebra and solver library.
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-2015, 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 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 
39 namespace device_specific
40 {
41  class custom_operation;
42  void enqueue_custom_op(viennacl::device_specific::custom_operation & op, viennacl::ocl::command_queue const & queue);
43 }
44 
45 namespace ocl
46 {
47 
49 template<typename KernelType>
50 void enqueue(KernelType & k, viennacl::ocl::command_queue const & queue)
51 {
52 #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
53  cl_event event;
54 #endif
55 
56  // 1D kernel:
57  if (k.local_work_size(1) == 0)
58  {
59 #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
60  std::cout << "ViennaCL: Starting 1D-kernel '" << k.name() << "'..." << std::endl;
61  std::cout << "ViennaCL: Global work size: '" << k.global_work_size() << "'..." << std::endl;
62  std::cout << "ViennaCL: Local work size: '" << k.local_work_size() << "'..." << std::endl;
63 #endif
64 
65  vcl_size_t tmp_global = k.global_work_size();
66  vcl_size_t tmp_local = k.local_work_size();
67 
68  cl_int err;
69  if (tmp_global == 1 && tmp_local == 1)
70 #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
71  err = clEnqueueTask(queue.handle().get(), k.handle().get(), 0, NULL, &event);
72 #else
73  err = clEnqueueTask(queue.handle().get(), k.handle().get(), 0, NULL, NULL);
74 #endif
75  else
76 #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
77  err = clEnqueueNDRangeKernel(queue.handle().get(), k.handle().get(), 1, NULL, &tmp_global, &tmp_local, 0, NULL, &event);
78 #else
79  err = clEnqueueNDRangeKernel(queue.handle().get(), k.handle().get(), 1, NULL, &tmp_global, &tmp_local, 0, NULL, NULL);
80 #endif
81 
82  if (err != CL_SUCCESS)
83  {
84  std::cerr << "ViennaCL: FATAL ERROR: Kernel start failed for '" << k.name() << "'." << std::endl;
85  std::cerr << "ViennaCL: Smaller work sizes could not solve the problem. " << std::endl;
86  VIENNACL_ERR_CHECK(err);
87  }
88  }
89  else //2D or 3D kernel
90  {
91 #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
92  std::cout << "ViennaCL: Starting 2D/3D-kernel '" << k.name() << "'..." << std::endl;
93  std::cout << "ViennaCL: Global work size: '" << k.global_work_size(0) << ", " << k.global_work_size(1) << ", " << k.global_work_size(2) << "'..." << std::endl;
94  std::cout << "ViennaCL: Local work size: '" << k.local_work_size(0) << ", " << k.local_work_size(1) << ", " << k.local_work_size(2) << "'..." << std::endl;
95 #endif
96 
97  vcl_size_t tmp_global[3];
98  tmp_global[0] = k.global_work_size(0);
99  tmp_global[1] = k.global_work_size(1);
100  tmp_global[2] = k.global_work_size(2);
101 
102  vcl_size_t tmp_local[3];
103  tmp_local[0] = k.local_work_size(0);
104  tmp_local[1] = k.local_work_size(1);
105  tmp_local[2] = k.local_work_size(2);
106 
107 #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
108  cl_int err = clEnqueueNDRangeKernel(queue.handle().get(), k.handle().get(), (tmp_global[2] == 0) ? 2 : 3, NULL, tmp_global, tmp_local, 0, NULL, &event);
109 #else
110  cl_int err = clEnqueueNDRangeKernel(queue.handle().get(), k.handle().get(), (tmp_global[2] == 0) ? 2 : 3, NULL, tmp_global, tmp_local, 0, NULL, NULL);
111 #endif
112  if (err != CL_SUCCESS)
113  {
114  //could not start kernel with any parameters
115  std::cerr << "ViennaCL: FATAL ERROR: Kernel start failed for '" << k.name() << "'." << std::endl;
116  VIENNACL_ERR_CHECK(err);
117  }
118  }
119 
120 #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
121  queue.finish();
122  cl_int execution_status;
123  clGetEventInfo(event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int), &execution_status, NULL);
124  std::cout << "ViennaCL: Kernel " << k.name() << " finished with status " << execution_status << "!" << std::endl;
125 #endif
126 } //enqueue()
127 
128 
130 template<typename KernelType>
131 void enqueue(KernelType & k)
132 {
133  enqueue(k, k.context().get_queue());
134 }
135 
136 inline void enqueue(viennacl::device_specific::custom_operation & op, viennacl::ocl::command_queue const & queue)
137 {
139 }
140 
141 inline void enqueue(viennacl::device_specific::custom_operation & op)
142 {
144 }
145 
146 } // namespace ocl
147 } // namespace viennacl
148 #endif
void finish() const
Waits until all kernels in the queue have finished their execution.
viennacl::ocl::context & current_context()
Convenience function for returning the current context.
Definition: backend.hpp:213
A class representing a command queue.
viennacl::ocl::handle< cl_command_queue > const & handle() const
Implementations of command queue representations.
#define VIENNACL_ERR_CHECK(err)
Definition: error.hpp:681
Main namespace in ViennaCL. Holds all the basic types such as vector, matrix, etc. and defines operations upon them.
Definition: cpu_ram.hpp:34
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:320
void enqueue_custom_op(viennacl::device_specific::custom_operation &op, viennacl::ocl::command_queue const &queue)
std::size_t vcl_size_t
Definition: forwards.h:75
Represents an OpenCL context within ViennaCL.
Implementations of the OpenCL backend, where all contexts are stored in.
void enqueue(KernelType &k, viennacl::ocl::command_queue const &queue)
Enqueues a kernel in the provided queue.
Definition: enqueue.hpp:50
Representation of an OpenCL kernel in ViennaCL.