Merge pull request #19 from Polytonic/development

Add Profiling Support

chlorine/chlorine.hpp

@@ -41,31 +41,32 @@ namespace ch
 
         // Handle the Base Case
         template<unsigned int const argn = 0>
-        void call(std::string const & kernel_function);
+        cl::Event call(std::string const & kernel_function);
 
         // Handle Primitive Types
-        template<unsigned int const argn = 0, typename T, typename ... Params>
-        typename std::enable_if<std::is_arithmetic<T>::value>::type
-        call(std::string const & kernel_function, T primitive, Params && ... parameters);
+        template<unsigned int const argn = 0, typename T,
+                 typename std::enable_if<std::is_arithmetic<T>::value>::type* = nullptr, typename ... Params>
+        cl::Event call(std::string const & kernel_function, T primitive, Params && ... parameters);
 
         // Handle C-Style Arrays
         template<unsigned int const argn = 0, class T, size_t const N, typename ... Params>
-        void call(std::string const & kernel_function, T (&array) [N], Params && ... parameters);
+        cl::Event call(std::string const & kernel_function, T (&array) [N], Params && ... parameters);
 
         // Handle STL Arrays
         template<unsigned int const argn = 0, class T, size_t const N, typename ... Params>
-        void call(std::string const & kernel_function, std::array<T, N> & array, Params && ... parameters);
+        cl::Event call(std::string const & kernel_function, std::array<T, N> & array, Params && ... parameters);
 
         // Handle Other STL Containers
         template<unsigned int const argn = 0, template<typename ...> class V, typename T, typename ... Params>
-        void call(std::string const & kernel_function, V<T> & array, Params && ... parameters);
+        cl::Event call(std::string const & kernel_function, V<T> & array, Params && ... parameters);
 
     private:
 
         // Components
         cl::CommandQueue mQueue;
         cl::Context      mContext;
         cl::Device       mDevice;
+        cl::Event        mEvent;
         cl::Platform     mPlatform;
         cl::Program      mProgram;
 
@@ -93,6 +94,13 @@ namespace ch
         return worker;
     }
 
+    // Determine the Elapsed Computation Time (ns)
+    unsigned int elapsed(cl::Event const & event)
+    {
+        return event.getProfilingInfo<CL_PROFILING_COMMAND_END>()
+             - event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
+    }
+
     // Read the Contents of the Given Filename
     std::string read(std::string const & filename)
     {
@@ -155,59 +163,62 @@ namespace ch
 
     // Handle the Base Case
     template<unsigned int const argn>
-    void Worker::call(std::string const & kernel_function)
+    cl::Event Worker::call(std::string const & kernel_function)
     {
         // Perform the Calculation and Read Data from Memory Buffers
-        mQueue.enqueueNDRangeKernel(mKernels[kernel_function], mOffset, mGlobal, mLocal);
+        mQueue.enqueueNDRangeKernel(mKernels[kernel_function], mOffset, mGlobal, mLocal, NULL, & mEvent);
         for (auto &i : mBuffers)
             mQueue.enqueueUnmapMemObject(i.first,
             mQueue.enqueueMapBuffer(i.first, CL_TRUE, CL_MAP_READ, 0, i.second));
             mBuffers.clear();
+
+        // Return OpenCL Event Object Containing Profiling Data
+        return mEvent;
     }
 
     // Handle Primitive Types
-    template<unsigned int const argn, typename T, typename ... Params>
-    typename std::enable_if<std::is_arithmetic<T>::value>::type
-    Worker::call(std::string const & kernel_function, T primitive, Params && ... parameters)
+    template<unsigned int const argn, typename T,
+             typename std::enable_if<std::is_arithmetic<T>::value>::type*, typename ... Params>
+    cl::Event Worker::call(std::string const & kernel_function, T primitive, Params && ... parameters)
     {
         mKernels[kernel_function].setArg(argn, primitive);
-        call<argn+1>(kernel_function, parameters...);
+        return call<argn+1>(kernel_function, parameters...);
     }
 
     // Handle C-Style Arrays
     template<unsigned int const argn, class T, size_t const N, typename ... Params>
-    void Worker::call(std::string const & kernel_function, T (&array) [N], Params && ... parameters)
+    cl::Event Worker::call(std::string const & kernel_function, T (&array) [N], Params && ... parameters)
     {
         size_t array_size = N * sizeof(array[0]);
         if (N > mGlobal[0]) { mGlobal = cl::NDRange(N); }
         cl::Buffer buffer = cl::Buffer(mContext, CL_MEM_USE_HOST_PTR, array_size, & array[0]);
         mBuffers.push_back(std::make_pair(buffer, array_size));
         mKernels[kernel_function].setArg(argn, buffer);
-        call<argn+1>(kernel_function, parameters...);
+        return call<argn+1>(kernel_function, parameters...);
     }
 
     // Handle STL Arrays
     template<unsigned int const argn, class T, size_t const N, typename ... Params>
-    void Worker::call(std::string const & kernel_function, std::array<T, N> & array, Params && ... parameters)
+    cl::Event Worker::call(std::string const & kernel_function, std::array<T, N> & array, Params && ... parameters)
     {
         size_t array_size = array.size() * sizeof(T);
         if (array.size() > mGlobal[0]) { mGlobal = cl::NDRange(array.size()); }
         cl::Buffer buffer = cl::Buffer(mContext, CL_MEM_USE_HOST_PTR, array_size, & array[0]);
         mBuffers.push_back(std::make_pair(buffer, array_size));
         mKernels[kernel_function].setArg(argn, buffer);
-        call<argn+1>(kernel_function, parameters...);
+        return call<argn+1>(kernel_function, parameters...);
     }
 
     // Handle Other STL Containers
     template<unsigned int const argn, template<typename ...> class V, typename T, typename ... Params>
-    void Worker::call(std::string const & kernel_function, V<T> & array, Params && ... parameters)
+    cl::Event Worker::call(std::string const & kernel_function, V<T> & array, Params && ... parameters)
     {
         size_t array_size = array.size() * sizeof(T);
         if (array.size() > mGlobal[0]) { mGlobal = cl::NDRange(array.size()); }
         cl::Buffer buffer = cl::Buffer(mContext, CL_MEM_USE_HOST_PTR, array_size, & array[0]);
         mBuffers.push_back(std::make_pair(buffer, array_size));
         mKernels[kernel_function].setArg(argn, buffer);
-        call<argn+1>(kernel_function, parameters...);
+        return call<argn+1>(kernel_function, parameters...);
     }
 };
 

