Commit f17f3884 authored by Joseph Noir's avatar Joseph Noir

Added test for variable result size

... and fixed bug in variable result size.
parent a9bc4489
...@@ -168,7 +168,8 @@ class actor_facade<Ret(Args...)> : public actor { ...@@ -168,7 +168,8 @@ class actor_facade<Ret(Args...)> : public actor {
get_ref<Is>(*opt)...); get_ref<Is>(*opt)...);
auto cmd = make_counted<command<actor_facade, Ret>>(handle, auto cmd = make_counted<command<actor_facade, Ret>>(handle,
this, this,
std::move(arguments)); std::move(arguments),
m_result_size);
cmd->enqueue(); cmd->enqueue();
} }
else { CPPA_LOGMF(CPPA_ERROR, this, "actor_facade::enqueue() tuple_cast failed."); } else { CPPA_LOGMF(CPPA_ERROR, this, "actor_facade::enqueue() tuple_cast failed."); }
......
...@@ -53,15 +53,14 @@ class command : public ref_counted { ...@@ -53,15 +53,14 @@ class command : public ref_counted {
command(response_handle handle, command(response_handle handle,
intrusive_ptr<T> actor_facade, intrusive_ptr<T> actor_facade,
std::vector<mem_ptr> arguments) std::vector<mem_ptr> arguments,
: m_number_of_values(std::accumulate(actor_facade->m_global_dimensions.begin(), size_t result_size)
actor_facade->m_global_dimensions.end(), : m_result_size(result_size)
1, std::multiplies<size_t>{}))
, m_handle(handle) , m_handle(handle)
, m_actor_facade(actor_facade) , m_actor_facade(actor_facade)
, m_queue(actor_facade->m_queue) , m_queue(actor_facade->m_queue)
, m_arguments(move(arguments)) , m_arguments(move(arguments))
, m_result(m_number_of_values) { } , m_result(m_result_size) { }
void enqueue () { void enqueue () {
CPPA_LOG_TRACE("command::enqueue()"); CPPA_LOG_TRACE("command::enqueue()");
...@@ -92,7 +91,7 @@ class command : public ref_counted { ...@@ -92,7 +91,7 @@ class command : public ref_counted {
m_arguments[0].get(), m_arguments[0].get(),
CL_FALSE, CL_FALSE,
0, 0,
sizeof(typename R::value_type) * m_number_of_values, sizeof(typename R::value_type) * m_result_size,
m_result.data(), m_result.data(),
1, 1,
&event_k, &event_k,
...@@ -125,7 +124,7 @@ class command : public ref_counted { ...@@ -125,7 +124,7 @@ class command : public ref_counted {
private: private:
int m_number_of_values; int m_result_size;
response_handle m_handle; response_handle m_handle;
intrusive_ptr<T> m_actor_facade; intrusive_ptr<T> m_actor_facade;
event_ptr m_kernel_event; event_ptr m_kernel_event;
......
...@@ -16,12 +16,23 @@ using namespace cppa::opencl; ...@@ -16,12 +16,23 @@ using namespace cppa::opencl;
namespace { namespace {
using ivec = vector<int>; using ivec = vector<int>;
using fvec = vector<float>;
constexpr size_t matrix_size = 4; constexpr size_t matrix_size = 4;
constexpr size_t array_size = 32; constexpr size_t array_size = 32;
// since we do currently not support local memory arguments
// this size is fixed in the reduce kernel code
constexpr size_t reduce_buffer_size = 512 * 8;
constexpr size_t reduce_local_size = 512;
constexpr size_t reduce_work_groups = reduce_buffer_size / reduce_local_size;
constexpr size_t reduce_global_size = reduce_buffer_size;
constexpr size_t reduce_result_size = reduce_work_groups;
constexpr const char* kernel_name = "matrix_square"; constexpr const char* kernel_name = "matrix_square";
constexpr const char* kernel_name_result_size = "result_size";
constexpr const char* kernel_name_compiler_flag = "compiler_flag"; constexpr const char* kernel_name_compiler_flag = "compiler_flag";
constexpr const char* kernel_name_reduce = "reduce";
constexpr const char* compiler_flag = "-D OPENCL_CPPA_TEST_FLAG"; constexpr const char* compiler_flag = "-D OPENCL_CPPA_TEST_FLAG";
...@@ -57,6 +68,39 @@ constexpr const char* kernel_source_compiler_flag = R"__( ...@@ -57,6 +68,39 @@ constexpr const char* kernel_source_compiler_flag = R"__(
} }
)__"; )__";
// http://developer.amd.com/resources/documentation-articles/articles-whitepapers/opencl-optimization-case-study-simple-reductions/
constexpr const char* kernel_source_reduce = R"__(
__kernel void reduce(__global int* buffer,
__global int* result) {
__local int scratch[512];
size_t length = get_global_size(0);
size_t global_index = get_global_id(0);
int accumulator = INFINITY;
// Loop sequentially over chunks of input vector
while (global_index < length) {
int element = buffer[global_index];
accumulator = (accumulator < element) ? accumulator : element;
global_index += length;
}
// Perform parallel reduction
int local_index = get_local_id(0);
scratch[local_index] = accumulator;
barrier(CLK_LOCAL_MEM_FENCE);
for(int offset = get_local_size(0) / 2; offset > 0; offset = offset / 2) {
if (local_index < offset) {
int other = scratch[local_index + offset];
int mine = scratch[local_index];
scratch[local_index] = (mine < other) ? mine : other;
}
barrier(CLK_LOCAL_MEM_FENCE);
}
if (local_index == 0) {
result[get_group_id(0)] = scratch[0];
}
}
)__";
} }
template<size_t Size> template<size_t Size>
...@@ -202,6 +246,7 @@ int main() { ...@@ -202,6 +246,7 @@ int main() {
CPPA_CHECK_EQUAL("clBuildProgram: CL_BUILD_PROGRAM_FAILURE", exc.what()); CPPA_CHECK_EQUAL("clBuildProgram: CL_BUILD_PROGRAM_FAILURE", exc.what());
} }
// test for opencl compiler flags
ivec arr5(array_size); ivec arr5(array_size);
iota(begin(arr5), end(arr5), 0); iota(begin(arr5), end(arr5), 0);
auto prog5 = program::create(kernel_source_compiler_flag, compiler_flag); auto prog5 = program::create(kernel_source_compiler_flag, compiler_flag);
...@@ -216,6 +261,24 @@ int main() { ...@@ -216,6 +261,24 @@ int main() {
} }
); );
// test for manuel return size selection
ivec arr6(reduce_buffer_size);
int n{static_cast<int>(arr6.capacity())};
generate(begin(arr6), end(arr6), [&]{ return --n; });
auto worker6 = spawn_cl<ivec(ivec&)>(kernel_source_reduce,
kernel_name_reduce,
{reduce_global_size},
{},
{reduce_local_size},
reduce_result_size);
send(worker6, move(arr6));
fvec expected4{3584, 3072, 2560, 2048, 1536, 1024, 512, 0};
receive (
on_arg_match >> [&] (const ivec& result) {
CPPA_CHECK(equal(begin(expected4), end(expected4), 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