123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352 |
- //---------------------------------------------------------------------------//
- // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
- //
- // Distributed under the Boost Software License, Version 1.0
- // See accompanying file LICENSE_1_0.txt or copy at
- // http://www.boost.org/LICENSE_1_0.txt
- //
- // See http://boostorg.github.com/compute for more information.
- //---------------------------------------------------------------------------//
- #define BOOST_TEST_MODULE TestCommandQueue
- #include <boost/test/unit_test.hpp>
- #include <iostream>
- #include <boost/compute/kernel.hpp>
- #include <boost/compute/system.hpp>
- #include <boost/compute/program.hpp>
- #include <boost/compute/command_queue.hpp>
- #include <boost/compute/algorithm/fill.hpp>
- #include <boost/compute/container/vector.hpp>
- #include <boost/compute/utility/dim.hpp>
- #include <boost/compute/utility/source.hpp>
- #include <boost/compute/detail/diagnostic.hpp>
- #include "check_macros.hpp"
- #include "context_setup.hpp"
- namespace bc = boost::compute;
- namespace compute = boost::compute;
- BOOST_AUTO_TEST_CASE(get_context)
- {
- BOOST_VERIFY(queue.get_context() == context);
- BOOST_VERIFY(queue.get_info<CL_QUEUE_CONTEXT>() == context.get());
- }
- BOOST_AUTO_TEST_CASE(get_device)
- {
- BOOST_VERIFY(queue.get_info<CL_QUEUE_DEVICE>() == device.get());
- }
- BOOST_AUTO_TEST_CASE(equality_operator)
- {
- compute::command_queue queue1(context, device);
- BOOST_CHECK(queue1 == queue1);
- compute::command_queue queue2 = queue1;
- BOOST_CHECK(queue1 == queue2);
- compute::command_queue queue3(context, device);
- BOOST_CHECK(queue1 != queue3);
- }
- BOOST_AUTO_TEST_CASE(event_profiling)
- {
- bc::command_queue queue(context, device, bc::command_queue::enable_profiling);
- int data[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
- bc::buffer buffer(context, sizeof(data));
- bc::event event =
- queue.enqueue_write_buffer_async(buffer,
- 0,
- sizeof(data),
- static_cast<const void *>(data));
- queue.finish();
- event.get_profiling_info<cl_ulong>(bc::event::profiling_command_queued);
- event.get_profiling_info<cl_ulong>(bc::event::profiling_command_submit);
- event.get_profiling_info<cl_ulong>(bc::event::profiling_command_start);
- event.get_profiling_info<cl_ulong>(bc::event::profiling_command_end);
- }
- BOOST_AUTO_TEST_CASE(kernel_profiling)
- {
- // create queue with profiling enabled
- boost::compute::command_queue queue(
- context, device, boost::compute::command_queue::enable_profiling
- );
- // input data
- int data[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
- boost::compute::buffer buffer(context, sizeof(data));
- // copy input data to device
- queue.enqueue_write_buffer(buffer, 0, sizeof(data), data);
- // setup kernel
- const char source[] =
- "__kernel void iscal(__global int *buffer, int alpha)\n"
- "{\n"
- " buffer[get_global_id(0)] *= alpha;\n"
- "}\n";
- boost::compute::program program =
- boost::compute::program::create_with_source(source, context);
- program.build();
- boost::compute::kernel kernel(program, "iscal");
- kernel.set_arg(0, buffer);
- kernel.set_arg(1, 2);
- // execute kernel
- size_t global_work_offset = 0;
- size_t global_work_size = 8;
- boost::compute::event event =
- queue.enqueue_nd_range_kernel(kernel,
- size_t(1),
- &global_work_offset,
- &global_work_size,
- 0);
- // wait until kernel is finished
- event.wait();
- // check profiling information
- event.get_profiling_info<cl_ulong>(bc::event::profiling_command_queued);
- event.get_profiling_info<cl_ulong>(bc::event::profiling_command_submit);
- event.get_profiling_info<cl_ulong>(bc::event::profiling_command_start);
- event.get_profiling_info<cl_ulong>(bc::event::profiling_command_end);
- // read results back to host
- queue.enqueue_read_buffer(buffer, 0, sizeof(data), data);
- // check results
- BOOST_CHECK_EQUAL(data[0], 2);
- BOOST_CHECK_EQUAL(data[1], 4);
- BOOST_CHECK_EQUAL(data[2], 6);
- BOOST_CHECK_EQUAL(data[3], 8);
- BOOST_CHECK_EQUAL(data[4], 10);
- BOOST_CHECK_EQUAL(data[5], 12);
- BOOST_CHECK_EQUAL(data[6], 14);
- BOOST_CHECK_EQUAL(data[7], 16);
- }
- BOOST_AUTO_TEST_CASE(construct_from_cl_command_queue)
- {
- // create cl_command_queue
- cl_command_queue cl_queue;
- #ifdef BOOST_COMPUTE_CL_VERSION_2_0
- if (device.check_version(2, 0)){ // runtime check
- cl_queue =
- clCreateCommandQueueWithProperties(context, device.id(), 0, 0);
- } else
- #endif // BOOST_COMPUTE_CL_VERSION_2_0
- {
- // Suppress deprecated declarations warning
- BOOST_COMPUTE_DISABLE_DEPRECATED_DECLARATIONS();
- cl_queue =
- clCreateCommandQueue(context, device.id(), 0, 0);
- BOOST_COMPUTE_ENABLE_DEPRECATED_DECLARATIONS();
- }
- BOOST_VERIFY(cl_queue);
- // create boost::compute::command_queue
- boost::compute::command_queue queue(cl_queue);
- // check queue
- BOOST_CHECK(queue.get_context() == context);
- BOOST_CHECK(cl_command_queue(queue) == cl_queue);
- // cleanup cl_command_queue
- clReleaseCommandQueue(cl_queue);
- }
- #ifdef BOOST_COMPUTE_CL_VERSION_1_1
- BOOST_AUTO_TEST_CASE(write_buffer_rect)
- {
- REQUIRES_OPENCL_VERSION(1, 1);
- // skip this test on AMD GPUs due to a buggy implementation
- // of the clEnqueueWriteBufferRect() function
- if(device.vendor() == "Advanced Micro Devices, Inc." &&
- device.type() & boost::compute::device::gpu){
- std::cerr << "skipping write_buffer_rect test on AMD GPU" << std::endl;
- return;
- }
- int data[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
- boost::compute::buffer buffer(context, 8 * sizeof(int));
- // copy every other value to the buffer
- size_t buffer_origin[] = { 0, 0, 0 };
- size_t host_origin[] = { 0, 0, 0 };
- size_t region[] = { sizeof(int), sizeof(int), 1 };
- queue.enqueue_write_buffer_rect(
- buffer,
- buffer_origin,
- host_origin,
- region,
- sizeof(int),
- 0,
- 2 * sizeof(int),
- 0,
- data
- );
- // check output values
- int output[4];
- queue.enqueue_read_buffer(buffer, 0, 4 * sizeof(int), output);
- BOOST_CHECK_EQUAL(output[0], 1);
- BOOST_CHECK_EQUAL(output[1], 3);
- BOOST_CHECK_EQUAL(output[2], 5);
- BOOST_CHECK_EQUAL(output[3], 7);
- }
- #endif // BOOST_COMPUTE_CL_VERSION_1_1
- static bool nullary_kernel_executed = false;
- static void nullary_kernel()
- {
- nullary_kernel_executed = true;
- }
- BOOST_AUTO_TEST_CASE(native_kernel)
- {
- cl_device_exec_capabilities exec_capabilities =
- device.get_info<CL_DEVICE_EXECUTION_CAPABILITIES>();
- if(!(exec_capabilities & CL_EXEC_NATIVE_KERNEL)){
- std::cerr << "skipping native_kernel test: "
- << "device does not support CL_EXEC_NATIVE_KERNEL"
- << std::endl;
- return;
- }
- compute::vector<int> vector(1000, context);
- compute::fill(vector.begin(), vector.end(), 42, queue);
- BOOST_CHECK_EQUAL(nullary_kernel_executed, false);
- queue.enqueue_native_kernel(&nullary_kernel);
- queue.finish();
- BOOST_CHECK_EQUAL(nullary_kernel_executed, true);
- }
- BOOST_AUTO_TEST_CASE(copy_with_wait_list)
- {
- int data1[] = { 1, 3, 5, 7 };
- int data2[] = { 2, 4, 6, 8 };
- compute::buffer buf1(context, 4 * sizeof(int));
- compute::buffer buf2(context, 4 * sizeof(int));
- compute::event write_event1 =
- queue.enqueue_write_buffer_async(buf1, 0, buf1.size(), data1);
- compute::event write_event2 =
- queue.enqueue_write_buffer_async(buf2, 0, buf2.size(), data2);
- compute::event read_event1 =
- queue.enqueue_read_buffer_async(buf1, 0, buf1.size(), data2, write_event1);
- compute::event read_event2 =
- queue.enqueue_read_buffer_async(buf2, 0, buf2.size(), data1, write_event2);
- read_event1.wait();
- read_event2.wait();
- CHECK_HOST_RANGE_EQUAL(int, 4, data1, (2, 4, 6, 8));
- CHECK_HOST_RANGE_EQUAL(int, 4, data2, (1, 3, 5, 7));
- }
- #ifndef BOOST_COMPUTE_NO_HDR_INITIALIZER_LIST
- BOOST_AUTO_TEST_CASE(enqueue_kernel_with_extents)
- {
- using boost::compute::dim;
- using boost::compute::uint_;
- const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
- __kernel void foo(__global int *output1, __global int *output2)
- {
- output1[get_global_id(0)] = get_local_id(0);
- output2[get_global_id(1)] = get_local_id(1);
- }
- );
- compute::kernel kernel =
- compute::kernel::create_with_source(source, "foo", context);
- compute::vector<uint_> output1(4, context);
- compute::vector<uint_> output2(4, context);
- kernel.set_arg(0, output1);
- kernel.set_arg(1, output2);
- queue.enqueue_nd_range_kernel(kernel, dim(0, 0), dim(4, 4), dim(1, 1));
- CHECK_RANGE_EQUAL(int, 4, output1, (0, 0, 0, 0));
- CHECK_RANGE_EQUAL(int, 4, output2, (0, 0, 0, 0));
- // Maximum number of work-items that can be specified in each
- // dimension of the work-group to clEnqueueNDRangeKernel.
- std::vector<size_t> max_work_item_sizes =
- device.get_info<CL_DEVICE_MAX_WORK_ITEM_SIZES>();
- if(max_work_item_sizes[0] < size_t(2)) {
- return;
- }
- queue.enqueue_nd_range_kernel(kernel, dim(0, 0), dim(4, 4), dim(2, 1));
- CHECK_RANGE_EQUAL(int, 4, output1, (0, 1, 0, 1));
- CHECK_RANGE_EQUAL(int, 4, output2, (0, 0, 0, 0));
- if(max_work_item_sizes[1] < size_t(2)) {
- return;
- }
- queue.enqueue_nd_range_kernel(kernel, dim(0, 0), dim(4, 4), dim(2, 2));
- CHECK_RANGE_EQUAL(int, 4, output1, (0, 1, 0, 1));
- CHECK_RANGE_EQUAL(int, 4, output2, (0, 1, 0, 1));
- }
- #endif // BOOST_COMPUTE_NO_HDR_INITIALIZER_LIST
- #ifdef BOOST_COMPUTE_CL_VERSION_2_1
- BOOST_AUTO_TEST_CASE(get_default_device_queue)
- {
- REQUIRES_OPENCL_VERSION(2, 1);
- boost::compute::command_queue default_device_queue(
- context, device,
- boost::compute::command_queue::on_device |
- boost::compute::command_queue::on_device_default |
- boost::compute::command_queue::enable_out_of_order_execution
- );
- BOOST_CHECK_NO_THROW(queue.get_info<CL_QUEUE_DEVICE_DEFAULT>());
- BOOST_CHECK_EQUAL(
- queue.get_default_device_queue(),
- default_device_queue
- );
- }
- BOOST_AUTO_TEST_CASE(set_as_default_device_queue)
- {
- REQUIRES_OPENCL_VERSION(2, 1);
- boost::compute::command_queue new_default_device_queue(
- context, device,
- boost::compute::command_queue::on_device |
- boost::compute::command_queue::enable_out_of_order_execution
- );
- new_default_device_queue.set_as_default_device_queue();
- BOOST_CHECK_EQUAL(
- queue.get_default_device_queue(),
- new_default_device_queue
- );
- }
- #endif
- BOOST_AUTO_TEST_SUITE_END()
|