forked from boostorg/compute
-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathtest_program.cpp
217 lines (170 loc) · 6.65 KB
/
test_program.cpp
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
//---------------------------------------------------------------------------//
// Copyright (c) 2013 Kyle Lutz <[email protected]>
//
// Distributed under the Boost Software License, Version 1.0
// See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt
//
// See http://boostorg.github.com/compute for more information.
//---------------------------------------------------------------------------//
#define BOOST_TEST_MODULE TestProgram
#include <boost/test/unit_test.hpp>
// disable the automatic kernel compilation debug messages. this allows the
// test for program to check that compilation error exceptions are properly
// thrown when invalid kernel code is passed to program::build().
#undef BOOST_COMPUTE_DEBUG_KERNEL_COMPILATION
#include <boost/compute/kernel.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/program.hpp>
#include <boost/compute/utility/source.hpp>
#include "quirks.hpp"
#include "context_setup.hpp"
namespace compute = boost::compute;
const char source[] =
"__kernel void foo(__global float *x, const uint n) { }\n"
"__kernel void bar(__global int *x, __global int *y) { }\n";
BOOST_AUTO_TEST_CASE(get_program_info)
{
// create program
boost::compute::program program =
boost::compute::program::create_with_source(source, context);
// build program
program.build();
// check program info
#ifndef BOOST_COMPUTE_USE_OFFLINE_CACHE
BOOST_CHECK(program.source().empty() == false);
#endif
BOOST_CHECK(program.get_context() == context);
}
BOOST_AUTO_TEST_CASE(program_source)
{
// create program from source
boost::compute::program program =
boost::compute::program::create_with_source(source, context);
BOOST_CHECK_EQUAL(std::string(source), program.source());
}
BOOST_AUTO_TEST_CASE(program_multiple_sources)
{
std::vector<std::string> sources;
sources.push_back("__kernel void foo(__global int* x) { }\n");
sources.push_back("__kernel void bar(__global float* y) { }\n");
// create program from sources
boost::compute::program program =
boost::compute::program::create_with_source(sources, context);
program.build();
boost::compute::kernel foo = program.create_kernel("foo");
boost::compute::kernel bar = program.create_kernel("bar");
}
BOOST_AUTO_TEST_CASE(program_source_no_file)
{
// create program from a non-existant source file
// and verifies it throws.
BOOST_CHECK_THROW(boost::compute::program program =
boost::compute::program::create_with_source_file
(std::string(), context),
std::ios_base::failure);
}
BOOST_AUTO_TEST_CASE(create_kernel)
{
boost::compute::program program =
boost::compute::program::create_with_source(source, context);
program.build();
boost::compute::kernel foo = program.create_kernel("foo");
boost::compute::kernel bar = program.create_kernel("bar");
// try to create a kernel that doesn't exist
BOOST_CHECK_THROW(program.create_kernel("baz"), boost::compute::opencl_error);
}
BOOST_AUTO_TEST_CASE(create_with_binary)
{
// create program from source
boost::compute::program source_program =
boost::compute::program::create_with_source(source, context);
source_program.build();
// create kernels in source program
boost::compute::kernel source_foo_kernel = source_program.create_kernel("foo");
boost::compute::kernel source_bar_kernel = source_program.create_kernel("bar");
// check source kernels
BOOST_CHECK_EQUAL(source_foo_kernel.name(), std::string("foo"));
BOOST_CHECK_EQUAL(source_bar_kernel.name(), std::string("bar"));
// get binary
std::vector<unsigned char> binary = source_program.binary();
// create program from binary
boost::compute::program binary_program =
boost::compute::program::create_with_binary(binary, context);
binary_program.build();
// create kernels in binary program
boost::compute::kernel binary_foo_kernel = binary_program.create_kernel("foo");
boost::compute::kernel binary_bar_kernel = binary_program.create_kernel("bar");
// check binary kernels
BOOST_CHECK_EQUAL(binary_foo_kernel.name(), std::string("foo"));
BOOST_CHECK_EQUAL(binary_bar_kernel.name(), std::string("bar"));
}
BOOST_AUTO_TEST_CASE(create_with_source_doctest)
{
//! [create_with_source]
std::string source = "__kernel void foo(__global int *data) { }";
boost::compute::program foo_program =
boost::compute::program::create_with_source(source, context);
//! [create_with_source]
foo_program.build();
}
#ifdef CL_VERSION_1_2
BOOST_AUTO_TEST_CASE(compile_and_link)
{
REQUIRES_OPENCL_VERSION(1,2);
if(!supports_compile_program(device) || !supports_link_program(device)) {
return;
}
// create the library program
const char library_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
// for some reason the apple opencl compilers complains if a prototype
// for the square() function is not available, so we add it here
T square(T);
// generic square function definition
T square(T x) { return x * x; }
);
compute::program library_program =
compute::program::create_with_source(library_source, context);
library_program.compile("-DT=int");
// create the kernel program
const char kernel_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
// forward declare square function
extern int square(int);
// square kernel definition
__kernel void square_kernel(__global int *x)
{
x[0] = square(x[0]);
}
);
compute::program square_program =
compute::program::create_with_source(kernel_source, context);
square_program.compile();
// link the programs
std::vector<compute::program> programs;
programs.push_back(library_program);
programs.push_back(square_program);
compute::program linked_program =
compute::program::link(programs, context);
// create the square kernel
compute::kernel square_kernel =
linked_program.create_kernel("square_kernel");
BOOST_CHECK_EQUAL(square_kernel.name(), "square_kernel");
}
#endif // CL_VERSION_1_2
BOOST_AUTO_TEST_CASE(build_log)
{
const char invalid_source[] =
"__kernel void foo(__global int *input) { !@#$%^&*() }";
compute::program invalid_program =
compute::program::create_with_source(invalid_source, context);
try {
invalid_program.build();
// should not get here
BOOST_CHECK(false);
}
catch(compute::opencl_error&){
std::string log = invalid_program.build_log();
BOOST_CHECK(!log.empty());
}
}
BOOST_AUTO_TEST_SUITE_END()