examples/swap/readme.md

@@ -53,3 +53,22 @@ $ clang++ -std=c++11 swap.cpp -lOpenCL
 ```
 
 Kernel files are written in a variant of the C programming language. While I won't go into detail about it here, I hope this serves as a valuable demonstration in how Chlorine may be used to easily port code to run in parallel.
+
+### Profiling
+
+```
+// Original Call
+worker.call("swap", spam, eggs);
+
+// New Call
+auto event = worker.call("swap", spam, eggs);
+```
+
+As of 1d776984a11466cfd742bec2af0ff7b278a4479a, Chlorine now returns the OpenCL event associated with the kernel function call. This allows you to recover profiling data, such as how much time was spent executing the kernel function, as well as the idling time. For simplicity, we return the entire OpenCL object rather than adding a new function call to our API. This should be a non-breaking change for the most part.
+
+```
+// Print Some Profiling Data
+std::cout << "Elapsed Time: " << ch::elapsed(event) << "ns\n";
+```
+
+To make things easier, we also added a new helper function `ch::elapsed()` which accepts an OpenCL event and returns the elapsed time spent on your kernel function. This helper preserves the nanosecond resolution offered by the OpenCL API and is merely a convenience wrapper.

examples/swap/swap.cpp

@@ -10,9 +10,12 @@ int main(int argc, char * argv[])
     ch::Worker worker("swap.cl");
 
     // Call the Swap Function in the Given Kernel
-    worker.call("swap", spam, eggs);
+    auto event = worker.call("swap", spam, eggs);
 
     // Host Containers Are Automatically Updated
     std::cout << "Spam: " << spam[0] << "\n"; // 2.7182
     std::cout << "Eggs: " << eggs[0] << "\n"; // 3.1415
+
+    // Print Some Profiling Data
+    std::cout << "Elapsed Time: " << ch::elapsed(event) << "ns\n";
 }

readme.md

@@ -23,11 +23,14 @@ int main(int argc, char * argv[])
     ch::Worker worker("swap.cl");
 
     // Call the Swap Function in the Given Kernel
-    worker.call("swap", spam, eggs);
+    auto event = worker.call("swap", spam, eggs);
 
     // Host Containers Are Automatically Updated
     std::cout << "Spam: " << spam[0] << "\n"; // 2.7182
     std::cout << "Eggs: " << eggs[0] << "\n"; // 3.1415
+
+    // Print Some Profiling Data
+    std::cout << "Elapsed Time: " << ch::elapsed(event) << "ns\n";
 }
 ```
 **swap.cl**

tests/test_helpers.cl

@@ -1,6 +1,6 @@
 // Test the Kernel Read Function
-__kernel void read(__global int * a)
+__kernel void fill(__global int * a)
 {
     unsigned int i = get_global_id(0);
-    a[i] = 0;
+    a[i] = 1;
 }

tests/test_main.cpp

@@ -78,6 +78,16 @@ TEST_CASE("Vectors", "[arithmetic, vectors]")
 // Test Chlorine Worker Helpers
 TEST_CASE("Helpers", "[helpers]")
 {
+    SECTION("ch::elapsed() Returns an Elapsed Time")
+    {
+        ch::Worker worker("tests/test_helpers.cl");
+        std::vector<int> a(100, 100);
+        auto event = worker.call("fill", a);
+        for (unsigned int i = 0; i < a.size(); i++)
+            CHECK(a[i] == 1);
+        CHECK(ch::elapsed(event) > 0);
+    }
+
     SECTION("ch::read() Returns File Contents")
     {
         // Read Kernel Source
@@ -86,10 +96,10 @@ TEST_CASE("Helpers", "[helpers]")
         // Define String Literal to Match
         std::string match =
             "// Test the Kernel Read Function\n"
-            "__kernel void read(__global int * a)\n"
+            "__kernel void fill(__global int * a)\n"
             "{\n"
             "    unsigned int i = get_global_id(0);\n"
-            "    a[i] = 0;\n"
+            "    a[i] = 1;\n"
             "}\n";
 
         // Assert Kernel Source is Equivalent