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 #define BOOST_TEST_MODULE TestCommandQueue
12 #include <boost/test/unit_test.hpp>
16 #include <boost/compute/kernel.hpp>
17 #include <boost/compute/system.hpp>
18 #include <boost/compute/program.hpp>
19 #include <boost/compute/command_queue.hpp>
20 #include <boost/compute/algorithm/fill.hpp>
21 #include <boost/compute/container/vector.hpp>
22 #include <boost/compute/utility/dim.hpp>
23 #include <boost/compute/utility/source.hpp>
24 #include <boost/compute/detail/diagnostic.hpp>
26 #include "check_macros.hpp"
27 #include "context_setup.hpp"
29 namespace bc
= boost::compute
;
30 namespace compute
= boost::compute
;
32 BOOST_AUTO_TEST_CASE(get_context
)
34 BOOST_VERIFY(queue
.get_context() == context
);
35 BOOST_VERIFY(queue
.get_info
<CL_QUEUE_CONTEXT
>() == context
.get());
38 BOOST_AUTO_TEST_CASE(get_device
)
40 BOOST_VERIFY(queue
.get_info
<CL_QUEUE_DEVICE
>() == device
.get());
43 BOOST_AUTO_TEST_CASE(equality_operator
)
45 compute::command_queue
queue1(context
, device
);
46 BOOST_CHECK(queue1
== queue1
);
48 compute::command_queue queue2
= queue1
;
49 BOOST_CHECK(queue1
== queue2
);
51 compute::command_queue
queue3(context
, device
);
52 BOOST_CHECK(queue1
!= queue3
);
55 BOOST_AUTO_TEST_CASE(event_profiling
)
57 bc::command_queue
queue(context
, device
, bc::command_queue::enable_profiling
);
59 int data
[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
60 bc::buffer
buffer(context
, sizeof(data
));
63 queue
.enqueue_write_buffer_async(buffer
,
66 static_cast<const void *>(data
));
69 event
.get_profiling_info
<cl_ulong
>(bc::event::profiling_command_queued
);
70 event
.get_profiling_info
<cl_ulong
>(bc::event::profiling_command_submit
);
71 event
.get_profiling_info
<cl_ulong
>(bc::event::profiling_command_start
);
72 event
.get_profiling_info
<cl_ulong
>(bc::event::profiling_command_end
);
75 BOOST_AUTO_TEST_CASE(kernel_profiling
)
77 // create queue with profiling enabled
78 boost::compute::command_queue
queue(
79 context
, device
, boost::compute::command_queue::enable_profiling
83 int data
[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
84 boost::compute::buffer
buffer(context
, sizeof(data
));
86 // copy input data to device
87 queue
.enqueue_write_buffer(buffer
, 0, sizeof(data
), data
);
91 "__kernel void iscal(__global int *buffer, int alpha)\n"
93 " buffer[get_global_id(0)] *= alpha;\n"
96 boost::compute::program program
=
97 boost::compute::program::create_with_source(source
, context
);
100 boost::compute::kernel
kernel(program
, "iscal");
101 kernel
.set_arg(0, buffer
);
102 kernel
.set_arg(1, 2);
105 size_t global_work_offset
= 0;
106 size_t global_work_size
= 8;
108 boost::compute::event event
=
109 queue
.enqueue_nd_range_kernel(kernel
,
115 // wait until kernel is finished
118 // check profiling information
119 event
.get_profiling_info
<cl_ulong
>(bc::event::profiling_command_queued
);
120 event
.get_profiling_info
<cl_ulong
>(bc::event::profiling_command_submit
);
121 event
.get_profiling_info
<cl_ulong
>(bc::event::profiling_command_start
);
122 event
.get_profiling_info
<cl_ulong
>(bc::event::profiling_command_end
);
124 // read results back to host
125 queue
.enqueue_read_buffer(buffer
, 0, sizeof(data
), data
);
128 BOOST_CHECK_EQUAL(data
[0], 2);
129 BOOST_CHECK_EQUAL(data
[1], 4);
130 BOOST_CHECK_EQUAL(data
[2], 6);
131 BOOST_CHECK_EQUAL(data
[3], 8);
132 BOOST_CHECK_EQUAL(data
[4], 10);
133 BOOST_CHECK_EQUAL(data
[5], 12);
134 BOOST_CHECK_EQUAL(data
[6], 14);
135 BOOST_CHECK_EQUAL(data
[7], 16);
138 BOOST_AUTO_TEST_CASE(construct_from_cl_command_queue
)
140 // create cl_command_queue
141 cl_command_queue cl_queue
;
142 #ifdef BOOST_COMPUTE_CL_VERSION_2_0
143 if (device
.check_version(2, 0)){ // runtime check
145 clCreateCommandQueueWithProperties(context
, device
.id(), 0, 0);
147 #endif // BOOST_COMPUTE_CL_VERSION_2_0
149 // Suppress deprecated declarations warning
150 BOOST_COMPUTE_DISABLE_DEPRECATED_DECLARATIONS();
152 clCreateCommandQueue(context
, device
.id(), 0, 0);
153 BOOST_COMPUTE_ENABLE_DEPRECATED_DECLARATIONS();
155 BOOST_VERIFY(cl_queue
);
157 // create boost::compute::command_queue
158 boost::compute::command_queue
queue(cl_queue
);
161 BOOST_CHECK(queue
.get_context() == context
);
162 BOOST_CHECK(cl_command_queue(queue
) == cl_queue
);
164 // cleanup cl_command_queue
165 clReleaseCommandQueue(cl_queue
);
168 #ifdef BOOST_COMPUTE_CL_VERSION_1_1
169 BOOST_AUTO_TEST_CASE(write_buffer_rect
)
171 REQUIRES_OPENCL_VERSION(1, 1);
173 // skip this test on AMD GPUs due to a buggy implementation
174 // of the clEnqueueWriteBufferRect() function
175 if(device
.vendor() == "Advanced Micro Devices, Inc." &&
176 device
.type() & boost::compute::device::gpu
){
177 std::cerr
<< "skipping write_buffer_rect test on AMD GPU" << std::endl
;
181 int data
[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
182 boost::compute::buffer
buffer(context
, 8 * sizeof(int));
184 // copy every other value to the buffer
185 size_t buffer_origin
[] = { 0, 0, 0 };
186 size_t host_origin
[] = { 0, 0, 0 };
187 size_t region
[] = { sizeof(int), sizeof(int), 1 };
189 queue
.enqueue_write_buffer_rect(
201 // check output values
203 queue
.enqueue_read_buffer(buffer
, 0, 4 * sizeof(int), output
);
204 BOOST_CHECK_EQUAL(output
[0], 1);
205 BOOST_CHECK_EQUAL(output
[1], 3);
206 BOOST_CHECK_EQUAL(output
[2], 5);
207 BOOST_CHECK_EQUAL(output
[3], 7);
209 #endif // BOOST_COMPUTE_CL_VERSION_1_1
211 static bool nullary_kernel_executed
= false;
213 static void nullary_kernel()
215 nullary_kernel_executed
= true;
218 BOOST_AUTO_TEST_CASE(native_kernel
)
220 cl_device_exec_capabilities exec_capabilities
=
221 device
.get_info
<CL_DEVICE_EXECUTION_CAPABILITIES
>();
222 if(!(exec_capabilities
& CL_EXEC_NATIVE_KERNEL
)){
223 std::cerr
<< "skipping native_kernel test: "
224 << "device does not support CL_EXEC_NATIVE_KERNEL"
229 compute::vector
<int> vector(1000, context
);
230 compute::fill(vector
.begin(), vector
.end(), 42, queue
);
231 BOOST_CHECK_EQUAL(nullary_kernel_executed
, false);
232 queue
.enqueue_native_kernel(&nullary_kernel
);
234 BOOST_CHECK_EQUAL(nullary_kernel_executed
, true);
237 BOOST_AUTO_TEST_CASE(copy_with_wait_list
)
239 int data1
[] = { 1, 3, 5, 7 };
240 int data2
[] = { 2, 4, 6, 8 };
242 compute::buffer
buf1(context
, 4 * sizeof(int));
243 compute::buffer
buf2(context
, 4 * sizeof(int));
245 compute::event write_event1
=
246 queue
.enqueue_write_buffer_async(buf1
, 0, buf1
.size(), data1
);
248 compute::event write_event2
=
249 queue
.enqueue_write_buffer_async(buf2
, 0, buf2
.size(), data2
);
251 compute::event read_event1
=
252 queue
.enqueue_read_buffer_async(buf1
, 0, buf1
.size(), data2
, write_event1
);
254 compute::event read_event2
=
255 queue
.enqueue_read_buffer_async(buf2
, 0, buf2
.size(), data1
, write_event2
);
260 CHECK_HOST_RANGE_EQUAL(int, 4, data1
, (2, 4, 6, 8));
261 CHECK_HOST_RANGE_EQUAL(int, 4, data2
, (1, 3, 5, 7));
264 #ifndef BOOST_COMPUTE_NO_HDR_INITIALIZER_LIST
265 BOOST_AUTO_TEST_CASE(enqueue_kernel_with_extents
)
267 using boost::compute::dim
;
268 using boost::compute::uint_
;
270 const char source
[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
271 __kernel
void foo(__global
int *output1
, __global
int *output2
)
273 output1
[get_global_id(0)] = get_local_id(0);
274 output2
[get_global_id(1)] = get_local_id(1);
278 compute::kernel kernel
=
279 compute::kernel::create_with_source(source
, "foo", context
);
281 compute::vector
<uint_
> output1(4, context
);
282 compute::vector
<uint_
> output2(4, context
);
284 kernel
.set_arg(0, output1
);
285 kernel
.set_arg(1, output2
);
287 queue
.enqueue_nd_range_kernel(kernel
, dim(0, 0), dim(4, 4), dim(1, 1));
289 CHECK_RANGE_EQUAL(int, 4, output1
, (0, 0, 0, 0));
290 CHECK_RANGE_EQUAL(int, 4, output2
, (0, 0, 0, 0));
292 // Maximum number of work-items that can be specified in each
293 // dimension of the work-group to clEnqueueNDRangeKernel.
294 std::vector
<size_t> max_work_item_sizes
=
295 device
.get_info
<CL_DEVICE_MAX_WORK_ITEM_SIZES
>();
297 if(max_work_item_sizes
[0] < size_t(2)) {
301 queue
.enqueue_nd_range_kernel(kernel
, dim(0, 0), dim(4, 4), dim(2, 1));
303 CHECK_RANGE_EQUAL(int, 4, output1
, (0, 1, 0, 1));
304 CHECK_RANGE_EQUAL(int, 4, output2
, (0, 0, 0, 0));
306 if(max_work_item_sizes
[1] < size_t(2)) {
310 queue
.enqueue_nd_range_kernel(kernel
, dim(0, 0), dim(4, 4), dim(2, 2));
312 CHECK_RANGE_EQUAL(int, 4, output1
, (0, 1, 0, 1));
313 CHECK_RANGE_EQUAL(int, 4, output2
, (0, 1, 0, 1));
315 #endif // BOOST_COMPUTE_NO_HDR_INITIALIZER_LIST
317 #ifdef BOOST_COMPUTE_CL_VERSION_2_1
318 BOOST_AUTO_TEST_CASE(get_default_device_queue
)
320 REQUIRES_OPENCL_VERSION(2, 1);
322 boost::compute::command_queue
default_device_queue(
324 boost::compute::command_queue::on_device
|
325 boost::compute::command_queue::on_device_default
|
326 boost::compute::command_queue::enable_out_of_order_execution
328 BOOST_CHECK_NO_THROW(queue
.get_info
<CL_QUEUE_DEVICE_DEFAULT
>());
330 queue
.get_default_device_queue(),
335 BOOST_AUTO_TEST_CASE(set_as_default_device_queue
)
337 REQUIRES_OPENCL_VERSION(2, 1);
339 boost::compute::command_queue
new_default_device_queue(
341 boost::compute::command_queue::on_device
|
342 boost::compute::command_queue::enable_out_of_order_execution
344 new_default_device_queue
.set_as_default_device_queue();
346 queue
.get_default_device_queue(),
347 new_default_device_queue
352 BOOST_AUTO_TEST_SUITE_END()