Commit a9bc4489 authored by Joseph Noir's avatar Joseph Noir

added compiler options to program

+ unit test in test_opencl
parent 0b9cf32c
...@@ -57,7 +57,7 @@ class program { ...@@ -57,7 +57,7 @@ class program {
* from a given @p kernel_source. * from a given @p kernel_source.
* @returns A program object. * @returns A program object.
*/ */
static program create(const char* kernel_source, uint32_t device_id = 0); static program create(const char* kernel_source, const char* options = nullptr, uint32_t device_id = 0);
private: private:
......
...@@ -45,7 +45,7 @@ namespace cppa { namespace opencl { ...@@ -45,7 +45,7 @@ namespace cppa { namespace opencl {
program::program(context_ptr context, command_queue_ptr queue, program_ptr program) program::program(context_ptr context, command_queue_ptr queue, program_ptr program)
: m_context(move(context)), m_program(move(program)), m_queue(move(queue)) { } : m_context(move(context)), m_program(move(program)), m_queue(move(queue)) { }
program program::create(const char* kernel_source, uint32_t device_id) { program program::create(const char* kernel_source, const char* options, uint32_t device_id) {
auto metainfo = get_opencl_metainfo(); auto metainfo = get_opencl_metainfo();
auto devices = metainfo->get_devices(); auto devices = metainfo->get_devices();
auto context = metainfo->m_context; auto context = metainfo->m_context;
...@@ -77,7 +77,7 @@ program program::create(const char* kernel_source, uint32_t device_id) { ...@@ -77,7 +77,7 @@ program program::create(const char* kernel_source, uint32_t device_id) {
// build programm from program object // build programm from program object
auto dev_tmp = devices[device_id].m_device.get(); auto dev_tmp = devices[device_id].m_device.get();
err = clBuildProgram(pptr.get(), 1, &dev_tmp, nullptr, nullptr, nullptr); err = clBuildProgram(pptr.get(), 1, &dev_tmp, options, nullptr, nullptr);
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
ostringstream oss; ostringstream oss;
oss << "clBuildProgram: " << get_opencl_error(err); oss << "clBuildProgram: " << get_opencl_error(err);
......
...@@ -18,7 +18,12 @@ namespace { ...@@ -18,7 +18,12 @@ namespace {
using ivec = vector<int>; using ivec = vector<int>;
constexpr size_t matrix_size = 4; constexpr size_t matrix_size = 4;
constexpr size_t array_size = 32;
constexpr const char* kernel_name = "matrix_square"; constexpr const char* kernel_name = "matrix_square";
constexpr const char* kernel_name_compiler_flag = "compiler_flag";
constexpr const char* compiler_flag = "-D OPENCL_CPPA_TEST_FLAG";
constexpr const char* kernel_source = R"__( constexpr const char* kernel_source = R"__(
__kernel void matrix_square(__global int* matrix, __kernel void matrix_square(__global int* matrix,
...@@ -35,11 +40,23 @@ constexpr const char* kernel_source = R"__( ...@@ -35,11 +40,23 @@ constexpr const char* kernel_source = R"__(
)__"; )__";
constexpr const char* kernel_source_error = R"__( constexpr const char* kernel_source_error = R"__(
__kernel void matrix_square(__global int*) { __kernel void missing(__global int*) {
size_t semicolon size_t semicolon
} }
)__"; )__";
constexpr const char* kernel_source_compiler_flag = R"__(
__kernel void compiler_flag(__global int* input,
__global int* output) {
size_t x = get_global_id(0);
#ifdef OPENCL_CPPA_TEST_FLAG
output[x] = input[x];
#else
output[x] = 0;
#endif
}
)__";
} }
template<size_t Size> template<size_t Size>
...@@ -108,7 +125,7 @@ int main() { ...@@ -108,7 +125,7 @@ int main() {
announce<ivec>(); announce<ivec>();
announce<matrix_type>(); announce<matrix_type>();
const ivec expected1{ 56, 62, 68, 74 const ivec expected1{ 56, 62, 68, 74
, 152, 174, 196, 218 , 152, 174, 196, 218
, 248, 286, 324, 362 , 248, 286, 324, 362
, 344, 398, 452, 506}; , 344, 398, 452, 506};
...@@ -117,7 +134,7 @@ int main() { ...@@ -117,7 +134,7 @@ int main() {
kernel_name, kernel_name,
{matrix_size, matrix_size}); {matrix_size, matrix_size});
ivec m1(matrix_size * matrix_size); ivec m1(matrix_size * matrix_size);
iota(m1.begin(), m1.end(), 0); iota(begin(m1), end(m1), 0);
send(worker1, move(m1)); send(worker1, move(m1));
receive ( receive (
on_arg_match >> [&] (const ivec& result) { on_arg_match >> [&] (const ivec& result) {
...@@ -129,7 +146,7 @@ int main() { ...@@ -129,7 +146,7 @@ int main() {
kernel_name, kernel_name,
{matrix_size, matrix_size}); {matrix_size, matrix_size});
ivec m2(matrix_size * matrix_size); ivec m2(matrix_size * matrix_size);
iota(m2.begin(), m2.end(), 0); iota(begin(m2), end(m2), 0);
send(worker2, move(m2)); send(worker2, move(m2));
receive ( receive (
on_arg_match >> [&] (const ivec& result) { on_arg_match >> [&] (const ivec& result) {
...@@ -185,6 +202,20 @@ int main() { ...@@ -185,6 +202,20 @@ int main() {
CPPA_CHECK_EQUAL("clBuildProgram: CL_BUILD_PROGRAM_FAILURE", exc.what()); CPPA_CHECK_EQUAL("clBuildProgram: CL_BUILD_PROGRAM_FAILURE", exc.what());
} }
ivec arr5(array_size);
iota(begin(arr5), end(arr5), 0);
auto prog5 = program::create(kernel_source_compiler_flag, compiler_flag);
auto worker5 = spawn_cl<ivec(ivec&)>(prog5, kernel_name_compiler_flag, {array_size});
send(worker5, move(arr5));
ivec expected3(array_size);
iota(begin(expected3), end(expected3), 0);
receive (
on_arg_match >> [&] (const ivec& result) {
CPPA_CHECK(equal(begin(expected3), end(expected3), begin(result)));
}
);
cppa::await_all_others_done(); cppa::await_all_others_done();
cppa::shutdown(); cppa::shutdown();
......
Markdown is supported
0%
or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment