A class to assist in readying buffers via the Node::readyBuffers() method. More...
#include <Kokkos_NodeHelpers.hpp>
A class to assist in readying buffers via the Node::readyBuffers() method.
Kokkos asks that you call the Node's readyBuffers() method on all compute buffers that you plan to pass into a kernel. In a debug build, this checks to make sure that the ArrayRCP objects are really compute buffers. ReadyBufferHelper provides a transaction-like interface to help you with this. The ReadyBufferHelper also keeps the ArrayRCP objects for you during its lifetime. This prevents them from falling out of scope and getting deleted, thus ensuring that when you extract the raw pointer to give to the Kokkos kernel, the pointer will be valid throughout the lifetime of the kernel.
Here is an example of how to use ReadyBufferHelper with a simple Kokkos kernel. Suppose you have a
parallel_for kernel Op that takes a const array and a nonconst array:
template<class T> class Op { public: Op (const T* x, T* y, size_t n) : x_ (x), y_ (y), n_ (n) {} inline KERNEL_PREFIX void execute (size_t i) { y_[i] += 2.0 * x_[i]; } private: const T* x_; T* y_; size_t n_; };
Here's how you would use ReadyBufferHelper when invoking the kernel in parallel:
using Teuchos::ArrayRCP; using Teuchos::RCP; // Suppose node_type is the Kokkos Node type. RCP<node_type> node = ...; // x and y are compute buffers that you got elsewhere. ArrayRCP<const T> x = ...; ArrayRCP<T> y = ...; const size_t n = x.size (); ReadyBufferHelper<node_type> rbh (node); rbh.begin (); // Start adding compute buffers // Create the Kokkos kernel "work-data pair" (Kokkos' phrase for // what other programming languages call "closure"). When adding // a const buffer (ArrayRCP<const T>), omit the "const" when naming // the type in the template method call. // // If T is a template parameter in this scope, you need the "template" // keyword when invoking ReadyBufferHelper template methods. You don't // need it if T is a concrete type in this scope. Op wdp (rbh.template addConstBuffer<T> (x), rbh.template addNonConstBuffer<T> (y), n); rbh.end (); // Done adding compute buffers // Invoke the kernel in parallel over the range 0, ..., n-1. // The ReadyBufferHelper will protect x and y from falling out // of scope. // // If T is a template parameter in this scope, you need the "template" // keyword. You don't need it if T is a concrete type in this scope. node->template parallel_for<Op<T> > (0, x.size (), wdp);
Definition at line 126 of file Kokkos_NodeHelpers.hpp.
The node via which buffers are being readied.
Definition at line 158 of file Kokkos_NodeHelpers.hpp.
Destructor.
Definition at line 163 of file Kokkos_NodeHelpers.hpp.
Tell this object that you are ready to start adding buffers.
Definition at line 167 of file Kokkos_NodeHelpers.hpp.
Add a const buffer, and return its raw pointer.
Definition at line 174 of file Kokkos_NodeHelpers.hpp.
Add a non-const buffer, and return its raw pointer.
Definition at line 181 of file Kokkos_NodeHelpers.hpp.
Tell this object that you are done adding buffers.
Definition at line 187 of file Kokkos_NodeHelpers.hpp.
The Kokkos Node instance for which to add buffers.
Definition at line 150 of file Kokkos_NodeHelpers.hpp.
List of compute buffers that were added with addConstBuffer().
Definition at line 152 of file Kokkos_NodeHelpers.hpp.
List of compute buffers that were added with addNonConstBuffer().
Definition at line 154 of file Kokkos_NodeHelpers.hpp.