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
84 #ifdef BOOST_COMPUTE_CL_VERSION_2_0
86 on_device = CL_QUEUE_ON_DEVICE,
87 on_device_default = CL_QUEUE_ON_DEVICE_DEFAULT
92 map_read = CL_MAP_READ,
93 map_write = CL_MAP_WRITE
94 #ifdef BOOST_COMPUTE_CL_VERSION_1_2
96 map_write_invalidate_region = CL_MAP_WRITE_INVALIDATE_REGION
100 #ifdef BOOST_COMPUTE_CL_VERSION_1_2
101 enum mem_migration_flags {
102 migrate_to_host = CL_MIGRATE_MEM_OBJECT_HOST,
103 migrate_content_undefined = CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED
105 #endif // BOOST_COMPUTE_CL_VERSION_1_2
107 /// Creates a null command queue.
113 explicit command_queue(cl_command_queue queue, bool retain = true)
116 if(m_queue && retain){
117 clRetainCommandQueue(m_queue);
121 /// Creates a command queue in \p context for \p device with
124 /// \see_opencl_ref{clCreateCommandQueue}
125 command_queue(const context &context,
126 const device &device,
127 cl_command_queue_properties properties = 0)
129 BOOST_ASSERT(device.id() != 0);
133 #ifdef BOOST_COMPUTE_CL_VERSION_2_0
134 if (device.check_version(2, 0)){
135 std::vector<cl_queue_properties> queue_properties;
137 queue_properties.push_back(CL_QUEUE_PROPERTIES);
138 queue_properties.push_back(cl_queue_properties(properties));
139 queue_properties.push_back(cl_queue_properties(0));
142 const cl_queue_properties *queue_properties_ptr =
143 queue_properties.empty() ? 0 : &queue_properties[0];
145 m_queue = clCreateCommandQueueWithProperties(
146 context, device.id(), queue_properties_ptr, &error
151 // Suppress deprecated declarations warning
152 BOOST_COMPUTE_DISABLE_DEPRECATED_DECLARATIONS();
153 m_queue = clCreateCommandQueue(
154 context, device.id(), properties, &error
156 BOOST_COMPUTE_ENABLE_DEPRECATED_DECLARATIONS();
160 BOOST_THROW_EXCEPTION(opencl_error(error));
164 /// Creates a new command queue object as a copy of \p other.
165 command_queue(const command_queue &other)
166 : m_queue(other.m_queue)
169 clRetainCommandQueue(m_queue);
173 /// Copies the command queue object from \p other to \c *this.
174 command_queue& operator=(const command_queue &other)
178 clReleaseCommandQueue(m_queue);
181 m_queue = other.m_queue;
184 clRetainCommandQueue(m_queue);
191 #ifndef BOOST_COMPUTE_NO_RVALUE_REFERENCES
192 /// Move-constructs a new command queue object from \p other.
193 command_queue(command_queue&& other) BOOST_NOEXCEPT
194 : m_queue(other.m_queue)
199 /// Move-assigns the command queue from \p other to \c *this.
200 command_queue& operator=(command_queue&& other) BOOST_NOEXCEPT
203 clReleaseCommandQueue(m_queue);
206 m_queue = other.m_queue;
211 #endif // BOOST_COMPUTE_NO_RVALUE_REFERENCES
213 /// Destroys the command queue.
215 /// \see_opencl_ref{clReleaseCommandQueue}
219 BOOST_COMPUTE_ASSERT_CL_SUCCESS(
220 clReleaseCommandQueue(m_queue)
225 /// Returns the underlying OpenCL command queue.
226 cl_command_queue& get() const
228 return const_cast<cl_command_queue &>(m_queue);
231 /// Returns the device that the command queue issues commands to.
232 device get_device() const
234 return device(get_info<cl_device_id>(CL_QUEUE_DEVICE));
237 /// Returns the context for the command queue.
238 context get_context() const
240 return context(get_info<cl_context>(CL_QUEUE_CONTEXT));
243 /// Returns information about the command queue.
245 /// \see_opencl_ref{clGetCommandQueueInfo}
247 T get_info(cl_command_queue_info info) const
249 return detail::get_object_info<T>(clGetCommandQueueInfo, m_queue, info);
254 typename detail::get_object_info_type<command_queue, Enum>::type
257 /// Returns the properties for the command queue.
258 cl_command_queue_properties get_properties() const
260 return get_info<cl_command_queue_properties>(CL_QUEUE_PROPERTIES);
263 #if defined(BOOST_COMPUTE_CL_VERSION_2_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
264 /// Returns the current default device command queue for the underlying device.
266 /// \opencl_version_warning{2,1}
267 command_queue get_default_device_queue() const
269 return command_queue(get_info<cl_command_queue>(CL_QUEUE_DEVICE_DEFAULT));
272 /// Replaces the default device command queue for the underlying device
273 /// with this command queue. Command queue must have been created
274 /// with CL_QUEUE_ON_DEVICE flag.
276 /// \see_opencl21_ref{clSetDefaultDeviceCommandQueue}
278 /// \opencl_version_warning{2,1}
279 void set_as_default_device_queue() const
281 cl_int ret = clSetDefaultDeviceCommandQueue(
282 this->get_context().get(),
283 this->get_device().get(),
286 if(ret != CL_SUCCESS){
287 BOOST_THROW_EXCEPTION(opencl_error(ret));
290 #endif // BOOST_COMPUTE_CL_VERSION_2_1
292 /// Enqueues a command to read data from \p buffer to host memory.
294 /// \see_opencl_ref{clEnqueueReadBuffer}
297 event enqueue_read_buffer(const buffer &buffer,
301 const wait_list &events = wait_list())
303 BOOST_ASSERT(m_queue != 0);
304 BOOST_ASSERT(size <= buffer.size());
305 BOOST_ASSERT(buffer.get_context() == this->get_context());
306 BOOST_ASSERT(host_ptr != 0);
310 cl_int ret = clEnqueueReadBuffer(
318 events.get_event_ptr(),
322 if(ret != CL_SUCCESS){
323 BOOST_THROW_EXCEPTION(opencl_error(ret));
329 /// Enqueues a command to read data from \p buffer to host memory. The
330 /// copy will be performed asynchronously.
332 /// \see_opencl_ref{clEnqueueReadBuffer}
334 /// \see copy_async()
335 event enqueue_read_buffer_async(const buffer &buffer,
339 const wait_list &events = wait_list())
341 BOOST_ASSERT(m_queue != 0);
342 BOOST_ASSERT(size <= buffer.size());
343 BOOST_ASSERT(buffer.get_context() == this->get_context());
344 BOOST_ASSERT(host_ptr != 0);
348 cl_int ret = clEnqueueReadBuffer(
356 events.get_event_ptr(),
360 if(ret != CL_SUCCESS){
361 BOOST_THROW_EXCEPTION(opencl_error(ret));
367 #if defined(BOOST_COMPUTE_CL_VERSION_1_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
368 /// Enqueues a command to read a rectangular region from \p buffer to
371 /// \see_opencl_ref{clEnqueueReadBufferRect}
373 /// \opencl_version_warning{1,1}
374 event enqueue_read_buffer_rect(const buffer &buffer,
375 const size_t buffer_origin[3],
376 const size_t host_origin[3],
377 const size_t region[3],
378 size_t buffer_row_pitch,
379 size_t buffer_slice_pitch,
380 size_t host_row_pitch,
381 size_t host_slice_pitch,
383 const wait_list &events = wait_list())
385 BOOST_ASSERT(m_queue != 0);
386 BOOST_ASSERT(buffer.get_context() == this->get_context());
387 BOOST_ASSERT(host_ptr != 0);
391 cl_int ret = clEnqueueReadBufferRect(
404 events.get_event_ptr(),
408 if(ret != CL_SUCCESS){
409 BOOST_THROW_EXCEPTION(opencl_error(ret));
415 /// Enqueues a command to read a rectangular region from \p buffer to
416 /// host memory. The copy will be performed asynchronously.
418 /// \see_opencl_ref{clEnqueueReadBufferRect}
420 /// \opencl_version_warning{1,1}
421 event enqueue_read_buffer_rect_async(const buffer &buffer,
422 const size_t buffer_origin[3],
423 const size_t host_origin[3],
424 const size_t region[3],
425 size_t buffer_row_pitch,
426 size_t buffer_slice_pitch,
427 size_t host_row_pitch,
428 size_t host_slice_pitch,
430 const wait_list &events = wait_list())
432 BOOST_ASSERT(m_queue != 0);
433 BOOST_ASSERT(buffer.get_context() == this->get_context());
434 BOOST_ASSERT(host_ptr != 0);
438 cl_int ret = clEnqueueReadBufferRect(
451 events.get_event_ptr(),
455 if(ret != CL_SUCCESS){
456 BOOST_THROW_EXCEPTION(opencl_error(ret));
461 #endif // BOOST_COMPUTE_CL_VERSION_1_1
463 /// Enqueues a command to write data from host memory to \p buffer.
465 /// \see_opencl_ref{clEnqueueWriteBuffer}
468 event enqueue_write_buffer(const buffer &buffer,
471 const void *host_ptr,
472 const wait_list &events = wait_list())
474 BOOST_ASSERT(m_queue != 0);
475 BOOST_ASSERT(size <= buffer.size());
476 BOOST_ASSERT(buffer.get_context() == this->get_context());
477 BOOST_ASSERT(host_ptr != 0);
481 cl_int ret = clEnqueueWriteBuffer(
489 events.get_event_ptr(),
493 if(ret != CL_SUCCESS){
494 BOOST_THROW_EXCEPTION(opencl_error(ret));
500 /// Enqueues a command to write data from host memory to \p buffer.
501 /// The copy is performed asynchronously.
503 /// \see_opencl_ref{clEnqueueWriteBuffer}
505 /// \see copy_async()
506 event enqueue_write_buffer_async(const buffer &buffer,
509 const void *host_ptr,
510 const wait_list &events = wait_list())
512 BOOST_ASSERT(m_queue != 0);
513 BOOST_ASSERT(size <= buffer.size());
514 BOOST_ASSERT(buffer.get_context() == this->get_context());
515 BOOST_ASSERT(host_ptr != 0);
519 cl_int ret = clEnqueueWriteBuffer(
527 events.get_event_ptr(),
531 if(ret != CL_SUCCESS){
532 BOOST_THROW_EXCEPTION(opencl_error(ret));
538 #if defined(BOOST_COMPUTE_CL_VERSION_1_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
539 /// Enqueues a command to write a rectangular region from host memory
542 /// \see_opencl_ref{clEnqueueWriteBufferRect}
544 /// \opencl_version_warning{1,1}
545 event enqueue_write_buffer_rect(const buffer &buffer,
546 const size_t buffer_origin[3],
547 const size_t host_origin[3],
548 const size_t region[3],
549 size_t buffer_row_pitch,
550 size_t buffer_slice_pitch,
551 size_t host_row_pitch,
552 size_t host_slice_pitch,
554 const wait_list &events = wait_list())
556 BOOST_ASSERT(m_queue != 0);
557 BOOST_ASSERT(buffer.get_context() == this->get_context());
558 BOOST_ASSERT(host_ptr != 0);
562 cl_int ret = clEnqueueWriteBufferRect(
575 events.get_event_ptr(),
579 if(ret != CL_SUCCESS){
580 BOOST_THROW_EXCEPTION(opencl_error(ret));
586 /// Enqueues a command to write a rectangular region from host memory
587 /// to \p buffer. The copy is performed asynchronously.
589 /// \see_opencl_ref{clEnqueueWriteBufferRect}
591 /// \opencl_version_warning{1,1}
592 event enqueue_write_buffer_rect_async(const buffer &buffer,
593 const size_t buffer_origin[3],
594 const size_t host_origin[3],
595 const size_t region[3],
596 size_t buffer_row_pitch,
597 size_t buffer_slice_pitch,
598 size_t host_row_pitch,
599 size_t host_slice_pitch,
601 const wait_list &events = wait_list())
603 BOOST_ASSERT(m_queue != 0);
604 BOOST_ASSERT(buffer.get_context() == this->get_context());
605 BOOST_ASSERT(host_ptr != 0);
609 cl_int ret = clEnqueueWriteBufferRect(
622 events.get_event_ptr(),
626 if(ret != CL_SUCCESS){
627 BOOST_THROW_EXCEPTION(opencl_error(ret));
632 #endif // BOOST_COMPUTE_CL_VERSION_1_1
634 /// Enqueues a command to copy data from \p src_buffer to
637 /// \see_opencl_ref{clEnqueueCopyBuffer}
640 event enqueue_copy_buffer(const buffer &src_buffer,
641 const buffer &dst_buffer,
645 const wait_list &events = wait_list())
647 BOOST_ASSERT(m_queue != 0);
648 BOOST_ASSERT(src_offset + size <= src_buffer.size());
649 BOOST_ASSERT(dst_offset + size <= dst_buffer.size());
650 BOOST_ASSERT(src_buffer.get_context() == this->get_context());
651 BOOST_ASSERT(dst_buffer.get_context() == this->get_context());
655 cl_int ret = clEnqueueCopyBuffer(
663 events.get_event_ptr(),
667 if(ret != CL_SUCCESS){
668 BOOST_THROW_EXCEPTION(opencl_error(ret));
674 #if defined(BOOST_COMPUTE_CL_VERSION_1_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
675 /// Enqueues a command to copy a rectangular region from
676 /// \p src_buffer to \p dst_buffer.
678 /// \see_opencl_ref{clEnqueueCopyBufferRect}
680 /// \opencl_version_warning{1,1}
681 event enqueue_copy_buffer_rect(const buffer &src_buffer,
682 const buffer &dst_buffer,
683 const size_t src_origin[3],
684 const size_t dst_origin[3],
685 const size_t region[3],
686 size_t buffer_row_pitch,
687 size_t buffer_slice_pitch,
688 size_t host_row_pitch,
689 size_t host_slice_pitch,
690 const wait_list &events = wait_list())
692 BOOST_ASSERT(m_queue != 0);
693 BOOST_ASSERT(src_buffer.get_context() == this->get_context());
694 BOOST_ASSERT(dst_buffer.get_context() == this->get_context());
698 cl_int ret = clEnqueueCopyBufferRect(
710 events.get_event_ptr(),
714 if(ret != CL_SUCCESS){
715 BOOST_THROW_EXCEPTION(opencl_error(ret));
720 #endif // BOOST_COMPUTE_CL_VERSION_1_1
722 #if defined(BOOST_COMPUTE_CL_VERSION_1_2) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
723 /// Enqueues a command to fill \p buffer with \p pattern.
725 /// \see_opencl_ref{clEnqueueFillBuffer}
727 /// \opencl_version_warning{1,2}
730 event enqueue_fill_buffer(const buffer &buffer,
735 const wait_list &events = wait_list())
737 BOOST_ASSERT(m_queue != 0);
738 BOOST_ASSERT(offset + size <= buffer.size());
739 BOOST_ASSERT(buffer.get_context() == this->get_context());
743 cl_int ret = clEnqueueFillBuffer(
751 events.get_event_ptr(),
755 if(ret != CL_SUCCESS){
756 BOOST_THROW_EXCEPTION(opencl_error(ret));
761 #endif // BOOST_COMPUTE_CL_VERSION_1_2
763 /// Enqueues a command to map \p buffer into the host address space.
764 /// Event associated with map operation is returned through
765 /// \p map_buffer_event parameter.
767 /// \see_opencl_ref{clEnqueueMapBuffer}
768 void* enqueue_map_buffer(const buffer &buffer,
772 event &map_buffer_event,
773 const wait_list &events = wait_list())
775 BOOST_ASSERT(m_queue != 0);
776 BOOST_ASSERT(offset + size <= buffer.size());
777 BOOST_ASSERT(buffer.get_context() == this->get_context());
780 void *pointer = clEnqueueMapBuffer(
788 events.get_event_ptr(),
789 &map_buffer_event.get(),
793 if(ret != CL_SUCCESS){
794 BOOST_THROW_EXCEPTION(opencl_error(ret));
801 void* enqueue_map_buffer(const buffer &buffer,
805 const wait_list &events = wait_list())
808 return enqueue_map_buffer(buffer, flags, offset, size, event_, events);
811 /// Enqueues a command to map \p buffer into the host address space.
812 /// Map operation is performed asynchronously. The pointer to the mapped
813 /// region cannot be used until the map operation has completed.
815 /// Event associated with map operation is returned through
816 /// \p map_buffer_event parameter.
818 /// \see_opencl_ref{clEnqueueMapBuffer}
819 void* enqueue_map_buffer_async(const buffer &buffer,
823 event &map_buffer_event,
824 const wait_list &events = wait_list())
826 BOOST_ASSERT(m_queue != 0);
827 BOOST_ASSERT(offset + size <= buffer.size());
828 BOOST_ASSERT(buffer.get_context() == this->get_context());
831 void *pointer = clEnqueueMapBuffer(
839 events.get_event_ptr(),
840 &map_buffer_event.get(),
844 if(ret != CL_SUCCESS){
845 BOOST_THROW_EXCEPTION(opencl_error(ret));
851 /// Enqueues a command to unmap \p buffer from the host memory space.
853 /// \see_opencl_ref{clEnqueueUnmapMemObject}
854 event enqueue_unmap_buffer(const buffer &buffer,
856 const wait_list &events = wait_list())
858 BOOST_ASSERT(buffer.get_context() == this->get_context());
860 return enqueue_unmap_mem_object(buffer.get(), mapped_ptr, events);
863 /// Enqueues a command to unmap \p mem from the host memory space.
865 /// \see_opencl_ref{clEnqueueUnmapMemObject}
866 event enqueue_unmap_mem_object(cl_mem mem,
868 const wait_list &events = wait_list())
870 BOOST_ASSERT(m_queue != 0);
874 cl_int ret = clEnqueueUnmapMemObject(
879 events.get_event_ptr(),
883 if(ret != CL_SUCCESS){
884 BOOST_THROW_EXCEPTION(opencl_error(ret));
890 /// Enqueues a command to read data from \p image to host memory.
892 /// \see_opencl_ref{clEnqueueReadImage}
893 event enqueue_read_image(const image_object& image,
894 const size_t *origin,
895 const size_t *region,
899 const wait_list &events = wait_list())
901 BOOST_ASSERT(m_queue != 0);
905 cl_int ret = clEnqueueReadImage(
915 events.get_event_ptr(),
919 if(ret != CL_SUCCESS){
920 BOOST_THROW_EXCEPTION(opencl_error(ret));
928 event enqueue_read_image(const image_object& image,
929 const extents<N> origin,
930 const extents<N> region,
932 size_t row_pitch = 0,
933 size_t slice_pitch = 0,
934 const wait_list &events = wait_list())
936 BOOST_ASSERT(image.get_context() == this->get_context());
938 size_t origin3[3] = { 0, 0, 0 };
939 size_t region3[3] = { 1, 1, 1 };
941 std::copy(origin.data(), origin.data() + N, origin3);
942 std::copy(region.data(), region.data() + N, region3);
944 return enqueue_read_image(
945 image, origin3, region3, row_pitch, slice_pitch, host_ptr, events
949 /// Enqueues a command to write data from host memory to \p image.
951 /// \see_opencl_ref{clEnqueueWriteImage}
952 event enqueue_write_image(image_object& image,
953 const size_t *origin,
954 const size_t *region,
955 const void *host_ptr,
956 size_t input_row_pitch = 0,
957 size_t input_slice_pitch = 0,
958 const wait_list &events = wait_list())
960 BOOST_ASSERT(m_queue != 0);
964 cl_int ret = clEnqueueWriteImage(
974 events.get_event_ptr(),
978 if(ret != CL_SUCCESS){
979 BOOST_THROW_EXCEPTION(opencl_error(ret));
987 event enqueue_write_image(image_object& image,
988 const extents<N> origin,
989 const extents<N> region,
990 const void *host_ptr,
991 const size_t input_row_pitch = 0,
992 const size_t input_slice_pitch = 0,
993 const wait_list &events = wait_list())
995 BOOST_ASSERT(image.get_context() == this->get_context());
997 size_t origin3[3] = { 0, 0, 0 };
998 size_t region3[3] = { 1, 1, 1 };
1000 std::copy(origin.data(), origin.data() + N, origin3);
1001 std::copy(region.data(), region.data() + N, region3);
1003 return enqueue_write_image(
1004 image, origin3, region3, host_ptr, input_row_pitch, input_slice_pitch, events
1008 /// Enqueues a command to map \p image into the host address space.
1010 /// Event associated with map operation is returned through
1011 /// \p map_image_event parameter.
1013 /// \see_opencl_ref{clEnqueueMapImage}
1014 void* enqueue_map_image(const image_object &image,
1016 const size_t *origin,
1017 const size_t *region,
1018 size_t &output_row_pitch,
1019 size_t &output_slice_pitch,
1020 event &map_image_event,
1021 const wait_list &events = wait_list())
1023 BOOST_ASSERT(m_queue != 0);
1024 BOOST_ASSERT(image.get_context() == this->get_context());
1027 void *pointer = clEnqueueMapImage(
1035 &output_slice_pitch,
1037 events.get_event_ptr(),
1038 &map_image_event.get(),
1042 if(ret != CL_SUCCESS){
1043 BOOST_THROW_EXCEPTION(opencl_error(ret));
1050 void* enqueue_map_image(const image_object &image,
1052 const size_t *origin,
1053 const size_t *region,
1054 size_t &output_row_pitch,
1055 size_t &output_slice_pitch,
1056 const wait_list &events = wait_list())
1059 return enqueue_map_image(
1060 image, flags, origin, region,
1061 output_row_pitch, output_slice_pitch, event_, events
1067 void* enqueue_map_image(image_object& image,
1069 const extents<N> origin,
1070 const extents<N> region,
1071 size_t &output_row_pitch,
1072 size_t &output_slice_pitch,
1073 event &map_image_event,
1074 const wait_list &events = wait_list())
1076 BOOST_ASSERT(image.get_context() == this->get_context());
1078 size_t origin3[3] = { 0, 0, 0 };
1079 size_t region3[3] = { 1, 1, 1 };
1081 std::copy(origin.data(), origin.data() + N, origin3);
1082 std::copy(region.data(), region.data() + N, region3);
1084 return enqueue_map_image(
1085 image, flags, origin3, region3,
1086 output_row_pitch, output_slice_pitch, map_image_event, events
1092 void* enqueue_map_image(image_object& image,
1094 const extents<N> origin,
1095 const extents<N> region,
1096 size_t &output_row_pitch,
1097 size_t &output_slice_pitch,
1098 const wait_list &events = wait_list())
1101 return enqueue_map_image(
1102 image, flags, origin, region,
1103 output_row_pitch, output_slice_pitch, event_, events
1107 /// Enqueues a command to map \p image into the host address space.
1108 /// Map operation is performed asynchronously. The pointer to the mapped
1109 /// region cannot be used until the map operation has completed.
1111 /// Event associated with map operation is returned through
1112 /// \p map_image_event parameter.
1114 /// \see_opencl_ref{clEnqueueMapImage}
1115 void* enqueue_map_image_async(const image_object &image,
1117 const size_t *origin,
1118 const size_t *region,
1119 size_t &output_row_pitch,
1120 size_t &output_slice_pitch,
1121 event &map_image_event,
1122 const wait_list &events = wait_list())
1124 BOOST_ASSERT(m_queue != 0);
1125 BOOST_ASSERT(image.get_context() == this->get_context());
1128 void *pointer = clEnqueueMapImage(
1136 &output_slice_pitch,
1138 events.get_event_ptr(),
1139 &map_image_event.get(),
1143 if(ret != CL_SUCCESS){
1144 BOOST_THROW_EXCEPTION(opencl_error(ret));
1152 void* enqueue_map_image_async(image_object& image,
1154 const extents<N> origin,
1155 const extents<N> region,
1156 size_t &output_row_pitch,
1157 size_t &output_slice_pitch,
1158 event &map_image_event,
1159 const wait_list &events = wait_list())
1161 BOOST_ASSERT(image.get_context() == this->get_context());
1163 size_t origin3[3] = { 0, 0, 0 };
1164 size_t region3[3] = { 1, 1, 1 };
1166 std::copy(origin.data(), origin.data() + N, origin3);
1167 std::copy(region.data(), region.data() + N, region3);
1169 return enqueue_map_image_async(
1170 image, flags, origin3, region3,
1171 output_row_pitch, output_slice_pitch, map_image_event, events
1175 /// Enqueues a command to unmap \p image from the host memory space.
1177 /// \see_opencl_ref{clEnqueueUnmapMemObject}
1178 event enqueue_unmap_image(const image_object &image,
1180 const wait_list &events = wait_list())
1182 BOOST_ASSERT(image.get_context() == this->get_context());
1184 return enqueue_unmap_mem_object(image.get(), mapped_ptr, events);
1187 /// Enqueues a command to copy data from \p src_image to \p dst_image.
1189 /// \see_opencl_ref{clEnqueueCopyImage}
1190 event enqueue_copy_image(const image_object& src_image,
1191 image_object& dst_image,
1192 const size_t *src_origin,
1193 const size_t *dst_origin,
1194 const size_t *region,
1195 const wait_list &events = wait_list())
1197 BOOST_ASSERT(m_queue != 0);
1201 cl_int ret = clEnqueueCopyImage(
1209 events.get_event_ptr(),
1213 if(ret != CL_SUCCESS){
1214 BOOST_THROW_EXCEPTION(opencl_error(ret));
1222 event enqueue_copy_image(const image_object& src_image,
1223 image_object& dst_image,
1224 const extents<N> src_origin,
1225 const extents<N> dst_origin,
1226 const extents<N> region,
1227 const wait_list &events = wait_list())
1229 BOOST_ASSERT(src_image.get_context() == this->get_context());
1230 BOOST_ASSERT(dst_image.get_context() == this->get_context());
1231 BOOST_ASSERT_MSG(src_image.format() == dst_image.format(),
1232 "Source and destination image formats must match.");
1234 size_t src_origin3[3] = { 0, 0, 0 };
1235 size_t dst_origin3[3] = { 0, 0, 0 };
1236 size_t region3[3] = { 1, 1, 1 };
1238 std::copy(src_origin.data(), src_origin.data() + N, src_origin3);
1239 std::copy(dst_origin.data(), dst_origin.data() + N, dst_origin3);
1240 std::copy(region.data(), region.data() + N, region3);
1242 return enqueue_copy_image(
1243 src_image, dst_image, src_origin3, dst_origin3, region3, events
1247 /// Enqueues a command to copy data from \p src_image to \p dst_buffer.
1249 /// \see_opencl_ref{clEnqueueCopyImageToBuffer}
1250 event enqueue_copy_image_to_buffer(const image_object& src_image,
1251 memory_object& dst_buffer,
1252 const size_t *src_origin,
1253 const size_t *region,
1255 const wait_list &events = wait_list())
1257 BOOST_ASSERT(m_queue != 0);
1261 cl_int ret = clEnqueueCopyImageToBuffer(
1269 events.get_event_ptr(),
1273 if(ret != CL_SUCCESS){
1274 BOOST_THROW_EXCEPTION(opencl_error(ret));
1280 /// Enqueues a command to copy data from \p src_buffer to \p dst_image.
1282 /// \see_opencl_ref{clEnqueueCopyBufferToImage}
1283 event enqueue_copy_buffer_to_image(const memory_object& src_buffer,
1284 image_object& dst_image,
1286 const size_t *dst_origin,
1287 const size_t *region,
1288 const wait_list &events = wait_list())
1290 BOOST_ASSERT(m_queue != 0);
1294 cl_int ret = clEnqueueCopyBufferToImage(
1302 events.get_event_ptr(),
1306 if(ret != CL_SUCCESS){
1307 BOOST_THROW_EXCEPTION(opencl_error(ret));
1313 #if defined(BOOST_COMPUTE_CL_VERSION_1_2) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1314 /// Enqueues a command to fill \p image with \p fill_color.
1316 /// \see_opencl_ref{clEnqueueFillImage}
1318 /// \opencl_version_warning{1,2}
1319 event enqueue_fill_image(image_object& image,
1320 const void *fill_color,
1321 const size_t *origin,
1322 const size_t *region,
1323 const wait_list &events = wait_list())
1325 BOOST_ASSERT(m_queue != 0);
1329 cl_int ret = clEnqueueFillImage(
1336 events.get_event_ptr(),
1340 if(ret != CL_SUCCESS){
1341 BOOST_THROW_EXCEPTION(opencl_error(ret));
1349 event enqueue_fill_image(image_object& image,
1350 const void *fill_color,
1351 const extents<N> origin,
1352 const extents<N> region,
1353 const wait_list &events = wait_list())
1355 BOOST_ASSERT(image.get_context() == this->get_context());
1357 size_t origin3[3] = { 0, 0, 0 };
1358 size_t region3[3] = { 1, 1, 1 };
1360 std::copy(origin.data(), origin.data() + N, origin3);
1361 std::copy(region.data(), region.data() + N, region3);
1363 return enqueue_fill_image(
1364 image, fill_color, origin3, region3, events
1368 /// Enqueues a command to migrate \p mem_objects.
1370 /// \see_opencl_ref{clEnqueueMigrateMemObjects}
1372 /// \opencl_version_warning{1,2}
1373 event enqueue_migrate_memory_objects(uint_ num_mem_objects,
1374 const cl_mem *mem_objects,
1375 cl_mem_migration_flags flags,
1376 const wait_list &events = wait_list())
1378 BOOST_ASSERT(m_queue != 0);
1382 cl_int ret = clEnqueueMigrateMemObjects(
1388 events.get_event_ptr(),
1392 if(ret != CL_SUCCESS){
1393 BOOST_THROW_EXCEPTION(opencl_error(ret));
1398 #endif // BOOST_COMPUTE_CL_VERSION_1_2
1400 /// Enqueues a kernel for execution.
1402 /// \see_opencl_ref{clEnqueueNDRangeKernel}
1403 event enqueue_nd_range_kernel(const kernel &kernel,
1405 const size_t *global_work_offset,
1406 const size_t *global_work_size,
1407 const size_t *local_work_size,
1408 const wait_list &events = wait_list())
1410 BOOST_ASSERT(m_queue != 0);
1411 BOOST_ASSERT(kernel.get_context() == this->get_context());
1415 cl_int ret = clEnqueueNDRangeKernel(
1418 static_cast<cl_uint>(work_dim),
1423 events.get_event_ptr(),
1427 if(ret != CL_SUCCESS){
1428 BOOST_THROW_EXCEPTION(opencl_error(ret));
1436 event enqueue_nd_range_kernel(const kernel &kernel,
1437 const extents<N> &global_work_offset,
1438 const extents<N> &global_work_size,
1439 const extents<N> &local_work_size,
1440 const wait_list &events = wait_list())
1442 return enqueue_nd_range_kernel(
1445 global_work_offset.data(),
1446 global_work_size.data(),
1447 local_work_size.data(),
1452 /// Convenience method which calls enqueue_nd_range_kernel() with a
1453 /// one-dimensional range.
1454 event enqueue_1d_range_kernel(const kernel &kernel,
1455 size_t global_work_offset,
1456 size_t global_work_size,
1457 size_t local_work_size,
1458 const wait_list &events = wait_list())
1460 return enqueue_nd_range_kernel(
1463 &global_work_offset,
1465 local_work_size ? &local_work_size : 0,
1470 /// Enqueues a kernel to execute using a single work-item.
1472 /// \see_opencl_ref{clEnqueueTask}
1473 event enqueue_task(const kernel &kernel, const wait_list &events = wait_list())
1475 BOOST_ASSERT(m_queue != 0);
1476 BOOST_ASSERT(kernel.get_context() == this->get_context());
1480 // clEnqueueTask() was deprecated in OpenCL 2.0. In that case we
1481 // just forward to the equivalent clEnqueueNDRangeKernel() call.
1482 #ifdef BOOST_COMPUTE_CL_VERSION_2_0
1484 cl_int ret = clEnqueueNDRangeKernel(
1485 m_queue, kernel, 1, 0, &one, &one,
1486 events.size(), events.get_event_ptr(), &event_.get()
1489 cl_int ret = clEnqueueTask(
1490 m_queue, kernel, events.size(), events.get_event_ptr(), &event_.get()
1494 if(ret != CL_SUCCESS){
1495 BOOST_THROW_EXCEPTION(opencl_error(ret));
1501 /// Enqueues a function to execute on the host.
1502 event enqueue_native_kernel(void (BOOST_COMPUTE_CL_CALLBACK *user_func)(void *),
1505 uint_ num_mem_objects,
1506 const cl_mem *mem_list,
1507 const void **args_mem_loc,
1508 const wait_list &events = wait_list())
1510 BOOST_ASSERT(m_queue != 0);
1513 cl_int ret = clEnqueueNativeKernel(
1522 events.get_event_ptr(),
1525 if(ret != CL_SUCCESS){
1526 BOOST_THROW_EXCEPTION(opencl_error(ret));
1532 /// Convenience overload for enqueue_native_kernel() which enqueues a
1533 /// native kernel on the host with a nullary function.
1534 event enqueue_native_kernel(void (BOOST_COMPUTE_CL_CALLBACK *user_func)(void),
1535 const wait_list &events = wait_list())
1537 return enqueue_native_kernel(
1538 detail::nullary_native_kernel_trampoline,
1539 reinterpret_cast<void *>(&user_func),
1548 /// Flushes the command queue.
1550 /// \see_opencl_ref{clFlush}
1553 BOOST_ASSERT(m_queue != 0);
1555 cl_int ret = clFlush(m_queue);
1556 if(ret != CL_SUCCESS){
1557 BOOST_THROW_EXCEPTION(opencl_error(ret));
1561 /// Blocks until all outstanding commands in the queue have finished.
1563 /// \see_opencl_ref{clFinish}
1566 BOOST_ASSERT(m_queue != 0);
1568 cl_int ret = clFinish(m_queue);
1569 if(ret != CL_SUCCESS){
1570 BOOST_THROW_EXCEPTION(opencl_error(ret));
1574 /// Enqueues a barrier in the queue.
1575 void enqueue_barrier()
1577 BOOST_ASSERT(m_queue != 0);
1578 cl_int ret = CL_SUCCESS;
1580 #ifdef BOOST_COMPUTE_CL_VERSION_1_2
1581 if(get_device().check_version(1, 2)){
1582 ret = clEnqueueBarrierWithWaitList(m_queue, 0, 0, 0);
1584 #endif // BOOST_COMPUTE_CL_VERSION_1_2
1586 // Suppress deprecated declarations warning
1587 BOOST_COMPUTE_DISABLE_DEPRECATED_DECLARATIONS();
1588 ret = clEnqueueBarrier(m_queue);
1589 BOOST_COMPUTE_ENABLE_DEPRECATED_DECLARATIONS();
1592 if(ret != CL_SUCCESS){
1593 BOOST_THROW_EXCEPTION(opencl_error(ret));
1597 #if defined(BOOST_COMPUTE_CL_VERSION_1_2) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1598 /// Enqueues a barrier in the queue after \p events.
1600 /// \opencl_version_warning{1,2}
1601 event enqueue_barrier(const wait_list &events)
1603 BOOST_ASSERT(m_queue != 0);
1606 cl_int ret = CL_SUCCESS;
1608 ret = clEnqueueBarrierWithWaitList(
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 // BOOST_COMPUTE_CL_VERSION_1_2
1620 /// Enqueues a marker in the queue and returns an event that can be
1621 /// used to track its progress.
1622 event enqueue_marker()
1625 cl_int ret = CL_SUCCESS;
1627 #ifdef BOOST_COMPUTE_CL_VERSION_1_2
1628 if(get_device().check_version(1, 2)){
1629 ret = clEnqueueMarkerWithWaitList(m_queue, 0, 0, &event_.get());
1633 // Suppress deprecated declarations warning
1634 BOOST_COMPUTE_DISABLE_DEPRECATED_DECLARATIONS();
1635 ret = clEnqueueMarker(m_queue, &event_.get());
1636 BOOST_COMPUTE_ENABLE_DEPRECATED_DECLARATIONS();
1639 if(ret != CL_SUCCESS){
1640 BOOST_THROW_EXCEPTION(opencl_error(ret));
1646 #if defined(BOOST_COMPUTE_CL_VERSION_1_2) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1647 /// Enqueues a marker after \p events in the queue and returns an
1648 /// event that can be used to track its progress.
1650 /// \opencl_version_warning{1,2}
1651 event enqueue_marker(const wait_list &events)
1655 cl_int ret = clEnqueueMarkerWithWaitList(
1656 m_queue, events.size(), events.get_event_ptr(), &event_.get()
1659 if(ret != CL_SUCCESS){
1660 BOOST_THROW_EXCEPTION(opencl_error(ret));
1665 #endif // BOOST_COMPUTE_CL_VERSION_1_2
1667 #if defined(BOOST_COMPUTE_CL_VERSION_2_0) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1668 /// Enqueues a command to copy \p size bytes of data from \p src_ptr to
1671 /// \opencl_version_warning{2,0}
1673 /// \see_opencl2_ref{clEnqueueSVMMemcpy}
1674 event enqueue_svm_memcpy(void *dst_ptr,
1675 const void *src_ptr,
1677 const wait_list &events = wait_list())
1681 cl_int ret = clEnqueueSVMMemcpy(
1688 events.get_event_ptr(),
1692 if(ret != CL_SUCCESS){
1693 BOOST_THROW_EXCEPTION(opencl_error(ret));
1699 /// Enqueues a command to copy \p size bytes of data from \p src_ptr to
1700 /// \p dst_ptr. The operation is performed asynchronously.
1702 /// \opencl_version_warning{2,0}
1704 /// \see_opencl2_ref{clEnqueueSVMMemcpy}
1705 event enqueue_svm_memcpy_async(void *dst_ptr,
1706 const void *src_ptr,
1708 const wait_list &events = wait_list())
1712 cl_int ret = clEnqueueSVMMemcpy(
1719 events.get_event_ptr(),
1723 if(ret != CL_SUCCESS){
1724 BOOST_THROW_EXCEPTION(opencl_error(ret));
1730 /// Enqueues a command to fill \p size bytes of data at \p svm_ptr with
1733 /// \opencl_version_warning{2,0}
1735 /// \see_opencl2_ref{clEnqueueSVMMemFill}
1736 event enqueue_svm_fill(void *svm_ptr,
1737 const void *pattern,
1738 size_t pattern_size,
1740 const wait_list &events = wait_list())
1745 cl_int ret = clEnqueueSVMMemFill(
1752 events.get_event_ptr(),
1756 if(ret != CL_SUCCESS){
1757 BOOST_THROW_EXCEPTION(opencl_error(ret));
1763 /// Enqueues a command to free \p svm_ptr.
1765 /// \opencl_version_warning{2,0}
1767 /// \see_opencl2_ref{clEnqueueSVMFree}
1770 event enqueue_svm_free(void *svm_ptr,
1771 const wait_list &events = wait_list())
1775 cl_int ret = clEnqueueSVMFree(
1782 events.get_event_ptr(),
1786 if(ret != CL_SUCCESS){
1787 BOOST_THROW_EXCEPTION(opencl_error(ret));
1793 /// Enqueues a command to map \p svm_ptr to the host memory space.
1795 /// \opencl_version_warning{2,0}
1797 /// \see_opencl2_ref{clEnqueueSVMMap}
1798 event enqueue_svm_map(void *svm_ptr,
1801 const wait_list &events = wait_list())
1805 cl_int ret = clEnqueueSVMMap(
1812 events.get_event_ptr(),
1816 if(ret != CL_SUCCESS){
1817 BOOST_THROW_EXCEPTION(opencl_error(ret));
1823 /// Enqueues a command to unmap \p svm_ptr from the host memory space.
1825 /// \opencl_version_warning{2,0}
1827 /// \see_opencl2_ref{clEnqueueSVMUnmap}
1828 event enqueue_svm_unmap(void *svm_ptr,
1829 const wait_list &events = wait_list())
1833 cl_int ret = clEnqueueSVMUnmap(
1837 events.get_event_ptr(),
1841 if(ret != CL_SUCCESS){
1842 BOOST_THROW_EXCEPTION(opencl_error(ret));
1847 #endif // BOOST_COMPUTE_CL_VERSION_2_0
1849 #if defined(BOOST_COMPUTE_CL_VERSION_2_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1850 /// Enqueues a command to indicate which device a set of ranges of SVM allocations
1851 /// should be associated with. The pair \p svm_ptrs[i] and \p sizes[i] together define
1852 /// the starting address and number of bytes in a range to be migrated.
1854 /// If \p sizes is empty, then that means every allocation containing any \p svm_ptrs[i]
1855 /// is to be migrated. Also, if \p sizes[i] is zero, then the entire allocation containing
1856 /// \p svm_ptrs[i] is migrated.
1858 /// \opencl_version_warning{2,1}
1860 /// \see_opencl21_ref{clEnqueueSVMMigrateMem}
1861 event enqueue_svm_migrate_memory(const std::vector<const void*> &svm_ptrs,
1862 const std::vector<size_t> &sizes,
1863 const cl_mem_migration_flags flags = 0,
1864 const wait_list &events = wait_list())
1866 BOOST_ASSERT(svm_ptrs.size() == sizes.size() || sizes.size() == 0);
1869 cl_int ret = clEnqueueSVMMigrateMem(
1871 static_cast<cl_uint>(svm_ptrs.size()),
1872 const_cast<void const **>(&svm_ptrs[0]),
1873 sizes.size() > 0 ? &sizes[0] : NULL,
1876 events.get_event_ptr(),
1880 if(ret != CL_SUCCESS){
1881 BOOST_THROW_EXCEPTION(opencl_error(ret));
1887 /// Enqueues a command to indicate which device a range of SVM allocation
1888 /// should be associated with. The pair \p svm_ptr and \p size together define
1889 /// the starting address and number of bytes in a range to be migrated.
1891 /// If \p size is 0, then the entire allocation containing \p svm_ptr is
1892 /// migrated. The default value for \p size is 0.
1894 /// \opencl_version_warning{2,1}
1896 /// \see_opencl21_ref{clEnqueueSVMMigrateMem}
1897 event enqueue_svm_migrate_memory(const void* svm_ptr,
1898 const size_t size = 0,
1899 const cl_mem_migration_flags flags = 0,
1900 const wait_list &events = wait_list())
1904 cl_int ret = clEnqueueSVMMigrateMem(
1911 events.get_event_ptr(),
1915 if(ret != CL_SUCCESS){
1916 BOOST_THROW_EXCEPTION(opencl_error(ret));
1921 #endif // BOOST_COMPUTE_CL_VERSION_2_1
1923 /// Returns \c true if the command queue is the same at \p other.
1924 bool operator==(const command_queue &other) const
1926 return m_queue == other.m_queue;
1929 /// Returns \c true if the command queue is different from \p other.
1930 bool operator!=(const command_queue &other) const
1932 return m_queue != other.m_queue;
1936 operator cl_command_queue() const
1942 bool check_device_version(int major, int minor) const
1944 return get_device().check_version(major, minor);
1948 cl_command_queue m_queue;
1951 inline buffer buffer::clone(command_queue &queue) const
1953 buffer copy(get_context(), size(), get_memory_flags());
1954 queue.enqueue_copy_buffer(*this, copy, 0, 0, size());
1958 inline image1d image1d::clone(command_queue &queue) const
1961 get_context(), width(), format(), get_memory_flags()
1964 queue.enqueue_copy_image(*this, copy, origin(), copy.origin(), size());
1969 inline image2d image2d::clone(command_queue &queue) const
1972 get_context(), width(), height(), format(), get_memory_flags()
1975 queue.enqueue_copy_image(*this, copy, origin(), copy.origin(), size());
1980 inline image3d image3d::clone(command_queue &queue) const
1983 get_context(), width(), height(), depth(), format(), get_memory_flags()
1986 queue.enqueue_copy_image(*this, copy, origin(), copy.origin(), size());
1991 /// \internal_ define get_info() specializations for command_queue
1992 BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_SPECIALIZATIONS(command_queue,
1993 ((cl_context, CL_QUEUE_CONTEXT))
1994 ((cl_device_id, CL_QUEUE_DEVICE))
1995 ((uint_, CL_QUEUE_REFERENCE_COUNT))
1996 ((cl_command_queue_properties, CL_QUEUE_PROPERTIES))
1999 #ifdef BOOST_COMPUTE_CL_VERSION_2_1
2000 BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_SPECIALIZATIONS(command_queue,
2001 ((cl_command_queue, CL_QUEUE_DEVICE_DEFAULT))
2003 #endif // BOOST_COMPUTE_CL_VERSION_2_1
2005 } // end compute namespace
2006 } // end boost namespace
2008 #endif // BOOST_COMPUTE_COMMAND_QUEUE_HPP