ViennaCL - The Vienna Computing Library  1.2.0
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-2011, Institute for Microelectronics,
6  Institute for Analysis and Scientific Computing,
7  TU Wien.
8 
9  -----------------
10  ViennaCL - The Vienna Computing Library
11  -----------------
12 
13  Project Head: Karl Rupp rupp@iue.tuwien.ac.at
14 
15  (A list of authors and contributors can be found in the PDF manual)
16 
17  License: MIT (X11), see file LICENSE in the base directory
18 ============================================================================= */
19 
24 #ifdef __APPLE__
25 #include <OpenCL/cl.h>
26 #else
27 #include <CL/cl.h>
28 #endif
29 
30 #include "viennacl/ocl/kernel.hpp"
32 
33 namespace viennacl
34 {
35  namespace ocl
36  {
37 
39  template <typename KernelType>
40  void enqueue(KernelType & k, viennacl::ocl::command_queue const & queue)
41  {
42  // 1D kernel:
43  if (k.local_work_size(1) == 0)
44  {
45  #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
46  std::cout << "ViennaCL: Starting 1D-kernel '" << k.name() << "'..." << std::endl;
47  std::cout << "ViennaCL: Global work size: '" << k.global_work_size() << "'..." << std::endl;
48  std::cout << "ViennaCL: Local work size: '" << k.local_work_size() << "'..." << std::endl;
49  #endif
50 
51  size_t tmp_global = k.global_work_size();
52  size_t tmp_local = k.local_work_size();
53 
54  cl_int err;
55  if (tmp_global == 1 && tmp_local == 1)
56  err = clEnqueueTask(queue.handle(), k.handle(), 0, NULL, NULL);
57  else
58  err = clEnqueueNDRangeKernel(queue.handle(), k.handle(), 1, NULL, &tmp_global, &tmp_local, 0, NULL, NULL);
59 
60  if (err != CL_SUCCESS) //if not successful, try to start with smaller work size
61  {
62  //std::cout << "FAIL: " << std::endl; exit(0);
63  while (err != CL_SUCCESS && tmp_local > 1)
64  {
65  //std::cout << "Flushing queue, then enqueuing again with half the size..." << std::endl;
66  //std::cout << "Error code: " << err << std::endl;
67 
68  tmp_global /= 2;
69  tmp_local /= 2;
70 
71  #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
72  std::cout << "ViennaCL: Kernel start failed for '" << k.name() << "'." << std::endl;
73  std::cout << "ViennaCL: Global work size: '" << tmp_global << "'..." << std::endl;
74  std::cout << "ViennaCL: Local work size: '" << tmp_local << "'..." << std::endl;
75  #endif
76 
77  queue.finish();
78  err = clEnqueueNDRangeKernel(queue.handle(), k.handle(), 1, NULL, &tmp_global, &tmp_local, 0, NULL, NULL);
79  }
80 
81  if (err != CL_SUCCESS)
82  {
83  //could not start kernel with any parameters
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  else
89  {
90  //remember parameters:
91  k.local_work_size(0, tmp_local);
92  k.global_work_size(0, tmp_global);
93  #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
94  std::cout << "ViennaCL: Kernel '" << k.name() << "' now uses global work size " << tmp_global << " and local work size " << tmp_local << "." << std::endl;
95  #endif
96  }
97  }
98  }
99  else //2D kernel
100  {
101  #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
102  std::cout << "ViennaCL: Starting 2D-kernel '" << k.name() << "'..." << std::endl;
103  std::cout << "ViennaCL: Global work size: '" << k.global_work_size(0) << ", " << k.global_work_size(1) << "'..." << std::endl;
104  std::cout << "ViennaCL: Local work size: '" << k.local_work_size(0) << ", " << k.local_work_size(1) << "'..." << std::endl;
105  #endif
106 
107  size_t tmp_global[2];
108  tmp_global[0] = k.global_work_size(0);
109  tmp_global[1] = k.global_work_size(1);
110 
111  size_t tmp_local[2];
112  tmp_local[0] = k.local_work_size(0);
113  tmp_local[1] = k.local_work_size(1);
114 
115  cl_int err = clEnqueueNDRangeKernel(queue.handle(), k.handle(), 2, NULL, tmp_global, tmp_local, 0, NULL, NULL);
116 
117  if (err != CL_SUCCESS)
118  {
119  //could not start kernel with any parameters
120  std::cerr << "ViennaCL: FATAL ERROR: Kernel start failed for '" << k.name() << "'." << std::endl;
121  VIENNACL_ERR_CHECK(err);
122  }
123 
124  }
125 
126  #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
127  queue.finish();
128  std::cout << "ViennaCL: Kernel " << k.name() << " finished!" << std::endl;
129  #endif
130  } //enqueue()
131 
132 
134  template <typename KernelType>
135  void enqueue(KernelType & k)
136  {
138  }
139  } // namespace ocl
140 } // namespace viennacl
141 #endif