1 #ifndef VIENNACL_OCL_ENQUEUE_HPP_
2 #define VIENNACL_OCL_ENQUEUE_HPP_
26 #include <OpenCL/cl.h>
39 class custom_operation;
47 template <
typename KernelType>
51 if (k.local_work_size(1) == 0)
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;
63 if (tmp_global == 1 && tmp_local == 1)
64 err = clEnqueueTask(queue.
handle().
get(), k.handle().get(), 0, NULL, NULL);
66 err = clEnqueueNDRangeKernel(queue.
handle().
get(), k.handle().get(), 1, NULL, &tmp_global, &tmp_local, 0, NULL, NULL);
68 if (err != CL_SUCCESS)
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;
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;
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);
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);
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);
95 if (err != CL_SUCCESS)
98 std::cerr <<
"ViennaCL: FATAL ERROR: Kernel start failed for '" << k.name() <<
"'." << std::endl;
103 #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
105 std::cout <<
"ViennaCL: Kernel " << k.name() <<
" finished!" << std::endl;
111 template <
typename KernelType>
114 enqueue(k, k.context().get_queue());
122 inline void enqueue(viennacl::generator::custom_operation & op)
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.