123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156 |
- //---------------------------------------------------------------------------//
- // Copyright (c) 2013-2014 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 TestSvmPtr
- #include <boost/test/unit_test.hpp>
- #include <iostream>
- #include <boost/compute/core.hpp>
- #include <boost/compute/svm.hpp>
- #include <boost/compute/container/vector.hpp>
- #include <boost/compute/utility/source.hpp>
- #include "quirks.hpp"
- #include "check_macros.hpp"
- #include "context_setup.hpp"
- namespace compute = boost::compute;
- BOOST_AUTO_TEST_CASE(empty)
- {
- }
- #ifdef BOOST_COMPUTE_CL_VERSION_2_0
- BOOST_AUTO_TEST_CASE(alloc)
- {
- REQUIRES_OPENCL_VERSION(2, 0);
- compute::svm_ptr<cl_int> ptr = compute::svm_alloc<cl_int>(context, 8);
- compute::svm_free(context, ptr);
- }
- BOOST_AUTO_TEST_CASE(svmmemcpy)
- {
- REQUIRES_OPENCL_VERSION(2, 0);
- if(bug_in_svmmemcpy(device)){
- std::cerr << "skipping svmmemcpy test case" << std::endl;
- return;
- }
- cl_int input[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
- cl_int output[] = { 0, 0, 0, 0, 0, 0, 0, 0 };
- compute::svm_ptr<cl_int> ptr = compute::svm_alloc<cl_int>(context, 8);
- compute::svm_ptr<cl_int> ptr2 = compute::svm_alloc<cl_int>(context, 8);
- // copying from and to host mem
- queue.enqueue_svm_memcpy(ptr.get(), input, 8 * sizeof(cl_int));
- queue.enqueue_svm_memcpy(output, ptr.get(), 8 * sizeof(cl_int));
- queue.finish();
- CHECK_HOST_RANGE_EQUAL(cl_int, 8, output, (1, 2, 3, 4, 5, 6, 7, 8));
- // copying between svm mem
- queue.enqueue_svm_memcpy(ptr2.get(), ptr.get(), 8 * sizeof(cl_int));
- queue.enqueue_svm_memcpy(output, ptr2.get(), 8 * sizeof(cl_int));
- queue.finish();
- CHECK_HOST_RANGE_EQUAL(cl_int, 8, output, (1, 2, 3, 4, 5, 6, 7, 8));
- compute::svm_free(context, ptr);
- compute::svm_free(context, ptr2);
- }
- BOOST_AUTO_TEST_CASE(sum_svm_kernel)
- {
- REQUIRES_OPENCL_VERSION(2, 0);
- const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
- __kernel void sum_svm_mem(__global const int *ptr, __global int *result)
- {
- int sum = 0;
- for(uint i = 0; i < 8; i++){
- sum += ptr[i];
- }
- *result = sum;
- }
- );
- compute::program program =
- compute::program::build_with_source(source, context, "-cl-std=CL2.0");
- compute::kernel sum_svm_mem_kernel = program.create_kernel("sum_svm_mem");
- cl_int data[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
- compute::svm_ptr<cl_int> ptr = compute::svm_alloc<cl_int>(context, 8);
- queue.enqueue_svm_map(ptr.get(), 8 * sizeof(cl_int), CL_MAP_WRITE);
- for(size_t i = 0; i < 8; i ++) {
- static_cast<cl_int*>(ptr.get())[i] = data[i];
- }
- queue.enqueue_svm_unmap(ptr.get());
- compute::vector<cl_int> result(1, context);
- sum_svm_mem_kernel.set_arg(0, ptr);
- sum_svm_mem_kernel.set_arg(1, result);
- queue.enqueue_task(sum_svm_mem_kernel);
- queue.finish();
- BOOST_CHECK_EQUAL(result[0], (36));
- compute::svm_free(context, ptr);
- }
- #endif // BOOST_COMPUTE_CL_VERSION_2_0
- #ifdef BOOST_COMPUTE_CL_VERSION_2_1
- BOOST_AUTO_TEST_CASE(migrate)
- {
- REQUIRES_OPENCL_VERSION(2, 1);
- compute::svm_ptr<cl_int> ptr =
- compute::svm_alloc<cl_int>(context, 8);
- // Migrate to device
- std::vector<const void*> ptrs(1, ptr.get());
- std::vector<size_t> sizes(1, 8 * sizeof(cl_int));
- queue.enqueue_svm_migrate_memory(ptrs, sizes).wait();
- // Set on device
- const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
- __kernel void foo(__global int *ptr)
- {
- for(int i = 0; i < 8; i++){
- ptr[i] = i;
- }
- }
- );
- compute::program program =
- compute::program::build_with_source(source, context, "-cl-std=CL2.0");
- compute::kernel foo_kernel = program.create_kernel("foo");
- foo_kernel.set_arg(0, ptr);
- queue.enqueue_task(foo_kernel).wait();
- // Migrate to host
- queue.enqueue_svm_migrate_memory(
- ptr.get(), 0, boost::compute::command_queue::migrate_to_host
- ).wait();
- // Check
- CHECK_HOST_RANGE_EQUAL(
- cl_int, 8,
- static_cast<cl_int*>(ptr.get()),
- (0, 1, 2, 3, 4, 5, 6, 7)
- );
- compute::svm_free(context, ptr);
- }
- #endif // BOOST_COMPUTE_CL_VERSION_2_1
- BOOST_AUTO_TEST_SUITE_END()
|