]>
Commit | Line | Data |
---|---|---|
7c673cae FG |
1 | //---------------------------------------------------------------------------// |
2 | // Copyright (c) 2013 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 TestKernel | |
12 | #include <boost/test/unit_test.hpp> | |
13 | ||
14 | #include <boost/compute/buffer.hpp> | |
15 | #include <boost/compute/kernel.hpp> | |
16 | #include <boost/compute/system.hpp> | |
17 | #include <boost/compute/utility/source.hpp> | |
18 | ||
19 | #include "context_setup.hpp" | |
b32b8144 | 20 | #include "check_macros.hpp" |
7c673cae FG |
21 | |
22 | namespace compute = boost::compute; | |
23 | ||
24 | BOOST_AUTO_TEST_CASE(name) | |
25 | { | |
26 | compute::kernel foo = compute::kernel::create_with_source( | |
27 | "__kernel void foo(int x) { }", "foo", context | |
28 | ); | |
29 | BOOST_CHECK_EQUAL(foo.name(), "foo"); | |
30 | ||
31 | compute::kernel bar = compute::kernel::create_with_source( | |
32 | "__kernel void bar(float x) { }", "bar", context | |
33 | ); | |
34 | BOOST_CHECK_EQUAL(bar.name(), "bar"); | |
35 | } | |
36 | ||
37 | BOOST_AUTO_TEST_CASE(arity) | |
38 | { | |
39 | compute::kernel foo = compute::kernel::create_with_source( | |
40 | "__kernel void foo(int x) { }", "foo", context | |
41 | ); | |
42 | BOOST_CHECK_EQUAL(foo.arity(), size_t(1)); | |
43 | ||
44 | compute::kernel bar = compute::kernel::create_with_source( | |
45 | "__kernel void bar(float x, float y) { }", "bar", context | |
46 | ); | |
47 | BOOST_CHECK_EQUAL(bar.arity(), size_t(2)); | |
48 | ||
49 | compute::kernel baz = compute::kernel::create_with_source( | |
50 | "__kernel void baz(char x, char y, char z) { }", "baz", context | |
51 | ); | |
52 | BOOST_CHECK_EQUAL(baz.arity(), size_t(3)); | |
53 | } | |
54 | ||
55 | BOOST_AUTO_TEST_CASE(set_buffer_arg) | |
56 | { | |
57 | const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE( | |
58 | __kernel void foo(__global int *x, __global int *y) | |
59 | { | |
60 | x[get_global_id(0)] = -y[get_global_id(0)]; | |
61 | } | |
62 | ); | |
63 | ||
64 | compute::kernel foo = | |
65 | compute::kernel::create_with_source(source, "foo", context); | |
66 | ||
67 | compute::buffer x(context, 16); | |
68 | compute::buffer y(context, 16); | |
69 | ||
70 | foo.set_arg(0, x); | |
71 | foo.set_arg(1, y.get()); | |
72 | } | |
73 | ||
74 | BOOST_AUTO_TEST_CASE(get_work_group_info) | |
75 | { | |
76 | const char source[] = | |
77 | "__kernel void sum(__global const float *input,\n" | |
78 | " __global float *output)\n" | |
79 | "{\n" | |
80 | " __local float scratch[16];\n" | |
81 | " const uint gid = get_global_id(0);\n" | |
82 | " const uint lid = get_local_id(0);\n" | |
83 | " if(lid < 16)\n" | |
84 | " scratch[lid] = input[gid];\n" | |
85 | "}\n"; | |
86 | ||
87 | compute::program program = | |
88 | compute::program::create_with_source(source, context); | |
89 | ||
90 | program.build(); | |
91 | ||
92 | compute::kernel kernel = program.create_kernel("sum"); | |
93 | ||
94 | using compute::ulong_; | |
95 | ||
96 | // get local memory size | |
97 | kernel.get_work_group_info<ulong_>(device, CL_KERNEL_LOCAL_MEM_SIZE); | |
98 | ||
99 | // check work group size | |
100 | size_t work_group_size = | |
101 | kernel.get_work_group_info<size_t>(device, CL_KERNEL_WORK_GROUP_SIZE); | |
102 | BOOST_CHECK(work_group_size >= 1); | |
103 | } | |
104 | ||
105 | #ifndef BOOST_COMPUTE_NO_VARIADIC_TEMPLATES | |
106 | BOOST_AUTO_TEST_CASE(kernel_set_args) | |
107 | { | |
108 | compute::kernel k = compute::kernel::create_with_source( | |
109 | "__kernel void test(int x, float y, char z) { }", "test", context | |
110 | ); | |
111 | ||
112 | k.set_args(4, 2.4f, 'a'); | |
113 | } | |
114 | #endif // BOOST_COMPUTE_NO_VARIADIC_TEMPLATES | |
115 | ||
b32b8144 | 116 | #ifdef BOOST_COMPUTE_CL_VERSION_1_2 |
7c673cae FG |
117 | BOOST_AUTO_TEST_CASE(get_arg_info) |
118 | { | |
119 | REQUIRES_OPENCL_VERSION(1, 2); | |
120 | ||
121 | const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE( | |
122 | __kernel void sum_kernel(__global const int *input, | |
123 | const uint size, | |
124 | __global int *result) | |
125 | { | |
126 | int sum = 0; | |
127 | for(uint i = 0; i < size; i++){ | |
128 | sum += input[i]; | |
129 | } | |
130 | *result = sum; | |
131 | } | |
132 | ); | |
133 | ||
134 | compute::program program = | |
135 | compute::program::create_with_source(source, context); | |
136 | ||
137 | program.build("-cl-kernel-arg-info"); | |
138 | ||
139 | compute::kernel kernel = program.create_kernel("sum_kernel"); | |
140 | ||
141 | BOOST_CHECK_EQUAL(kernel.get_info<CL_KERNEL_NUM_ARGS>(), compute::uint_(3)); | |
142 | ||
143 | BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(0, CL_KERNEL_ARG_TYPE_NAME), "int*"); | |
144 | BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(0, CL_KERNEL_ARG_NAME), "input"); | |
145 | BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(1, CL_KERNEL_ARG_TYPE_NAME), "uint"); | |
146 | BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(1, CL_KERNEL_ARG_NAME), "size"); | |
147 | BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(2, CL_KERNEL_ARG_TYPE_NAME), "int*"); | |
148 | BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(2, CL_KERNEL_ARG_NAME), "result"); | |
149 | ||
150 | BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_TYPE_NAME>(0), "int*"); | |
151 | BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_NAME>(0), "input"); | |
152 | BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_TYPE_NAME>(1), "uint"); | |
153 | BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_NAME>(1), "size"); | |
154 | BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_TYPE_NAME>(2), "int*"); | |
155 | BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_NAME>(2), "result"); | |
156 | } | |
b32b8144 FG |
157 | #endif // BOOST_COMPUTE_CL_VERSION_1_2 |
158 | ||
159 | #ifdef BOOST_COMPUTE_CL_VERSION_2_0 | |
160 | #ifndef CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE_KHR | |
161 | #define CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE_KHR CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE | |
162 | #endif | |
163 | #ifndef CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR | |
164 | #define CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE | |
165 | #endif | |
166 | BOOST_AUTO_TEST_CASE(get_sub_group_info_ext) | |
167 | { | |
168 | compute::kernel k = compute::kernel::create_with_source( | |
169 | "__kernel void test(float x) { }", "test", context | |
170 | ); | |
171 | ||
172 | // get_sub_group_info(const device&, cl_kernel_sub_group_info, const std::vector<size_t>) | |
173 | std::vector<size_t> local_work_size(2, size_t(64)); | |
174 | boost::optional<size_t> count = k.get_sub_group_info<size_t>( | |
175 | device, | |
176 | CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR, | |
177 | local_work_size | |
178 | ); | |
179 | ||
180 | if(device.check_version(2, 1)) | |
181 | { | |
182 | BOOST_CHECK(count); | |
183 | } | |
184 | else if(device.check_version(2, 0) && device.supports_extension("cl_khr_subgroups")) | |
185 | { | |
186 | // for device with cl_khr_subgroups it should return some value | |
187 | BOOST_CHECK(count); | |
188 | } | |
189 | else | |
190 | { | |
191 | // for device without cl_khr_subgroups ext it should return null optional | |
192 | BOOST_CHECK(count == boost::none); | |
193 | } | |
194 | ||
195 | // get_sub_group_info(const device&, cl_kernel_sub_group_info, const size_t, const void *) | |
196 | count = k.get_sub_group_info<size_t>( | |
197 | device, | |
198 | CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR, | |
199 | 2 * sizeof(size_t), | |
200 | &local_work_size[0] | |
201 | ); | |
202 | ||
203 | if(device.check_version(2, 1)) | |
204 | { | |
205 | BOOST_CHECK(count); | |
206 | } | |
207 | else if(device.check_version(2, 0) && device.supports_extension("cl_khr_subgroups")) | |
208 | { | |
209 | // for device with cl_khr_subgroups it should return some value | |
210 | BOOST_CHECK(count); | |
211 | } | |
212 | else | |
213 | { | |
214 | // for device without cl_khr_subgroups ext it should return null optional | |
215 | BOOST_CHECK(count == boost::none); | |
216 | } | |
217 | } | |
218 | #endif // BOOST_COMPUTE_CL_VERSION_2_0 | |
219 | ||
220 | #ifdef BOOST_COMPUTE_CL_VERSION_2_1 | |
221 | BOOST_AUTO_TEST_CASE(get_sub_group_info_core) | |
222 | { | |
223 | compute::kernel k = compute::kernel::create_with_source( | |
224 | "__kernel void test(float x) { }", "test", context | |
225 | ); | |
226 | ||
227 | // get_sub_group_info(const device&, cl_kernel_sub_group_info, const size_t) | |
228 | boost::optional<std::vector<size_t>> local_size = | |
229 | k.get_sub_group_info<std::vector<size_t> >( | |
230 | device, | |
231 | CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT, | |
232 | size_t(1) | |
233 | ); | |
234 | ||
235 | if(device.check_version(2, 1)) | |
236 | { | |
237 | // for 2.1 devices it should return some value | |
238 | BOOST_CHECK(local_size); | |
239 | BOOST_CHECK(local_size.value().size() == 3); | |
240 | } | |
241 | else | |
242 | { | |
243 | // for 1.x and 2.0 devices it should return null optional, | |
244 | // because CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT is not | |
245 | // supported by cl_khr_subgroups (2.0 ext) | |
246 | BOOST_CHECK(local_size == boost::none); | |
247 | } | |
248 | ||
249 | // get_sub_group_info(const device&, cl_kernel_sub_group_info, const size_t) | |
250 | boost::optional<size_t> local_size_simple = | |
251 | k.get_sub_group_info<size_t>( | |
252 | device, | |
253 | CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT, | |
254 | size_t(1) | |
255 | ); | |
256 | ||
257 | if(device.check_version(2, 1)) | |
258 | { | |
259 | // for 2.1 devices it should return some value | |
260 | BOOST_CHECK(local_size_simple); | |
261 | } | |
262 | else | |
263 | { | |
264 | // for 1.x and 2.0 devices it should return null optional, | |
265 | // because CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT is not | |
266 | // supported by cl_khr_subgroups (2.0 ext) | |
267 | BOOST_CHECK(local_size_simple == boost::none); | |
268 | } | |
269 | ||
270 | // get_sub_group_info(const device&, cl_kernel_sub_group_info) | |
271 | boost::optional<size_t> max = | |
272 | k.get_sub_group_info<size_t>( | |
273 | device, | |
274 | CL_KERNEL_MAX_NUM_SUB_GROUPS | |
275 | ); | |
276 | ||
277 | if(device.check_version(2, 1)) | |
278 | { | |
279 | // for 2.1 devices it should return some value | |
280 | BOOST_CHECK(max); | |
281 | } | |
282 | else | |
283 | { | |
284 | // for 1.x and 2.0 devices it should return null optional, | |
285 | // because CL_KERNEL_MAX_NUM_SUB_GROUPS is not | |
286 | // supported by cl_khr_subgroups (2.0 ext) | |
287 | BOOST_CHECK(max == boost::none); | |
288 | } | |
289 | } | |
290 | #endif // BOOST_COMPUTE_CL_VERSION_2_1 | |
291 | ||
292 | #ifdef BOOST_COMPUTE_CL_VERSION_2_1 | |
293 | BOOST_AUTO_TEST_CASE(clone_kernel) | |
294 | { | |
295 | REQUIRES_OPENCL_PLATFORM_VERSION(2, 1); | |
296 | ||
297 | compute::kernel k1 = compute::kernel::create_with_source( | |
298 | "__kernel void test(__global int * x) { x[get_global_id(0)] = get_global_id(0); }", | |
299 | "test", context | |
300 | ); | |
301 | ||
302 | compute::buffer x(context, 5 * sizeof(compute::int_)); | |
303 | k1.set_arg(0, x); | |
304 | ||
305 | // Clone k1 kernel | |
306 | compute::kernel k2 = k1.clone(); | |
307 | // After clone k2 0th argument (__global float * x) should be set, | |
308 | // so we should be able to enqueue k2 kernel without problems | |
309 | queue.enqueue_1d_range_kernel(k2, 0, x.size() / sizeof(compute::int_), 0).wait(); | |
310 | } | |
311 | #endif // BOOST_COMPUTE_CL_VERSION_2_1 | |
7c673cae FG |
312 | |
313 | BOOST_AUTO_TEST_SUITE_END() |