]>
Commit | Line | Data |
---|---|---|
7c673cae FG |
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 | ||
27 | BOOST_AUTO_TEST_CASE(empty) | |
28 | { | |
29 | } | |
30 | ||
b32b8144 | 31 | #ifdef BOOST_COMPUTE_CL_VERSION_2_0 |
7c673cae FG |
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 | ||
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 | ||
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 | } | |
b32b8144 FG |
111 | #endif // BOOST_COMPUTE_CL_VERSION_2_0 |
112 | ||
113 | #ifdef BOOST_COMPUTE_CL_VERSION_2_1 | |
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 | |
7c673cae FG |
155 | |
156 | BOOST_AUTO_TEST_SUITE_END() |