1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013-2014 Kyle Lutz <kyle.r.lutz@gmail.com>
3 //
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
7 //
8 // See http://boostorg.github.com/compute for more information.
9 //---------------------------------------------------------------------------//
10
11 #define BOOST_TEST_MODULE TestSvmPtr
12 #include <boost/test/unit_test.hpp>
13
14 #include <iostream>
15
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>
20
21 #include "quirks.hpp"
22 #include "check_macros.hpp"
23 #include "context_setup.hpp"
24
25 namespace compute = boost::compute;
26
BOOST_AUTO_TEST_CASE(empty)27 BOOST_AUTO_TEST_CASE(empty)
28 {
29 }
30
31 #ifdef BOOST_COMPUTE_CL_VERSION_2_0
BOOST_AUTO_TEST_CASE(alloc)32 BOOST_AUTO_TEST_CASE(alloc)
33 {
34 REQUIRES_OPENCL_VERSION(2, 0);
35
36 compute::svm_ptr<cl_int> ptr = compute::svm_alloc<cl_int>(context, 8);
37 compute::svm_free(context, ptr);
38 }
39
BOOST_AUTO_TEST_CASE(svmmemcpy)40 BOOST_AUTO_TEST_CASE(svmmemcpy)
41 {
42 REQUIRES_OPENCL_VERSION(2, 0);
43
44 if(bug_in_svmmemcpy(device)){
45 std::cerr << "skipping svmmemcpy test case" << std::endl;
46 return;
47 }
48
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);
53
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));
57 queue.finish();
58
59 CHECK_HOST_RANGE_EQUAL(cl_int, 8, output, (1, 2, 3, 4, 5, 6, 7, 8));
60
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));
64 queue.finish();
65
66 CHECK_HOST_RANGE_EQUAL(cl_int, 8, output, (1, 2, 3, 4, 5, 6, 7, 8));
67
68 compute::svm_free(context, ptr);
69 compute::svm_free(context, ptr2);
70 }
71
BOOST_AUTO_TEST_CASE(sum_svm_kernel)72 BOOST_AUTO_TEST_CASE(sum_svm_kernel)
73 {
74 REQUIRES_OPENCL_VERSION(2, 0);
75
76 const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
77 __kernel void sum_svm_mem(__global const int *ptr, __global int *result)
78 {
79 int sum = 0;
80 for(uint i = 0; i < 8; i++){
81 sum += ptr[i];
82 }
83 *result = sum;
84 }
85 );
86
87 compute::program program =
88 compute::program::build_with_source(source, context, "-cl-std=CL2.0");
89
90 compute::kernel sum_svm_mem_kernel = program.create_kernel("sum_svm_mem");
91
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];
97 }
98 queue.enqueue_svm_unmap(ptr.get());
99
100 compute::vector<cl_int> result(1, context);
101
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);
105
106 queue.finish();
107 BOOST_CHECK_EQUAL(result[0], (36));
108
109 compute::svm_free(context, ptr);
110 }
111 #endif // BOOST_COMPUTE_CL_VERSION_2_0
112
113 #ifdef BOOST_COMPUTE_CL_VERSION_2_1
BOOST_AUTO_TEST_CASE(migrate)114 BOOST_AUTO_TEST_CASE(migrate)
115 {
116 REQUIRES_OPENCL_VERSION(2, 1);
117
118 compute::svm_ptr<cl_int> ptr =
119 compute::svm_alloc<cl_int>(context, 8);
120
121 // Migrate to device
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();
125
126 // Set on device
127 const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
128 __kernel void foo(__global int *ptr)
129 {
130 for(int i = 0; i < 8; i++){
131 ptr[i] = i;
132 }
133 }
134 );
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();
140
141 // Migrate to host
142 queue.enqueue_svm_migrate_memory(
143 ptr.get(), 0, boost::compute::command_queue::migrate_to_host
144 ).wait();
145
146 // Check
147 CHECK_HOST_RANGE_EQUAL(
148 cl_int, 8,
149 static_cast<cl_int*>(ptr.get()),
150 (0, 1, 2, 3, 4, 5, 6, 7)
151 );
152 compute::svm_free(context, ptr);
153 }
154 #endif // BOOST_COMPUTE_CL_VERSION_2_1
155
156 BOOST_AUTO_TEST_SUITE_END()
157