353 lines
11 KiB
C++
353 lines
11 KiB
C++
//---------------------------------------------------------------------------//
|
|
// 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()
|