1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013-2014 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 TestSvmPtr
12 #include <boost/test/unit_test.hpp>
16 #include <boost/compute/core.hpp>
17 #include <boost/compute/svm.hpp>
18 #include <boost/compute/container/vector.hpp>
19 #include <boost/compute/utility/source.hpp>
22 #include "check_macros.hpp"
23 #include "context_setup.hpp"
25 namespace compute
= boost::compute
;
27 BOOST_AUTO_TEST_CASE(empty
)
31 #ifdef BOOST_COMPUTE_CL_VERSION_2_0
32 BOOST_AUTO_TEST_CASE(alloc
)
34 REQUIRES_OPENCL_VERSION(2, 0);
36 compute::svm_ptr
<cl_int
> ptr
= compute::svm_alloc
<cl_int
>(context
, 8);
37 compute::svm_free(context
, ptr
);
40 BOOST_AUTO_TEST_CASE(svmmemcpy
)
42 REQUIRES_OPENCL_VERSION(2, 0);
44 if(bug_in_svmmemcpy(device
)){
45 std::cerr
<< "skipping svmmemcpy test case" << std::endl
;
49 cl_int input
[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
50 cl_int output
[] = { 0, 0, 0, 0, 0, 0, 0, 0 };
51 compute::svm_ptr
<cl_int
> ptr
= compute::svm_alloc
<cl_int
>(context
, 8);
52 compute::svm_ptr
<cl_int
> ptr2
= compute::svm_alloc
<cl_int
>(context
, 8);
54 // copying from and to host mem
55 queue
.enqueue_svm_memcpy(ptr
.get(), input
, 8 * sizeof(cl_int
));
56 queue
.enqueue_svm_memcpy(output
, ptr
.get(), 8 * sizeof(cl_int
));
59 CHECK_HOST_RANGE_EQUAL(cl_int
, 8, output
, (1, 2, 3, 4, 5, 6, 7, 8));
61 // copying between svm mem
62 queue
.enqueue_svm_memcpy(ptr2
.get(), ptr
.get(), 8 * sizeof(cl_int
));
63 queue
.enqueue_svm_memcpy(output
, ptr2
.get(), 8 * sizeof(cl_int
));
66 CHECK_HOST_RANGE_EQUAL(cl_int
, 8, output
, (1, 2, 3, 4, 5, 6, 7, 8));
68 compute::svm_free(context
, ptr
);
69 compute::svm_free(context
, ptr2
);
72 BOOST_AUTO_TEST_CASE(sum_svm_kernel
)
74 REQUIRES_OPENCL_VERSION(2, 0);
76 const char source
[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
77 __kernel
void sum_svm_mem(__global
const int *ptr
, __global
int *result
)
80 for(uint i
= 0; i
< 8; i
++){
87 compute::program program
=
88 compute::program::build_with_source(source
, context
, "-cl-std=CL2.0");
90 compute::kernel sum_svm_mem_kernel
= program
.create_kernel("sum_svm_mem");
92 cl_int data
[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
93 compute::svm_ptr
<cl_int
> ptr
= compute::svm_alloc
<cl_int
>(context
, 8);
94 queue
.enqueue_svm_map(ptr
.get(), 8 * sizeof(cl_int
), CL_MAP_WRITE
);
95 for(size_t i
= 0; i
< 8; i
++) {
96 static_cast<cl_int
*>(ptr
.get())[i
] = data
[i
];
98 queue
.enqueue_svm_unmap(ptr
.get());
100 compute::vector
<cl_int
> result(1, context
);
102 sum_svm_mem_kernel
.set_arg(0, ptr
);
103 sum_svm_mem_kernel
.set_arg(1, result
);
104 queue
.enqueue_task(sum_svm_mem_kernel
);
107 BOOST_CHECK_EQUAL(result
[0], (36));
109 compute::svm_free(context
, ptr
);
111 #endif // BOOST_COMPUTE_CL_VERSION_2_0
113 #ifdef BOOST_COMPUTE_CL_VERSION_2_1
114 BOOST_AUTO_TEST_CASE(migrate
)
116 REQUIRES_OPENCL_VERSION(2, 1);
118 compute::svm_ptr
<cl_int
> ptr
=
119 compute::svm_alloc
<cl_int
>(context
, 8);
122 std::vector
<const void*> ptrs(1, ptr
.get());
123 std::vector
<size_t> sizes(1, 8 * sizeof(cl_int
));
124 queue
.enqueue_svm_migrate_memory(ptrs
, sizes
).wait();
127 const char source
[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
128 __kernel
void foo(__global
int *ptr
)
130 for(int i
= 0; i
< 8; i
++){
135 compute::program program
=
136 compute::program::build_with_source(source
, context
, "-cl-std=CL2.0");
137 compute::kernel foo_kernel
= program
.create_kernel("foo");
138 foo_kernel
.set_arg(0, ptr
);
139 queue
.enqueue_task(foo_kernel
).wait();
142 queue
.enqueue_svm_migrate_memory(
143 ptr
.get(), 0, boost::compute::command_queue::migrate_to_host
147 CHECK_HOST_RANGE_EQUAL(
149 static_cast<cl_int
*>(ptr
.get()),
150 (0, 1, 2, 3, 4, 5, 6, 7)
152 compute::svm_free(context
, ptr
);
154 #endif // BOOST_COMPUTE_CL_VERSION_2_1
156 BOOST_AUTO_TEST_SUITE_END()