1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
4 // Distributed under the Boost Software License, Version 1.0
5 // See accompanying file LICENSE_1_0.txt or copy at
6 // http://www.boost.org/LICENSE_1_0.txt
8 // See http://boostorg.github.com/compute for more information.
9 //---------------------------------------------------------------------------//
11 #ifndef BOOST_COMPUTE_COMMAND_QUEUE_HPP
12 #define BOOST_COMPUTE_COMMAND_QUEUE_HPP
17 #include <boost/assert.hpp>
19 #include <boost/compute/config.hpp>
20 #include <boost/compute/event.hpp>
21 #include <boost/compute/buffer.hpp>
22 #include <boost/compute/device.hpp>
23 #include <boost/compute/kernel.hpp>
24 #include <boost/compute/context.hpp>
25 #include <boost/compute/exception.hpp>
26 #include <boost/compute/image/image1d.hpp>
27 #include <boost/compute/image/image2d.hpp>
28 #include <boost/compute/image/image3d.hpp>
29 #include <boost/compute/image/image_object.hpp>
30 #include <boost/compute/utility/wait_list.hpp>
31 #include <boost/compute/detail/get_object_info.hpp>
32 #include <boost/compute/detail/assert_cl_success.hpp>
33 #include <boost/compute/detail/diagnostic.hpp>
34 #include <boost/compute/utility/extents.hpp>
40 inline void BOOST_COMPUTE_CL_CALLBACK
41 nullary_native_kernel_trampoline(void *user_func_ptr)
44 std::memcpy(&user_func, user_func_ptr, sizeof(user_func));
48 } // end detail namespace
50 /// \class command_queue
51 /// \brief A command queue.
53 /// Command queues provide the interface for interacting with compute
54 /// devices. The command_queue class provides methods to copy data to
55 /// and from a compute device as well as execute compute kernels.
57 /// Command queues are created for a compute device within a compute
60 /// For example, to create a context and command queue for the default device
61 /// on the system (this is the normal set up code used by almost all OpenCL
64 /// #include <boost/compute/core.hpp>
66 /// // get the default compute device
67 /// boost::compute::device device = boost::compute::system::default_device();
69 /// // set up a compute context and command queue
70 /// boost::compute::context context(device);
71 /// boost::compute::command_queue queue(context, device);
74 /// The default command queue for the system can be obtained with the
75 /// system::default_queue() method.
77 /// \see buffer, context, kernel
82 enable_profiling = CL_QUEUE_PROFILING_ENABLE,
83 enable_out_of_order_execution = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
87 map_read = CL_MAP_READ,
88 map_write = CL_MAP_WRITE
91 map_write_invalidate_region = CL_MAP_WRITE_INVALIDATE_REGION
95 /// Creates a null command queue.
101 explicit command_queue(cl_command_queue queue, bool retain = true)
104 if(m_queue && retain){
105 clRetainCommandQueue(m_queue);
109 /// Creates a command queue in \p context for \p device with
112 /// \see_opencl_ref{clCreateCommandQueue}
113 command_queue(const context &context,
114 const device &device,
115 cl_command_queue_properties properties = 0)
117 BOOST_ASSERT(device.id() != 0);
121 #ifdef CL_VERSION_2_0
122 if (device.check_version(2, 0)){
123 std::vector<cl_queue_properties> queue_properties;
125 queue_properties.push_back(CL_QUEUE_PROPERTIES);
126 queue_properties.push_back(cl_queue_properties(properties));
127 queue_properties.push_back(cl_queue_properties(0));
130 const cl_queue_properties *queue_properties_ptr =
131 queue_properties.empty() ? 0 : &queue_properties[0];
133 m_queue = clCreateCommandQueueWithProperties(
134 context, device.id(), queue_properties_ptr, &error
139 // Suppress deprecated declarations warning
140 BOOST_COMPUTE_DISABLE_DEPRECATED_DECLARATIONS();
141 m_queue = clCreateCommandQueue(
142 context, device.id(), properties, &error
144 BOOST_COMPUTE_ENABLE_DEPRECATED_DECLARATIONS();
148 BOOST_THROW_EXCEPTION(opencl_error(error));
152 /// Creates a new command queue object as a copy of \p other.
153 command_queue(const command_queue &other)
154 : m_queue(other.m_queue)
157 clRetainCommandQueue(m_queue);
161 /// Copies the command queue object from \p other to \c *this.
162 command_queue& operator=(const command_queue &other)
166 clReleaseCommandQueue(m_queue);
169 m_queue = other.m_queue;
172 clRetainCommandQueue(m_queue);
179 #ifndef BOOST_COMPUTE_NO_RVALUE_REFERENCES
180 /// Move-constructs a new command queue object from \p other.
181 command_queue(command_queue&& other) BOOST_NOEXCEPT
182 : m_queue(other.m_queue)
187 /// Move-assigns the command queue from \p other to \c *this.
188 command_queue& operator=(command_queue&& other) BOOST_NOEXCEPT
191 clReleaseCommandQueue(m_queue);
194 m_queue = other.m_queue;
199 #endif // BOOST_COMPUTE_NO_RVALUE_REFERENCES
201 /// Destroys the command queue.
203 /// \see_opencl_ref{clReleaseCommandQueue}
207 BOOST_COMPUTE_ASSERT_CL_SUCCESS(
208 clReleaseCommandQueue(m_queue)
213 /// Returns the underlying OpenCL command queue.
214 cl_command_queue& get() const
216 return const_cast<cl_command_queue &>(m_queue);
219 /// Returns the device that the command queue issues commands to.
220 device get_device() const
222 return device(get_info<cl_device_id>(CL_QUEUE_DEVICE));
225 /// Returns the context for the command queue.
226 context get_context() const
228 return context(get_info<cl_context>(CL_QUEUE_CONTEXT));
231 /// Returns information about the command queue.
233 /// \see_opencl_ref{clGetCommandQueueInfo}
235 T get_info(cl_command_queue_info info) const
237 return detail::get_object_info<T>(clGetCommandQueueInfo, m_queue, info);
242 typename detail::get_object_info_type<command_queue, Enum>::type
245 /// Returns the properties for the command queue.
246 cl_command_queue_properties get_properties() const
248 return get_info<cl_command_queue_properties>(CL_QUEUE_PROPERTIES);
251 /// Enqueues a command to read data from \p buffer to host memory.
253 /// \see_opencl_ref{clEnqueueReadBuffer}
256 event enqueue_read_buffer(const buffer &buffer,
260 const wait_list &events = wait_list())
262 BOOST_ASSERT(m_queue != 0);
263 BOOST_ASSERT(size <= buffer.size());
264 BOOST_ASSERT(buffer.get_context() == this->get_context());
265 BOOST_ASSERT(host_ptr != 0);
269 cl_int ret = clEnqueueReadBuffer(
277 events.get_event_ptr(),
281 if(ret != CL_SUCCESS){
282 BOOST_THROW_EXCEPTION(opencl_error(ret));
288 /// Enqueues a command to read data from \p buffer to host memory. The
289 /// copy will be performed asynchronously.
291 /// \see_opencl_ref{clEnqueueReadBuffer}
293 /// \see copy_async()
294 event enqueue_read_buffer_async(const buffer &buffer,
298 const wait_list &events = wait_list())
300 BOOST_ASSERT(m_queue != 0);
301 BOOST_ASSERT(size <= buffer.size());
302 BOOST_ASSERT(buffer.get_context() == this->get_context());
303 BOOST_ASSERT(host_ptr != 0);
307 cl_int ret = clEnqueueReadBuffer(
315 events.get_event_ptr(),
319 if(ret != CL_SUCCESS){
320 BOOST_THROW_EXCEPTION(opencl_error(ret));
326 #if defined(CL_VERSION_1_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
327 /// Enqueues a command to read a rectangular region from \p buffer to
330 /// \see_opencl_ref{clEnqueueReadBufferRect}
332 /// \opencl_version_warning{1,1}
333 event enqueue_read_buffer_rect(const buffer &buffer,
334 const size_t buffer_origin[3],
335 const size_t host_origin[3],
336 const size_t region[3],
337 size_t buffer_row_pitch,
338 size_t buffer_slice_pitch,
339 size_t host_row_pitch,
340 size_t host_slice_pitch,
342 const wait_list &events = wait_list())
344 BOOST_ASSERT(m_queue != 0);
345 BOOST_ASSERT(buffer.get_context() == this->get_context());
346 BOOST_ASSERT(host_ptr != 0);
350 cl_int ret = clEnqueueReadBufferRect(
363 events.get_event_ptr(),
367 if(ret != CL_SUCCESS){
368 BOOST_THROW_EXCEPTION(opencl_error(ret));
374 /// Enqueues a command to read a rectangular region from \p buffer to
375 /// host memory. The copy will be performed asynchronously.
377 /// \see_opencl_ref{clEnqueueReadBufferRect}
379 /// \opencl_version_warning{1,1}
380 event enqueue_read_buffer_rect_async(const buffer &buffer,
381 const size_t buffer_origin[3],
382 const size_t host_origin[3],
383 const size_t region[3],
384 size_t buffer_row_pitch,
385 size_t buffer_slice_pitch,
386 size_t host_row_pitch,
387 size_t host_slice_pitch,
389 const wait_list &events = wait_list())
391 BOOST_ASSERT(m_queue != 0);
392 BOOST_ASSERT(buffer.get_context() == this->get_context());
393 BOOST_ASSERT(host_ptr != 0);
397 cl_int ret = clEnqueueReadBufferRect(
410 events.get_event_ptr(),
414 if(ret != CL_SUCCESS){
415 BOOST_THROW_EXCEPTION(opencl_error(ret));
420 #endif // CL_VERSION_1_1
422 /// Enqueues a command to write data from host memory to \p buffer.
424 /// \see_opencl_ref{clEnqueueWriteBuffer}
427 event enqueue_write_buffer(const buffer &buffer,
430 const void *host_ptr,
431 const wait_list &events = wait_list())
433 BOOST_ASSERT(m_queue != 0);
434 BOOST_ASSERT(size <= buffer.size());
435 BOOST_ASSERT(buffer.get_context() == this->get_context());
436 BOOST_ASSERT(host_ptr != 0);
440 cl_int ret = clEnqueueWriteBuffer(
448 events.get_event_ptr(),
452 if(ret != CL_SUCCESS){
453 BOOST_THROW_EXCEPTION(opencl_error(ret));
459 /// Enqueues a command to write data from host memory to \p buffer.
460 /// The copy is performed asynchronously.
462 /// \see_opencl_ref{clEnqueueWriteBuffer}
464 /// \see copy_async()
465 event enqueue_write_buffer_async(const buffer &buffer,
468 const void *host_ptr,
469 const wait_list &events = wait_list())
471 BOOST_ASSERT(m_queue != 0);
472 BOOST_ASSERT(size <= buffer.size());
473 BOOST_ASSERT(buffer.get_context() == this->get_context());
474 BOOST_ASSERT(host_ptr != 0);
478 cl_int ret = clEnqueueWriteBuffer(
486 events.get_event_ptr(),
490 if(ret != CL_SUCCESS){
491 BOOST_THROW_EXCEPTION(opencl_error(ret));
497 #if defined(CL_VERSION_1_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
498 /// Enqueues a command to write a rectangular region from host memory
501 /// \see_opencl_ref{clEnqueueWriteBufferRect}
503 /// \opencl_version_warning{1,1}
504 event enqueue_write_buffer_rect(const buffer &buffer,
505 const size_t buffer_origin[3],
506 const size_t host_origin[3],
507 const size_t region[3],
508 size_t buffer_row_pitch,
509 size_t buffer_slice_pitch,
510 size_t host_row_pitch,
511 size_t host_slice_pitch,
513 const wait_list &events = wait_list())
515 BOOST_ASSERT(m_queue != 0);
516 BOOST_ASSERT(buffer.get_context() == this->get_context());
517 BOOST_ASSERT(host_ptr != 0);
521 cl_int ret = clEnqueueWriteBufferRect(
534 events.get_event_ptr(),
538 if(ret != CL_SUCCESS){
539 BOOST_THROW_EXCEPTION(opencl_error(ret));
545 /// Enqueues a command to write a rectangular region from host memory
546 /// to \p buffer. The copy is performed asynchronously.
548 /// \see_opencl_ref{clEnqueueWriteBufferRect}
550 /// \opencl_version_warning{1,1}
551 event enqueue_write_buffer_rect_async(const buffer &buffer,
552 const size_t buffer_origin[3],
553 const size_t host_origin[3],
554 const size_t region[3],
555 size_t buffer_row_pitch,
556 size_t buffer_slice_pitch,
557 size_t host_row_pitch,
558 size_t host_slice_pitch,
560 const wait_list &events = wait_list())
562 BOOST_ASSERT(m_queue != 0);
563 BOOST_ASSERT(buffer.get_context() == this->get_context());
564 BOOST_ASSERT(host_ptr != 0);
568 cl_int ret = clEnqueueWriteBufferRect(
581 events.get_event_ptr(),
585 if(ret != CL_SUCCESS){
586 BOOST_THROW_EXCEPTION(opencl_error(ret));
591 #endif // CL_VERSION_1_1
593 /// Enqueues a command to copy data from \p src_buffer to
596 /// \see_opencl_ref{clEnqueueCopyBuffer}
599 event enqueue_copy_buffer(const buffer &src_buffer,
600 const buffer &dst_buffer,
604 const wait_list &events = wait_list())
606 BOOST_ASSERT(m_queue != 0);
607 BOOST_ASSERT(src_offset + size <= src_buffer.size());
608 BOOST_ASSERT(dst_offset + size <= dst_buffer.size());
609 BOOST_ASSERT(src_buffer.get_context() == this->get_context());
610 BOOST_ASSERT(dst_buffer.get_context() == this->get_context());
614 cl_int ret = clEnqueueCopyBuffer(
622 events.get_event_ptr(),
626 if(ret != CL_SUCCESS){
627 BOOST_THROW_EXCEPTION(opencl_error(ret));
633 #if defined(CL_VERSION_1_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
634 /// Enqueues a command to copy a rectangular region from
635 /// \p src_buffer to \p dst_buffer.
637 /// \see_opencl_ref{clEnqueueCopyBufferRect}
639 /// \opencl_version_warning{1,1}
640 event enqueue_copy_buffer_rect(const buffer &src_buffer,
641 const buffer &dst_buffer,
642 const size_t src_origin[3],
643 const size_t dst_origin[3],
644 const size_t region[3],
645 size_t buffer_row_pitch,
646 size_t buffer_slice_pitch,
647 size_t host_row_pitch,
648 size_t host_slice_pitch,
649 const wait_list &events = wait_list())
651 BOOST_ASSERT(m_queue != 0);
652 BOOST_ASSERT(src_buffer.get_context() == this->get_context());
653 BOOST_ASSERT(dst_buffer.get_context() == this->get_context());
657 cl_int ret = clEnqueueCopyBufferRect(
669 events.get_event_ptr(),
673 if(ret != CL_SUCCESS){
674 BOOST_THROW_EXCEPTION(opencl_error(ret));
679 #endif // CL_VERSION_1_1
681 #if defined(CL_VERSION_1_2) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
682 /// Enqueues a command to fill \p buffer with \p pattern.
684 /// \see_opencl_ref{clEnqueueFillBuffer}
686 /// \opencl_version_warning{1,2}
689 event enqueue_fill_buffer(const buffer &buffer,
694 const wait_list &events = wait_list())
696 BOOST_ASSERT(m_queue != 0);
697 BOOST_ASSERT(offset + size <= buffer.size());
698 BOOST_ASSERT(buffer.get_context() == this->get_context());
702 cl_int ret = clEnqueueFillBuffer(
710 events.get_event_ptr(),
714 if(ret != CL_SUCCESS){
715 BOOST_THROW_EXCEPTION(opencl_error(ret));
720 #endif // CL_VERSION_1_2
722 /// Enqueues a command to map \p buffer into the host address space.
723 /// Event associated with map operation is returned through
724 /// \p map_buffer_event parameter.
726 /// \see_opencl_ref{clEnqueueMapBuffer}
727 void* enqueue_map_buffer(const buffer &buffer,
731 event &map_buffer_event,
732 const wait_list &events = wait_list())
734 BOOST_ASSERT(m_queue != 0);
735 BOOST_ASSERT(offset + size <= buffer.size());
736 BOOST_ASSERT(buffer.get_context() == this->get_context());
739 void *pointer = clEnqueueMapBuffer(
747 events.get_event_ptr(),
748 &map_buffer_event.get(),
752 if(ret != CL_SUCCESS){
753 BOOST_THROW_EXCEPTION(opencl_error(ret));
760 void* enqueue_map_buffer(const buffer &buffer,
764 const wait_list &events = wait_list())
767 return enqueue_map_buffer(buffer, flags, offset, size, event_, events);
770 /// Enqueues a command to map \p buffer into the host address space.
771 /// Map operation is performed asynchronously. The pointer to the mapped
772 /// region cannot be used until the map operation has completed.
774 /// Event associated with map operation is returned through
775 /// \p map_buffer_event parameter.
777 /// \see_opencl_ref{clEnqueueMapBuffer}
778 void* enqueue_map_buffer_async(const buffer &buffer,
782 event &map_buffer_event,
783 const wait_list &events = wait_list())
785 BOOST_ASSERT(m_queue != 0);
786 BOOST_ASSERT(offset + size <= buffer.size());
787 BOOST_ASSERT(buffer.get_context() == this->get_context());
790 void *pointer = clEnqueueMapBuffer(
798 events.get_event_ptr(),
799 &map_buffer_event.get(),
803 if(ret != CL_SUCCESS){
804 BOOST_THROW_EXCEPTION(opencl_error(ret));
810 /// Enqueues a command to unmap \p buffer from the host memory space.
812 /// \see_opencl_ref{clEnqueueUnmapMemObject}
813 event enqueue_unmap_buffer(const buffer &buffer,
815 const wait_list &events = wait_list())
817 BOOST_ASSERT(buffer.get_context() == this->get_context());
819 return enqueue_unmap_mem_object(buffer.get(), mapped_ptr, events);
822 /// Enqueues a command to unmap \p mem from the host memory space.
824 /// \see_opencl_ref{clEnqueueUnmapMemObject}
825 event enqueue_unmap_mem_object(cl_mem mem,
827 const wait_list &events = wait_list())
829 BOOST_ASSERT(m_queue != 0);
833 cl_int ret = clEnqueueUnmapMemObject(
838 events.get_event_ptr(),
842 if(ret != CL_SUCCESS){
843 BOOST_THROW_EXCEPTION(opencl_error(ret));
849 /// Enqueues a command to read data from \p image to host memory.
851 /// \see_opencl_ref{clEnqueueReadImage}
852 event enqueue_read_image(const image_object& image,
853 const size_t *origin,
854 const size_t *region,
858 const wait_list &events = wait_list())
860 BOOST_ASSERT(m_queue != 0);
864 cl_int ret = clEnqueueReadImage(
874 events.get_event_ptr(),
878 if(ret != CL_SUCCESS){
879 BOOST_THROW_EXCEPTION(opencl_error(ret));
887 event enqueue_read_image(const image_object& image,
888 const extents<N> origin,
889 const extents<N> region,
891 size_t row_pitch = 0,
892 size_t slice_pitch = 0,
893 const wait_list &events = wait_list())
895 BOOST_ASSERT(image.get_context() == this->get_context());
897 size_t origin3[3] = { 0, 0, 0 };
898 size_t region3[3] = { 1, 1, 1 };
900 std::copy(origin.data(), origin.data() + N, origin3);
901 std::copy(region.data(), region.data() + N, region3);
903 return enqueue_read_image(
904 image, origin3, region3, row_pitch, slice_pitch, host_ptr, events
908 /// Enqueues a command to write data from host memory to \p image.
910 /// \see_opencl_ref{clEnqueueWriteImage}
911 event enqueue_write_image(image_object& image,
912 const size_t *origin,
913 const size_t *region,
914 const void *host_ptr,
915 size_t input_row_pitch = 0,
916 size_t input_slice_pitch = 0,
917 const wait_list &events = wait_list())
919 BOOST_ASSERT(m_queue != 0);
923 cl_int ret = clEnqueueWriteImage(
933 events.get_event_ptr(),
937 if(ret != CL_SUCCESS){
938 BOOST_THROW_EXCEPTION(opencl_error(ret));
946 event enqueue_write_image(image_object& image,
947 const extents<N> origin,
948 const extents<N> region,
949 const void *host_ptr,
950 const size_t input_row_pitch = 0,
951 const size_t input_slice_pitch = 0,
952 const wait_list &events = wait_list())
954 BOOST_ASSERT(image.get_context() == this->get_context());
956 size_t origin3[3] = { 0, 0, 0 };
957 size_t region3[3] = { 1, 1, 1 };
959 std::copy(origin.data(), origin.data() + N, origin3);
960 std::copy(region.data(), region.data() + N, region3);
962 return enqueue_write_image(
963 image, origin3, region3, host_ptr, input_row_pitch, input_slice_pitch, events
967 /// Enqueues a command to map \p image into the host address space.
969 /// Event associated with map operation is returned through
970 /// \p map_image_event parameter.
972 /// \see_opencl_ref{clEnqueueMapImage}
973 void* enqueue_map_image(const image_object &image,
975 const size_t *origin,
976 const size_t *region,
977 size_t &output_row_pitch,
978 size_t &output_slice_pitch,
979 event &map_image_event,
980 const wait_list &events = wait_list())
982 BOOST_ASSERT(m_queue != 0);
983 BOOST_ASSERT(image.get_context() == this->get_context());
986 void *pointer = clEnqueueMapImage(
996 events.get_event_ptr(),
997 &map_image_event.get(),
1001 if(ret != CL_SUCCESS){
1002 BOOST_THROW_EXCEPTION(opencl_error(ret));
1009 void* enqueue_map_image(const image_object &image,
1011 const size_t *origin,
1012 const size_t *region,
1013 size_t &output_row_pitch,
1014 size_t &output_slice_pitch,
1015 const wait_list &events = wait_list())
1018 return enqueue_map_image(
1019 image, flags, origin, region,
1020 output_row_pitch, output_slice_pitch, event_, events
1026 void* enqueue_map_image(image_object& image,
1028 const extents<N> origin,
1029 const extents<N> region,
1030 size_t &output_row_pitch,
1031 size_t &output_slice_pitch,
1032 event &map_image_event,
1033 const wait_list &events = wait_list())
1035 BOOST_ASSERT(image.get_context() == this->get_context());
1037 size_t origin3[3] = { 0, 0, 0 };
1038 size_t region3[3] = { 1, 1, 1 };
1040 std::copy(origin.data(), origin.data() + N, origin3);
1041 std::copy(region.data(), region.data() + N, region3);
1043 return enqueue_map_image(
1044 image, flags, origin3, region3,
1045 output_row_pitch, output_slice_pitch, map_image_event, events
1051 void* enqueue_map_image(image_object& image,
1053 const extents<N> origin,
1054 const extents<N> region,
1055 size_t &output_row_pitch,
1056 size_t &output_slice_pitch,
1057 const wait_list &events = wait_list())
1060 return enqueue_map_image(
1061 image, flags, origin, region,
1062 output_row_pitch, output_slice_pitch, event_, events
1066 /// Enqueues a command to map \p image into the host address space.
1067 /// Map operation is performed asynchronously. The pointer to the mapped
1068 /// region cannot be used until the map operation has completed.
1070 /// Event associated with map operation is returned through
1071 /// \p map_image_event parameter.
1073 /// \see_opencl_ref{clEnqueueMapImage}
1074 void* enqueue_map_image_async(const image_object &image,
1076 const size_t *origin,
1077 const size_t *region,
1078 size_t &output_row_pitch,
1079 size_t &output_slice_pitch,
1080 event &map_image_event,
1081 const wait_list &events = wait_list())
1083 BOOST_ASSERT(m_queue != 0);
1084 BOOST_ASSERT(image.get_context() == this->get_context());
1087 void *pointer = clEnqueueMapImage(
1095 &output_slice_pitch,
1097 events.get_event_ptr(),
1098 &map_image_event.get(),
1102 if(ret != CL_SUCCESS){
1103 BOOST_THROW_EXCEPTION(opencl_error(ret));
1111 void* enqueue_map_image_async(image_object& image,
1113 const extents<N> origin,
1114 const extents<N> region,
1115 size_t &output_row_pitch,
1116 size_t &output_slice_pitch,
1117 event &map_image_event,
1118 const wait_list &events = wait_list())
1120 BOOST_ASSERT(image.get_context() == this->get_context());
1122 size_t origin3[3] = { 0, 0, 0 };
1123 size_t region3[3] = { 1, 1, 1 };
1125 std::copy(origin.data(), origin.data() + N, origin3);
1126 std::copy(region.data(), region.data() + N, region3);
1128 return enqueue_map_image_async(
1129 image, flags, origin3, region3,
1130 output_row_pitch, output_slice_pitch, map_image_event, events
1134 /// Enqueues a command to unmap \p image from the host memory space.
1136 /// \see_opencl_ref{clEnqueueUnmapMemObject}
1137 event enqueue_unmap_image(const image_object &image,
1139 const wait_list &events = wait_list())
1141 BOOST_ASSERT(image.get_context() == this->get_context());
1143 return enqueue_unmap_mem_object(image.get(), mapped_ptr, events);
1146 /// Enqueues a command to copy data from \p src_image to \p dst_image.
1148 /// \see_opencl_ref{clEnqueueCopyImage}
1149 event enqueue_copy_image(const image_object& src_image,
1150 image_object& dst_image,
1151 const size_t *src_origin,
1152 const size_t *dst_origin,
1153 const size_t *region,
1154 const wait_list &events = wait_list())
1156 BOOST_ASSERT(m_queue != 0);
1160 cl_int ret = clEnqueueCopyImage(
1168 events.get_event_ptr(),
1172 if(ret != CL_SUCCESS){
1173 BOOST_THROW_EXCEPTION(opencl_error(ret));
1181 event enqueue_copy_image(const image_object& src_image,
1182 image_object& dst_image,
1183 const extents<N> src_origin,
1184 const extents<N> dst_origin,
1185 const extents<N> region,
1186 const wait_list &events = wait_list())
1188 BOOST_ASSERT(src_image.get_context() == this->get_context());
1189 BOOST_ASSERT(dst_image.get_context() == this->get_context());
1190 BOOST_ASSERT_MSG(src_image.format() == dst_image.format(),
1191 "Source and destination image formats must match.");
1193 size_t src_origin3[3] = { 0, 0, 0 };
1194 size_t dst_origin3[3] = { 0, 0, 0 };
1195 size_t region3[3] = { 1, 1, 1 };
1197 std::copy(src_origin.data(), src_origin.data() + N, src_origin3);
1198 std::copy(dst_origin.data(), dst_origin.data() + N, dst_origin3);
1199 std::copy(region.data(), region.data() + N, region3);
1201 return enqueue_copy_image(
1202 src_image, dst_image, src_origin3, dst_origin3, region3, events
1206 /// Enqueues a command to copy data from \p src_image to \p dst_buffer.
1208 /// \see_opencl_ref{clEnqueueCopyImageToBuffer}
1209 event enqueue_copy_image_to_buffer(const image_object& src_image,
1210 memory_object& dst_buffer,
1211 const size_t *src_origin,
1212 const size_t *region,
1214 const wait_list &events = wait_list())
1216 BOOST_ASSERT(m_queue != 0);
1220 cl_int ret = clEnqueueCopyImageToBuffer(
1228 events.get_event_ptr(),
1232 if(ret != CL_SUCCESS){
1233 BOOST_THROW_EXCEPTION(opencl_error(ret));
1239 /// Enqueues a command to copy data from \p src_buffer to \p dst_image.
1241 /// \see_opencl_ref{clEnqueueCopyBufferToImage}
1242 event enqueue_copy_buffer_to_image(const memory_object& src_buffer,
1243 image_object& dst_image,
1245 const size_t *dst_origin,
1246 const size_t *region,
1247 const wait_list &events = wait_list())
1249 BOOST_ASSERT(m_queue != 0);
1253 cl_int ret = clEnqueueCopyBufferToImage(
1261 events.get_event_ptr(),
1265 if(ret != CL_SUCCESS){
1266 BOOST_THROW_EXCEPTION(opencl_error(ret));
1272 #if defined(CL_VERSION_1_2) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1273 /// Enqueues a command to fill \p image with \p fill_color.
1275 /// \see_opencl_ref{clEnqueueFillImage}
1277 /// \opencl_version_warning{1,2}
1278 event enqueue_fill_image(image_object& image,
1279 const void *fill_color,
1280 const size_t *origin,
1281 const size_t *region,
1282 const wait_list &events = wait_list())
1284 BOOST_ASSERT(m_queue != 0);
1288 cl_int ret = clEnqueueFillImage(
1295 events.get_event_ptr(),
1299 if(ret != CL_SUCCESS){
1300 BOOST_THROW_EXCEPTION(opencl_error(ret));
1308 event enqueue_fill_image(image_object& image,
1309 const void *fill_color,
1310 const extents<N> origin,
1311 const extents<N> region,
1312 const wait_list &events = wait_list())
1314 BOOST_ASSERT(image.get_context() == this->get_context());
1316 size_t origin3[3] = { 0, 0, 0 };
1317 size_t region3[3] = { 1, 1, 1 };
1319 std::copy(origin.data(), origin.data() + N, origin3);
1320 std::copy(region.data(), region.data() + N, region3);
1322 return enqueue_fill_image(
1323 image, fill_color, origin3, region3, events
1327 /// Enqueues a command to migrate \p mem_objects.
1329 /// \see_opencl_ref{clEnqueueMigrateMemObjects}
1331 /// \opencl_version_warning{1,2}
1332 event enqueue_migrate_memory_objects(uint_ num_mem_objects,
1333 const cl_mem *mem_objects,
1334 cl_mem_migration_flags flags,
1335 const wait_list &events = wait_list())
1337 BOOST_ASSERT(m_queue != 0);
1341 cl_int ret = clEnqueueMigrateMemObjects(
1347 events.get_event_ptr(),
1351 if(ret != CL_SUCCESS){
1352 BOOST_THROW_EXCEPTION(opencl_error(ret));
1357 #endif // CL_VERSION_1_2
1359 /// Enqueues a kernel for execution.
1361 /// \see_opencl_ref{clEnqueueNDRangeKernel}
1362 event enqueue_nd_range_kernel(const kernel &kernel,
1364 const size_t *global_work_offset,
1365 const size_t *global_work_size,
1366 const size_t *local_work_size,
1367 const wait_list &events = wait_list())
1369 BOOST_ASSERT(m_queue != 0);
1370 BOOST_ASSERT(kernel.get_context() == this->get_context());
1374 cl_int ret = clEnqueueNDRangeKernel(
1377 static_cast<cl_uint>(work_dim),
1382 events.get_event_ptr(),
1386 if(ret != CL_SUCCESS){
1387 BOOST_THROW_EXCEPTION(opencl_error(ret));
1395 event enqueue_nd_range_kernel(const kernel &kernel,
1396 const extents<N> &global_work_offset,
1397 const extents<N> &global_work_size,
1398 const extents<N> &local_work_size,
1399 const wait_list &events = wait_list())
1401 return enqueue_nd_range_kernel(
1404 global_work_offset.data(),
1405 global_work_size.data(),
1406 local_work_size.data(),
1411 /// Convenience method which calls enqueue_nd_range_kernel() with a
1412 /// one-dimensional range.
1413 event enqueue_1d_range_kernel(const kernel &kernel,
1414 size_t global_work_offset,
1415 size_t global_work_size,
1416 size_t local_work_size,
1417 const wait_list &events = wait_list())
1419 return enqueue_nd_range_kernel(
1422 &global_work_offset,
1424 local_work_size ? &local_work_size : 0,
1429 /// Enqueues a kernel to execute using a single work-item.
1431 /// \see_opencl_ref{clEnqueueTask}
1432 event enqueue_task(const kernel &kernel, const wait_list &events = wait_list())
1434 BOOST_ASSERT(m_queue != 0);
1435 BOOST_ASSERT(kernel.get_context() == this->get_context());
1439 // clEnqueueTask() was deprecated in OpenCL 2.0. In that case we
1440 // just forward to the equivalent clEnqueueNDRangeKernel() call.
1441 #ifdef CL_VERSION_2_0
1443 cl_int ret = clEnqueueNDRangeKernel(
1444 m_queue, kernel, 1, 0, &one, &one,
1445 events.size(), events.get_event_ptr(), &event_.get()
1448 cl_int ret = clEnqueueTask(
1449 m_queue, kernel, events.size(), events.get_event_ptr(), &event_.get()
1453 if(ret != CL_SUCCESS){
1454 BOOST_THROW_EXCEPTION(opencl_error(ret));
1460 /// Enqueues a function to execute on the host.
1461 event enqueue_native_kernel(void (BOOST_COMPUTE_CL_CALLBACK *user_func)(void *),
1464 uint_ num_mem_objects,
1465 const cl_mem *mem_list,
1466 const void **args_mem_loc,
1467 const wait_list &events = wait_list())
1469 BOOST_ASSERT(m_queue != 0);
1472 cl_int ret = clEnqueueNativeKernel(
1481 events.get_event_ptr(),
1484 if(ret != CL_SUCCESS){
1485 BOOST_THROW_EXCEPTION(opencl_error(ret));
1491 /// Convenience overload for enqueue_native_kernel() which enqueues a
1492 /// native kernel on the host with a nullary function.
1493 event enqueue_native_kernel(void (BOOST_COMPUTE_CL_CALLBACK *user_func)(void),
1494 const wait_list &events = wait_list())
1496 return enqueue_native_kernel(
1497 detail::nullary_native_kernel_trampoline,
1498 reinterpret_cast<void *>(&user_func),
1507 /// Flushes the command queue.
1509 /// \see_opencl_ref{clFlush}
1512 BOOST_ASSERT(m_queue != 0);
1517 /// Blocks until all outstanding commands in the queue have finished.
1519 /// \see_opencl_ref{clFinish}
1522 BOOST_ASSERT(m_queue != 0);
1527 /// Enqueues a barrier in the queue.
1528 void enqueue_barrier()
1530 BOOST_ASSERT(m_queue != 0);
1531 cl_int ret = CL_SUCCESS;
1533 #ifdef CL_VERSION_1_2
1534 if(get_device().check_version(1, 2)){
1535 ret = clEnqueueBarrierWithWaitList(m_queue, 0, 0, 0);
1537 #endif // CL_VERSION_1_2
1539 // Suppress deprecated declarations warning
1540 BOOST_COMPUTE_DISABLE_DEPRECATED_DECLARATIONS();
1541 ret = clEnqueueBarrier(m_queue);
1542 BOOST_COMPUTE_ENABLE_DEPRECATED_DECLARATIONS();
1545 if(ret != CL_SUCCESS){
1546 BOOST_THROW_EXCEPTION(opencl_error(ret));
1550 #if defined(CL_VERSION_1_2) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1551 /// Enqueues a barrier in the queue after \p events.
1553 /// \opencl_version_warning{1,2}
1554 event enqueue_barrier(const wait_list &events)
1556 BOOST_ASSERT(m_queue != 0);
1559 cl_int ret = CL_SUCCESS;
1561 ret = clEnqueueBarrierWithWaitList(
1562 m_queue, events.size(), events.get_event_ptr(), &event_.get()
1565 if(ret != CL_SUCCESS){
1566 BOOST_THROW_EXCEPTION(opencl_error(ret));
1571 #endif // CL_VERSION_1_2
1573 /// Enqueues a marker in the queue and returns an event that can be
1574 /// used to track its progress.
1575 event enqueue_marker()
1578 cl_int ret = CL_SUCCESS;
1580 #ifdef CL_VERSION_1_2
1581 if(get_device().check_version(1, 2)){
1582 ret = clEnqueueMarkerWithWaitList(m_queue, 0, 0, &event_.get());
1586 // Suppress deprecated declarations warning
1587 BOOST_COMPUTE_DISABLE_DEPRECATED_DECLARATIONS();
1588 ret = clEnqueueMarker(m_queue, &event_.get());
1589 BOOST_COMPUTE_ENABLE_DEPRECATED_DECLARATIONS();
1592 if(ret != CL_SUCCESS){
1593 BOOST_THROW_EXCEPTION(opencl_error(ret));
1599 #if defined(CL_VERSION_1_2) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1600 /// Enqueues a marker after \p events in the queue and returns an
1601 /// event that can be used to track its progress.
1603 /// \opencl_version_warning{1,2}
1604 event enqueue_marker(const wait_list &events)
1608 cl_int ret = clEnqueueMarkerWithWaitList(
1609 m_queue, events.size(), events.get_event_ptr(), &event_.get()
1612 if(ret != CL_SUCCESS){
1613 BOOST_THROW_EXCEPTION(opencl_error(ret));
1618 #endif // CL_VERSION_1_2
1620 #if defined(CL_VERSION_2_0) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1621 /// Enqueues a command to copy \p size bytes of data from \p src_ptr to
1624 /// \opencl_version_warning{2,0}
1626 /// \see_opencl2_ref{clEnqueueSVMMemcpy}
1627 event enqueue_svm_memcpy(void *dst_ptr,
1628 const void *src_ptr,
1630 const wait_list &events = wait_list())
1634 cl_int ret = clEnqueueSVMMemcpy(
1641 events.get_event_ptr(),
1645 if(ret != CL_SUCCESS){
1646 BOOST_THROW_EXCEPTION(opencl_error(ret));
1652 /// Enqueues a command to copy \p size bytes of data from \p src_ptr to
1653 /// \p dst_ptr. The operation is performed asynchronously.
1655 /// \opencl_version_warning{2,0}
1657 /// \see_opencl2_ref{clEnqueueSVMMemcpy}
1658 event enqueue_svm_memcpy_async(void *dst_ptr,
1659 const void *src_ptr,
1661 const wait_list &events = wait_list())
1665 cl_int ret = clEnqueueSVMMemcpy(
1672 events.get_event_ptr(),
1676 if(ret != CL_SUCCESS){
1677 BOOST_THROW_EXCEPTION(opencl_error(ret));
1683 /// Enqueues a command to fill \p size bytes of data at \p svm_ptr with
1686 /// \opencl_version_warning{2,0}
1688 /// \see_opencl2_ref{clEnqueueSVMMemFill}
1689 event enqueue_svm_fill(void *svm_ptr,
1690 const void *pattern,
1691 size_t pattern_size,
1693 const wait_list &events = wait_list())
1698 cl_int ret = clEnqueueSVMMemFill(
1705 events.get_event_ptr(),
1709 if(ret != CL_SUCCESS){
1710 BOOST_THROW_EXCEPTION(opencl_error(ret));
1716 /// Enqueues a command to free \p svm_ptr.
1718 /// \opencl_version_warning{2,0}
1720 /// \see_opencl2_ref{clEnqueueSVMFree}
1723 event enqueue_svm_free(void *svm_ptr,
1724 const wait_list &events = wait_list())
1728 cl_int ret = clEnqueueSVMFree(
1735 events.get_event_ptr(),
1739 if(ret != CL_SUCCESS){
1740 BOOST_THROW_EXCEPTION(opencl_error(ret));
1746 /// Enqueues a command to map \p svm_ptr to the host memory space.
1748 /// \opencl_version_warning{2,0}
1750 /// \see_opencl2_ref{clEnqueueSVMMap}
1751 event enqueue_svm_map(void *svm_ptr,
1754 const wait_list &events = wait_list())
1758 cl_int ret = clEnqueueSVMMap(
1765 events.get_event_ptr(),
1769 if(ret != CL_SUCCESS){
1770 BOOST_THROW_EXCEPTION(opencl_error(ret));
1776 /// Enqueues a command to unmap \p svm_ptr from the host memory space.
1778 /// \opencl_version_warning{2,0}
1780 /// \see_opencl2_ref{clEnqueueSVMUnmap}
1781 event enqueue_svm_unmap(void *svm_ptr,
1782 const wait_list &events = wait_list())
1786 cl_int ret = clEnqueueSVMUnmap(
1790 events.get_event_ptr(),
1794 if(ret != CL_SUCCESS){
1795 BOOST_THROW_EXCEPTION(opencl_error(ret));
1800 #endif // CL_VERSION_2_0
1802 /// Returns \c true if the command queue is the same at \p other.
1803 bool operator==(const command_queue &other) const
1805 return m_queue == other.m_queue;
1808 /// Returns \c true if the command queue is different from \p other.
1809 bool operator!=(const command_queue &other) const
1811 return m_queue != other.m_queue;
1815 operator cl_command_queue() const
1821 bool check_device_version(int major, int minor) const
1823 return get_device().check_version(major, minor);
1827 cl_command_queue m_queue;
1830 inline buffer buffer::clone(command_queue &queue) const
1832 buffer copy(get_context(), size(), get_memory_flags());
1833 queue.enqueue_copy_buffer(*this, copy, 0, 0, size());
1837 inline image1d image1d::clone(command_queue &queue) const
1840 get_context(), width(), format(), get_memory_flags()
1843 queue.enqueue_copy_image(*this, copy, origin(), copy.origin(), size());
1848 inline image2d image2d::clone(command_queue &queue) const
1851 get_context(), width(), height(), format(), get_memory_flags()
1854 queue.enqueue_copy_image(*this, copy, origin(), copy.origin(), size());
1859 inline image3d image3d::clone(command_queue &queue) const
1862 get_context(), width(), height(), depth(), format(), get_memory_flags()
1865 queue.enqueue_copy_image(*this, copy, origin(), copy.origin(), size());
1870 /// \internal_ define get_info() specializations for command_queue
1871 BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_SPECIALIZATIONS(command_queue,
1872 ((cl_context, CL_QUEUE_CONTEXT))
1873 ((cl_device_id, CL_QUEUE_DEVICE))
1874 ((uint_, CL_QUEUE_REFERENCE_COUNT))
1875 ((cl_command_queue_properties, CL_QUEUE_PROPERTIES))
1878 } // end compute namespace
1879 } // end boost namespace
1881 #endif // BOOST_COMPUTE_COMMAND_QUEUE_HPP