]>
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 TestProgram | |
12 | #include <boost/test/unit_test.hpp> | |
13 | ||
14 | // disable the automatic kernel compilation debug messages. this allows the | |
15 | // test for program to check that compilation error exceptions are properly | |
16 | // thrown when invalid kernel code is passed to program::build(). | |
17 | #undef BOOST_COMPUTE_DEBUG_KERNEL_COMPILATION | |
18 | ||
92f5a8d4 | 19 | #include <boost/compute/exception/program_build_failure.hpp> |
7c673cae FG |
20 | #include <boost/compute/kernel.hpp> |
21 | #include <boost/compute/system.hpp> | |
22 | #include <boost/compute/program.hpp> | |
23 | #include <boost/compute/utility/source.hpp> | |
24 | ||
25 | #include "quirks.hpp" | |
26 | #include "context_setup.hpp" | |
27 | ||
28 | namespace compute = boost::compute; | |
29 | ||
30 | const char source[] = | |
31 | "__kernel void foo(__global float *x, const uint n) { }\n" | |
32 | "__kernel void bar(__global int *x, __global int *y) { }\n"; | |
33 | ||
34 | ||
35 | BOOST_AUTO_TEST_CASE(get_program_info) | |
36 | { | |
37 | // create program | |
38 | boost::compute::program program = | |
39 | boost::compute::program::create_with_source(source, context); | |
40 | ||
41 | // build program | |
42 | program.build(); | |
43 | ||
44 | // check program info | |
45 | #ifndef BOOST_COMPUTE_USE_OFFLINE_CACHE | |
46 | BOOST_CHECK(program.source().empty() == false); | |
47 | #endif | |
48 | BOOST_CHECK(program.get_context() == context); | |
49 | } | |
50 | ||
51 | BOOST_AUTO_TEST_CASE(program_source) | |
52 | { | |
53 | // create program from source | |
54 | boost::compute::program program = | |
55 | boost::compute::program::create_with_source(source, context); | |
56 | ||
57 | BOOST_CHECK_EQUAL(std::string(source), program.source()); | |
58 | } | |
59 | ||
60 | BOOST_AUTO_TEST_CASE(program_multiple_sources) | |
61 | { | |
62 | std::vector<std::string> sources; | |
63 | sources.push_back("__kernel void foo(__global int* x) { }\n"); | |
64 | sources.push_back("__kernel void bar(__global float* y) { }\n"); | |
65 | ||
66 | // create program from sources | |
67 | boost::compute::program program = | |
68 | boost::compute::program::create_with_source(sources, context); | |
69 | program.build(); | |
70 | ||
71 | boost::compute::kernel foo = program.create_kernel("foo"); | |
72 | boost::compute::kernel bar = program.create_kernel("bar"); | |
73 | } | |
74 | ||
75 | BOOST_AUTO_TEST_CASE(program_source_no_file) | |
76 | { | |
77 | // create program from a non-existant source file | |
78 | // and verifies it throws. | |
79 | BOOST_CHECK_THROW(boost::compute::program program = | |
80 | boost::compute::program::create_with_source_file | |
81 | (std::string(), context), | |
82 | std::ios_base::failure); | |
83 | } | |
84 | ||
85 | BOOST_AUTO_TEST_CASE(create_kernel) | |
86 | { | |
87 | boost::compute::program program = | |
88 | boost::compute::program::create_with_source(source, context); | |
89 | program.build(); | |
90 | ||
91 | boost::compute::kernel foo = program.create_kernel("foo"); | |
92 | boost::compute::kernel bar = program.create_kernel("bar"); | |
93 | ||
94 | // try to create a kernel that doesn't exist | |
95 | BOOST_CHECK_THROW(program.create_kernel("baz"), boost::compute::opencl_error); | |
96 | } | |
97 | ||
98 | BOOST_AUTO_TEST_CASE(create_with_binary) | |
99 | { | |
100 | // create program from source | |
101 | boost::compute::program source_program = | |
102 | boost::compute::program::create_with_source(source, context); | |
103 | source_program.build(); | |
104 | ||
105 | // create kernels in source program | |
106 | boost::compute::kernel source_foo_kernel = source_program.create_kernel("foo"); | |
107 | boost::compute::kernel source_bar_kernel = source_program.create_kernel("bar"); | |
108 | ||
109 | // check source kernels | |
110 | BOOST_CHECK_EQUAL(source_foo_kernel.name(), std::string("foo")); | |
111 | BOOST_CHECK_EQUAL(source_bar_kernel.name(), std::string("bar")); | |
112 | ||
113 | // get binary | |
114 | std::vector<unsigned char> binary = source_program.binary(); | |
115 | ||
116 | // create program from binary | |
117 | boost::compute::program binary_program = | |
118 | boost::compute::program::create_with_binary(binary, context); | |
119 | binary_program.build(); | |
120 | ||
121 | // create kernels in binary program | |
122 | boost::compute::kernel binary_foo_kernel = binary_program.create_kernel("foo"); | |
123 | boost::compute::kernel binary_bar_kernel = binary_program.create_kernel("bar"); | |
124 | ||
125 | // check binary kernels | |
126 | BOOST_CHECK_EQUAL(binary_foo_kernel.name(), std::string("foo")); | |
127 | BOOST_CHECK_EQUAL(binary_bar_kernel.name(), std::string("bar")); | |
128 | } | |
129 | ||
b32b8144 FG |
130 | #ifdef BOOST_COMPUTE_CL_VERSION_2_1 |
131 | BOOST_AUTO_TEST_CASE(create_with_il) | |
132 | { | |
92f5a8d4 TL |
133 | REQUIRES_OPENCL_VERSION(2, 1); |
134 | ||
b32b8144 FG |
135 | size_t device_address_space_size = device.address_bits(); |
136 | std::string file_path(BOOST_COMPUTE_TEST_DATA_PATH); | |
137 | if(device_address_space_size == 64) | |
138 | { | |
139 | file_path += "/program.spirv64"; | |
140 | } | |
141 | else | |
142 | { | |
143 | file_path += "/program.spirv32"; | |
144 | } | |
145 | ||
146 | // create program from il | |
147 | boost::compute::program il_program; | |
148 | BOOST_CHECK_NO_THROW( | |
149 | il_program = boost::compute::program::create_with_il_file( | |
150 | file_path, context | |
151 | ) | |
152 | ); | |
153 | BOOST_CHECK_NO_THROW(il_program.build()); | |
154 | ||
155 | // create kernel (to check if program was loaded correctly) | |
156 | BOOST_CHECK_NO_THROW(il_program.create_kernel("foobar")); | |
157 | } | |
158 | ||
159 | BOOST_AUTO_TEST_CASE(get_program_il_binary) | |
160 | { | |
92f5a8d4 TL |
161 | REQUIRES_OPENCL_VERSION(2, 1); |
162 | ||
b32b8144 FG |
163 | size_t device_address_space_size = device.address_bits(); |
164 | std::string file_path(BOOST_COMPUTE_TEST_DATA_PATH); | |
165 | if(device_address_space_size == 64) | |
166 | { | |
167 | file_path += "/program.spirv64"; | |
168 | } | |
169 | else | |
170 | { | |
171 | file_path += "/program.spirv32"; | |
172 | } | |
173 | ||
174 | // create program from il | |
175 | boost::compute::program il_program; | |
176 | BOOST_CHECK_NO_THROW( | |
177 | il_program = boost::compute::program::create_with_il_file( | |
178 | file_path, context | |
179 | ) | |
180 | ); | |
181 | BOOST_CHECK_NO_THROW(il_program.build()); | |
182 | ||
183 | std::vector<unsigned char> il_binary; | |
184 | BOOST_CHECK_NO_THROW(il_binary = il_program.il_binary()); | |
185 | ||
186 | // create program from loaded il binary | |
187 | BOOST_CHECK_NO_THROW( | |
188 | il_program = boost::compute::program::create_with_il(il_binary, context) | |
189 | ); | |
190 | BOOST_CHECK_NO_THROW(il_program.build()); | |
191 | ||
192 | // create kernel (to check if program was loaded correctly) | |
193 | BOOST_CHECK_NO_THROW(il_program.create_kernel("foobar")); | |
194 | } | |
195 | ||
196 | BOOST_AUTO_TEST_CASE(get_program_il_binary_empty) | |
197 | { | |
92f5a8d4 TL |
198 | REQUIRES_OPENCL_VERSION(2, 1); |
199 | ||
b32b8144 FG |
200 | boost::compute::program program; |
201 | BOOST_CHECK_NO_THROW( | |
202 | program = boost::compute::program::create_with_source(source, context) | |
203 | ); | |
204 | BOOST_CHECK_NO_THROW(program.build()); | |
205 | ||
206 | std::vector<unsigned char> il_binary; | |
207 | il_binary = program.il_binary(); | |
208 | BOOST_CHECK(il_binary.empty()); | |
209 | } | |
210 | #endif // BOOST_COMPUTE_CL_VERSION_2_1 | |
211 | ||
7c673cae FG |
212 | BOOST_AUTO_TEST_CASE(create_with_source_doctest) |
213 | { | |
214 | //! [create_with_source] | |
215 | std::string source = "__kernel void foo(__global int *data) { }"; | |
216 | ||
217 | boost::compute::program foo_program = | |
218 | boost::compute::program::create_with_source(source, context); | |
219 | //! [create_with_source] | |
220 | ||
221 | foo_program.build(); | |
222 | } | |
223 | ||
b32b8144 | 224 | #ifdef BOOST_COMPUTE_CL_VERSION_1_2 |
7c673cae FG |
225 | BOOST_AUTO_TEST_CASE(compile_and_link) |
226 | { | |
227 | REQUIRES_OPENCL_VERSION(1,2); | |
228 | ||
229 | if(!supports_compile_program(device) || !supports_link_program(device)) { | |
230 | return; | |
231 | } | |
232 | ||
233 | // create the library program | |
234 | const char library_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE( | |
235 | // for some reason the apple opencl compilers complains if a prototype | |
236 | // for the square() function is not available, so we add it here | |
237 | T square(T); | |
238 | ||
239 | // generic square function definition | |
240 | T square(T x) { return x * x; } | |
241 | ); | |
242 | ||
243 | compute::program library_program = | |
244 | compute::program::create_with_source(library_source, context); | |
245 | ||
246 | library_program.compile("-DT=int"); | |
247 | ||
248 | // create the kernel program | |
249 | const char kernel_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE( | |
250 | // forward declare square function | |
251 | extern int square(int); | |
252 | ||
253 | // square kernel definition | |
254 | __kernel void square_kernel(__global int *x) | |
255 | { | |
256 | x[0] = square(x[0]); | |
257 | } | |
258 | ); | |
259 | ||
260 | compute::program square_program = | |
261 | compute::program::create_with_source(kernel_source, context); | |
262 | ||
263 | square_program.compile(); | |
264 | ||
265 | // link the programs | |
266 | std::vector<compute::program> programs; | |
267 | programs.push_back(library_program); | |
268 | programs.push_back(square_program); | |
269 | ||
270 | compute::program linked_program = | |
271 | compute::program::link(programs, context); | |
272 | ||
273 | // create the square kernel | |
274 | compute::kernel square_kernel = | |
275 | linked_program.create_kernel("square_kernel"); | |
276 | BOOST_CHECK_EQUAL(square_kernel.name(), "square_kernel"); | |
277 | } | |
b32b8144 FG |
278 | |
279 | BOOST_AUTO_TEST_CASE(compile_and_link_with_headers) | |
280 | { | |
281 | REQUIRES_OPENCL_VERSION(1,2); | |
282 | ||
283 | if(!supports_compile_program(device) || !supports_link_program(device)) { | |
284 | return; | |
285 | } | |
286 | ||
287 | // create the header programs | |
288 | const char square_header_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE( | |
289 | T square(T x) { return x * x; } | |
290 | ); | |
291 | const char div2_header_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE( | |
292 | T div2(T x) { return x / 2; } | |
293 | ); | |
294 | ||
295 | compute::program square_header_program = | |
296 | compute::program::create_with_source(square_header_source, context); | |
297 | compute::program div2_header_program = | |
298 | compute::program::create_with_source(div2_header_source, context); | |
299 | ||
300 | // create the kernel program | |
301 | const char kernel_source[] = | |
302 | "#include \"square.h\"\n" | |
303 | "#include \"div2.h\"\n" | |
304 | "__kernel void squareby2_kernel(__global int *x)" | |
305 | "{" | |
306 | " x[0] = div2(square(x[0]));" | |
307 | "}"; | |
308 | ||
309 | compute::program square_program = | |
310 | compute::program::create_with_source(kernel_source, context); | |
311 | ||
312 | std::vector<std::pair<std::string, compute::program> > header_programs; | |
313 | header_programs.push_back(std::make_pair("square.h", square_header_program)); | |
314 | header_programs.push_back(std::make_pair("div2.h", div2_header_program)); | |
315 | ||
316 | square_program.compile("-DT=int", header_programs); | |
317 | ||
318 | // link program | |
319 | std::vector<compute::program> programs; | |
320 | programs.push_back(square_program); | |
321 | ||
322 | compute::program linked_program = | |
323 | compute::program::link(programs, context); | |
324 | ||
325 | // create the square kernel | |
326 | compute::kernel square_kernel = | |
327 | linked_program.create_kernel("squareby2_kernel"); | |
328 | BOOST_CHECK_EQUAL(square_kernel.name(), "squareby2_kernel"); | |
329 | } | |
330 | #endif // BOOST_COMPUTE_CL_VERSION_1_2 | |
7c673cae FG |
331 | |
332 | BOOST_AUTO_TEST_CASE(build_log) | |
333 | { | |
334 | const char invalid_source[] = | |
335 | "__kernel void foo(__global int *input) { !@#$%^&*() }"; | |
336 | ||
337 | compute::program invalid_program = | |
338 | compute::program::create_with_source(invalid_source, context); | |
339 | ||
340 | try { | |
341 | invalid_program.build(); | |
342 | ||
343 | // should not get here | |
344 | BOOST_CHECK(false); | |
345 | } | |
346 | catch(compute::opencl_error&){ | |
347 | std::string log = invalid_program.build_log(); | |
348 | BOOST_CHECK(!log.empty()); | |
349 | } | |
350 | } | |
351 | ||
92f5a8d4 TL |
352 | BOOST_AUTO_TEST_CASE(program_build_exception) |
353 | { | |
354 | const char invalid_source[] = | |
355 | "__kernel void foo(__global int *input) { !@#$%^&*() }"; | |
356 | ||
357 | compute::program invalid_program = | |
358 | compute::program::create_with_source(invalid_source, context); | |
359 | ||
360 | BOOST_CHECK_THROW(invalid_program.build(), | |
361 | compute::program_build_failure); | |
362 | ||
363 | try { | |
364 | // POCL bug: https://github.com/pocl/pocl/issues/577 | |
365 | if(pocl_bug_issue_577(device)) | |
366 | { | |
367 | invalid_program = | |
368 | compute::program::create_with_source(invalid_source, context); | |
369 | } | |
370 | invalid_program.build(); | |
371 | ||
372 | // should not get here | |
373 | BOOST_CHECK(false); | |
374 | } | |
375 | catch(compute::program_build_failure& e){ | |
376 | BOOST_CHECK(e.build_log() == invalid_program.build_log()); | |
377 | } | |
378 | catch(...) | |
379 | { | |
380 | // should not get here | |
381 | BOOST_CHECK(false); | |
382 | } | |
383 | } | |
384 | ||
385 | BOOST_AUTO_TEST_CASE(build_with_source_exception) | |
386 | { | |
387 | const char invalid_source[] = | |
388 | "__kernel void foo(__global int *input) { !@#$%^&*() }"; | |
389 | ||
390 | BOOST_CHECK_THROW(compute::program::build_with_source(invalid_source, context), | |
391 | compute::program_build_failure); | |
392 | } | |
393 | ||
394 | BOOST_AUTO_TEST_CASE(build_with_source_file_exception) | |
395 | { | |
396 | std::string file_path(BOOST_COMPUTE_TEST_DATA_PATH "/invalid_program.cl"); | |
397 | BOOST_CHECK_THROW(compute::program::build_with_source_file(file_path, context), | |
398 | compute::program_build_failure); | |
399 | } | |
400 | ||
7c673cae | 401 | BOOST_AUTO_TEST_SUITE_END() |