Initial attempt at exposing events to the user

This only works with vector expressions (no additive expressions, no
multiexpressions) for now.

refs #196

tests/CMakeLists.txt

@@ -64,6 +64,7 @@ add_vexcl_test(reduce_by_key            reduce_by_key.cpp)
 add_vexcl_test(logical                  logical.cpp)
 add_vexcl_test(threads                  threads.cpp)
 add_vexcl_test(svm                      svm.cpp)
+add_vexcl_test(events                   events.cpp)
 add_vexcl_test(multiple_objects         "dummy1.cpp;dummy2.cpp")
 
 if (NOT DEFINED ENV{APPVEYOR})

tests/events.cpp

@@ -0,0 +1,37 @@
+#define BOOST_TEST_MODULE Let
+#include <boost/test/unit_test.hpp>
+#include <vexcl/vector.hpp>
+#include <vexcl/let.hpp>
+#include <vexcl/reductor.hpp>
+#include "context_setup.hpp"
+
+BOOST_AUTO_TEST_CASE(let_vector_expr)
+{
+    const size_t n = 16 * 1024;
+
+    std::vector<vex::command_queue> q1(1, ctx.queue(0));
+    std::vector<vex::command_queue> q2(1, vex::backend::duplicate_queue(ctx.queue(0)));
+
+    vex::vector<int> x(q1, n);
+    vex::vector<int> y(q2, n);
+
+    vex::Reductor<size_t> count(q2);
+
+    x = 1;
+    q1[0].finish();
+
+    auto e = vex::let(x) = 2;
+    let(y, e) = x;
+
+    BOOST_CHECK_EQUAL(count(y != 2), 0);
+
+    q1[0].finish();
+
+    x = 3; e[0] = vex::backend::enqueue_marker(q1[0]);
+    let(y, e) = x;
+
+    BOOST_CHECK_EQUAL(count(y != 3), 0);
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+

vexcl/backend/compute/event.hpp

@@ -0,0 +1,61 @@
+#ifndef VEXCL_BACKEND_COMPUTE_EVENT_HPP
+#define VEXCL_BACKEND_COMPUTE_EVENT_HPP
+
+/*
+The MIT License
+
+Copyright (c) 2012-2016 Denis Demidov <[email protected]>
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/**
+ * \file   vexcl/backend/compute/event.hpp
+ * \author Denis Demidov <[email protected]>
+ * \brief  Bring Boost.Compute events into vex::backend::compute namespace.
+ */
+
+#include <boost/compute/event.hpp>
+#include <boost/compute/wait_list.hpp>
+
+namespace vex {
+namespace backend {
+namespace compute {
+
+using boost::compute::event;
+using boost::compute::wait_list;
+
+inline void wait_list_append(wait_list &dst, const event &e) {
+    dst.insert(e);
+}
+
+inline void wait_list_append(wait_list &dst, const wait_list &src) {
+    for(size_t i = 0; i < src.size(); ++i)
+        dst.insert(src[i]);
+}
+
+inline boost::compute::event enqueue_marker(const boost::compute::command_queue &q) {
+    return q.enqueue_marker();
+}
+
+} // namespace compute
+} // namespace backend
+} // namespace vex
+
+#endif

vexcl/backend/compute/kernel.hpp

@@ -36,6 +36,7 @@ THE SOFTWARE.
 #include <boost/compute/core.hpp>
 #include <boost/compute/memory/local_buffer.hpp>
 #include <vexcl/backend/compute/compiler.hpp>
+#include <vexcl/backend/compute/event.hpp>
 
 namespace vex {
 namespace backend {
@@ -100,17 +101,37 @@ class kernel {
         }
 
         /// Enqueue the kernel to the specified command queue.
-        void operator()(boost::compute::command_queue q) {
-            q.enqueue_nd_range_kernel(K, 3, NULL, g_size.dim, w_size.dim);
+        boost::compute::event operator()(boost::compute::command_queue q) {
             argpos = 0;
+            return q.enqueue_nd_range_kernel(K, 3, NULL, g_size.dim, w_size.dim);
+        }
+
+        /// Enqueue the kernel to the specified command queue, provide wait list.
+        boost::compute::event operator()(
+                boost::compute::command_queue q,
+                const boost::compute::wait_list &events)
+        {
+            argpos = 0;
+            return q.enqueue_nd_range_kernel(K, 3, NULL, g_size.dim, w_size.dim, events);
         }
 
 #ifndef BOOST_NO_VARIADIC_TEMPLATES
         /// Enqueue the kernel to the specified command queue with the given arguments
         template <class... Args>
-        void operator()(boost::compute::command_queue q, Args&&... args) {
-            K.set_args(std::forward<Args>(args)...);
-            (*this)(q);
+        boost::compute::event operator()(boost::compute::command_queue q, const Args&... args) {
+            K.set_args(args...);
+            return (*this)(q);
+        }
+
+        /// Enqueue the kernel to the specified command queue with the given arguments
+        template <class... Args>
+        boost::compute::event operator()(
+                boost::compute::command_queue q,
+                const boost::compute::wait_list &events,
+                const Args&... args)
+        {
+            K.set_args(args...);
+            return (*this)(q, events);
         }
 #endif
 

vexcl/backend/cuda/context.hpp

@@ -101,6 +101,9 @@ class device {
         /// Returns raw CUdevice handle.
         CUdevice raw() const { return d; }
 
+        /// Returns raw CUdevice handle.
+        operator CUdevice() const { return d; }
+
         /// Returns name of the device.
         std::string name() const {
             char name[256];
@@ -166,6 +169,11 @@ class context {
             return c.get();
         }
 
+        /// Returns raw CUcontext handle.
+        operator CUcontext() const {
+            return c.get();
+        }
+
         /// Binds the context to the calling CPU thread.
         void set_current() const {
             cuda_check( cuCtxSetCurrent( c.get() ) );
@@ -223,6 +231,12 @@ class command_queue {
         CUstream raw() const {
             return s.get();
         }
+
+        /// Returns raw CUstream handle for the command queue.
+        operator CUstream() const {
+            return s.get();
+        }
+
     private:
         vex::backend::context  ctx;
         vex::backend::device   dev;

vexcl/backend/cuda/event.hpp

@@ -0,0 +1,97 @@
+#ifndef VEXCL_BACKEND_CUDA_EVENT_HPP
+#define VEXCL_BACKEND_CUDA_EVENT_HPP
+
+/*
+The MIT License
+
+Copyright (c) 2012-2016 Denis Demidov <[email protected]>
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/**
+ * \file   vexcl/backend/compute/event.hpp
+ * \author Denis Demidov <[email protected]>
+ * \brief  Bring Boost.Compute events into vex::backend::compute namespace.
+ */
+
+#include <vexcl/backend/cuda/context.hpp>
+
+namespace vex {
+namespace backend {
+namespace cuda {
+
+namespace detail {
+
+template <>
+struct deleter_impl<CUevent> {
+    static void dispose(CUevent e) {
+        cuda_check( cuEventDestroy(e) );
+    }
+};
+
+} // namespace detail
+
+class event {
+    public:
+        event(const command_queue &q)
+            : q(q), e( create(q), detail::deleter(q.context().raw()) ) { }
+
+        CUevent raw() const { return e.get(); }
+
+        operator CUevent() const { return e.get(); }
+
+        void wait() const {
+            cuda_check( cuStreamWaitEvent(q.raw(), e.get(), 0) );
+        }
+    private:
+        command_queue q;
+        std::shared_ptr<std::remove_pointer<CUevent>::type> e;
+
+        static CUevent create(const command_queue &q) {
+            CUevent e;
+            q.context().set_current();
+
+            cuda_check( cuEventCreate(&e, CU_EVENT_DEFAULT) );
+            cuda_check( cuEventRecord(e, q.raw()) );
+
+            return e;
+        }
+};
+
+typedef std::vector<event> wait_list;
+
+inline void wait_list_append(wait_list &dst, const event &e) {
+    dst.push_back(e);
+}
+
+inline void wait_list_append(wait_list &dst, const wait_list &src) {
+    dst.insert(dst.begin(), src.begin(), src.end());
+}
+
+inline event enqueue_marker(const command_queue &q) {
+    return event(q);
+}
+
+} // namespace cuda
+} // namespace backend
+} // namespace vex
+
+
+#endif

vexcl/backend/cuda/kernel.hpp

@@ -36,6 +36,7 @@ THE SOFTWARE.
 #include <cuda.h>
 
 #include <vexcl/backend/cuda/compiler.hpp>
+#include <vexcl/backend/cuda/event.hpp>
 
 namespace vex {
 namespace backend {
@@ -98,11 +99,14 @@ class kernel {
         }
 
         /// Enqueue the kernel to the specified command queue.
-        void operator()(const command_queue &q) {
+        event operator()(const command_queue &q, const wait_list &events) {
             prm_addr.clear();
             for(auto p = prm_pos.begin(); p != prm_pos.end(); ++p)
                 prm_addr.push_back(stack.data() + *p);
 
+            for(auto e = events.begin(); e != events.end(); ++e)
+                cuda_check( cuEventSynchronize(e->raw()) );
+
             cuda_check(
                     cuLaunchKernel(
                         K,
@@ -117,15 +121,32 @@ class kernel {
 
             stack.clear();
             prm_pos.clear();
+
+            return event(q);
+        }
+
+        /// Enqueue the kernel to the specified command queue.
+        event operator()(const command_queue &q) {
+            return (*this)(q, wait_list());
         }
 
 #ifndef BOOST_NO_VARIADIC_TEMPLATES
         /// Enqueue the kernel to the specified command queue with the given arguments
         template <class Arg1, class... OtherArgs>
-        void operator()(const command_queue &q, Arg1 &&arg1, OtherArgs&&... other_args) {
-            push_arg(std::forward<Arg1>(arg1));
+        event operator()(const command_queue &q,
+                const Arg1 &arg1, const OtherArgs&... other_args)
+        {
+            push_arg(arg1);
+            return (*this)(q, other_args...);
+        }
 
-            (*this)(q, std::forward<OtherArgs>(other_args)...);
+        /// Enqueue the kernel to the specified command queue with the given arguments
+        template <class Arg1, class... OtherArgs>
+        event operator()(const command_queue &q, const wait_list &events,
+                const Arg1 &arg1, const OtherArgs&... other_args)
+        {
+            push_arg(arg1);
+            return (*this)(q, events, other_args...);
         }
 #endif