1 #ifndef VIENNACL_OCL_ENQUEUE_HPP_
2 #define VIENNACL_OCL_ENQUEUE_HPP_
25 #include <OpenCL/cl.h>
39 template <
typename KernelType>
43 if (k.local_work_size(1) == 0)
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;
51 size_t tmp_global = k.global_work_size();
52 size_t tmp_local = k.local_work_size();
55 if (tmp_global == 1 && tmp_local == 1)
56 err = clEnqueueTask(queue.
handle(), k.
handle(), 0, NULL, NULL);
58 err = clEnqueueNDRangeKernel(queue.
handle(), k.
handle(), 1, NULL, &tmp_global, &tmp_local, 0, NULL, NULL);
60 if (err != CL_SUCCESS)
63 while (err != CL_SUCCESS && tmp_local > 1)
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;
78 err = clEnqueueNDRangeKernel(queue.
handle(), k.
handle(), 1, NULL, &tmp_global, &tmp_local, 0, NULL, NULL);
81 if (err != CL_SUCCESS)
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;
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;
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;
107 size_t tmp_global[2];
108 tmp_global[0] = k.global_work_size(0);
109 tmp_global[1] = k.global_work_size(1);
112 tmp_local[0] = k.local_work_size(0);
113 tmp_local[1] = k.local_work_size(1);
115 cl_int err = clEnqueueNDRangeKernel(queue.
handle(), k.
handle(), 2, NULL, tmp_global, tmp_local, 0, NULL, NULL);
117 if (err != CL_SUCCESS)
120 std::cerr <<
"ViennaCL: FATAL ERROR: Kernel start failed for '" << k.name() <<
"'." << std::endl;
126 #if defined(VIENNACL_DEBUG_ALL) || defined(VIENNACL_DEBUG_KERNEL)
128 std::cout <<
"ViennaCL: Kernel " << k.name() <<
" finished!" << std::endl;
134 template <
typename KernelType>