Merge pull request #84 from intelfx/device-specific-kernels

Groundwork for device-specific routines

9 files changed, 92 insertions, 83 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 77d1cd08..95d1d500 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -95,7 +95,7 @@ set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${CFLAGS}")
 # ==================================================================================================
 
 # Package scripts location
-set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_SOURCE_DIR}/cmake/Modules/")
+set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${clblast_SOURCE_DIR}/cmake/Modules/")
 
 # Requires OpenCL. It is found through the included "FindOpenCL.cmake" in CMAKE_MODULE_PATH.
 find_package(OpenCL REQUIRED)
diff --git a/src/clpp11.hpp b/src/clpp11.hpp
index af9d2ea4..d57223dd 100644
--- a/src/clpp11.hpp
+++ b/src/clpp11.hpp
@@ -109,7 +109,9 @@ class Event {
 
   // Accessor to the private data-member
   cl_event& operator()() { return *event_; }
+  const cl_event& operator()() const { return *event_; }
   cl_event* pointer() { return &(*event_); }
+  const cl_event* pointer() const { return &(*event_); }
  private:
   std::shared_ptr<cl_event> event_;
 };
@@ -686,30 +688,21 @@ class Kernel {
   // As above, but with an event waiting list
   void Launch(const Queue &queue, const std::vector<size_t> &global,
               const std::vector<size_t> &local, EventPointer event,
-              std::vector<Event>& waitForEvents) {
-    if (waitForEvents.size() == 0) { return Launch(queue, global, local, event); }
-
+              const std::vector<Event> &waitForEvents) {
     // Builds a plain version of the events waiting list
     auto waitForEventsPlain = std::vector<cl_event>();
     for (auto &waitEvent : waitForEvents) {
-      waitForEventsPlain.push_back(waitEvent());
+      if (waitEvent()) { waitForEventsPlain.push_back(waitEvent()); }
     }
 
     // Launches the kernel while waiting for other events
     CheckError(clEnqueueNDRangeKernel(queue(), *kernel_, static_cast<cl_uint>(global.size()),
-                                      nullptr, global.data(), local.data(),
+                                      nullptr, global.data(), !local.empty() ? local.data() : nullptr,
                                       static_cast<cl_uint>(waitForEventsPlain.size()),
-                                      waitForEventsPlain.data(),
+                                      !waitForEventsPlain.empty() ? waitForEventsPlain.data() : nullptr,
                                       event));
   }
 
-  // As above, but with the default local workgroup size
-  void Launch(const Queue &queue, const std::vector<size_t> &global, EventPointer event) {
-    CheckError(clEnqueueNDRangeKernel(queue(), *kernel_, static_cast<cl_uint>(global.size()),
-                                      nullptr, global.data(), nullptr,
-                                      0, nullptr, event));
-  }
-
   // Accessor to the private data-member
   const cl_kernel& operator()() const { return *kernel_; }
  private:
diff --git a/src/database/database.cpp b/src/database/database.cpp
index 6ec93731..ea1557b9 100644
--- a/src/database/database.cpp
+++ b/src/database/database.cpp
@@ -44,7 +44,7 @@ const std::vector<Database::DatabaseEntry> Database::database = {
 
 // Constructor, computing device properties and populating the parameter-vector from the database
 Database::Database(const Queue &queue, const std::vector<std::string> &kernels,
-                   const Precision precision):
+                   const Precision precision, const std::vector<DatabaseEntry> &overlay):
   parameters_{} {
 
   // Finds information of the current device
@@ -53,10 +53,23 @@ Database::Database(const Queue &queue, const std::vector<std::string> &kernels,
   auto device_vendor = device.Vendor();
   auto device_name = device.Name();
 
+  // Set the short vendor name
+  for (auto &combination : kVendorNames) {
+    if (device_vendor == combination.first) {
+      device_vendor = combination.second;
+    }
+  }
+
   // Iterates over all kernels to include, and retrieves the parameters for each of them
   for (auto &kernel: kernels) {
-    auto search_result = Search(kernel, device_type, device_vendor, device_name, precision);
-    parameters_.insert(search_result.begin(), search_result.end());
+    auto search_result = ParametersPtr{};
+
+    for (auto db: { &overlay, &database }) {
+      search_result = Search(kernel, device_type, device_vendor, device_name, precision, *db);
+      if (search_result) { parameters_.insert(search_result->begin(), search_result->end()); break; }
+    }
+
+    if (!search_result) { throw std::runtime_error("Database error, could not find a suitable entry"); }
   }
 }
 
@@ -74,27 +87,21 @@ std::string Database::GetDefines() const {
 // =================================================================================================
 
 // Searches the database for the right kernel and precision
-Database::Parameters Database::Search(const std::string &this_kernel,
-                                      const std::string &this_type,
-                                      const std::string &this_vendor,
-                                      const std::string &this_device,
-                                      const Precision this_precision) const {
-  // Set the short vendor name
-  auto this_short_vendor = this_vendor;
-  for (auto &combination : kVendorNames) {
-    if (this_vendor == combination.first) {
-      this_short_vendor = combination.second;
-    }
-  }
+Database::ParametersPtr Database::Search(const std::string &this_kernel,
+                                         const std::string &this_type,
+                                         const std::string &this_vendor,
+                                         const std::string &this_device,
+                                         const Precision this_precision,
+                                         const std::vector<DatabaseEntry> &this_database) const {
 
   // Selects the right kernel
-  for (auto &db: database) {
+  for (auto &db: this_database) {
     if (db.kernel == this_kernel && db.precision == this_precision) {
 
       // Searches for the right vendor and device type, or selects the default if unavailable. This
       // assumes that the default vendor / device type is last in the database.
       for (auto &vendor: db.vendors) {
-        if ((vendor.name == this_short_vendor || vendor.name == kDeviceVendorAll) &&
+        if ((vendor.name == this_vendor || vendor.name == kDeviceVendorAll) &&
             (vendor.type == this_type || vendor.type == kDeviceTypeAll)) {
 
           // Searches for the right device. If the current device is unavailable, selects the vendor
@@ -104,7 +111,7 @@ Database::Parameters Database::Search(const std::string &this_kernel,
             if (device.name == this_device || device.name == "default") {
 
               // Sets the parameters accordingly
-              return device.parameters;
+              return &device.parameters;
             }
           }
         }
@@ -113,7 +120,7 @@ Database::Parameters Database::Search(const std::string &this_kernel,
   }
 
   // If we reached this point, something is wrong
-  throw std::runtime_error("Database error, could not find a suitable entry");
+  return nullptr;
 }
 
 // =================================================================================================
diff --git a/src/database/database.hpp b/src/database/database.hpp
index 0987cbed..5a61fad9 100644
--- a/src/database/database.hpp
+++ b/src/database/database.hpp
@@ -32,6 +32,7 @@ class Database {
 
   // Type alias for the database parameters
   using Parameters = std::unordered_map<std::string,size_t>;
+  using ParametersPtr = const Parameters*;
 
   // Structures for content inside the database
   struct DatabaseDevice {
@@ -78,9 +79,9 @@ class Database {
   static const DatabaseEntry PadtransposeHalf, PadtransposeSingle, PadtransposeDouble, PadtransposeComplexSingle, PadtransposeComplexDouble;
   static const std::vector<DatabaseEntry> database;
 
-  // The constructor
+  // The constructor with a user-provided database overlay
   explicit Database(const Queue &queue, const std::vector<std::string> &routines,
-                    const Precision precision);
+                    const Precision precision, const std::vector<DatabaseEntry> &overlay);
 
   // Accessor of values by key
   size_t operator[](const std::string key) const { return parameters_.find(key)->second; }
@@ -93,6 +94,11 @@ class Database {
                     const std::string &this_vendor, const std::string &this_device,
                     const Precision this_precision) const;
 
+  // Alternate search method in a specified database, returning pointer (possibly NULL)
+  ParametersPtr Search(const std::string &this_kernel, const std::string &this_type,
+                       const std::string &this_vendor, const std::string &this_device,
+                       const Precision this_precision, const std::vector<DatabaseEntry> &db) const;
+
   // Found parameters suitable for this device/kernel
   Parameters parameters_;
 };
diff --git a/src/routine.cpp b/src/routine.cpp
index 3c3343da..189ae190 100644
--- a/src/routine.cpp
+++ b/src/routine.cpp
@@ -22,7 +22,8 @@ namespace clblast {
 
 // Constructor: not much here, because no status codes can be returned
 Routine::Routine(Queue &queue, EventPointer event, const std::string &name,
-                 const std::vector<std::string> &routines, const Precision precision):
+                 const std::vector<std::string> &routines, const Precision precision,
+                 const std::vector<Database::DatabaseEntry> &userDatabase):
     precision_(precision),
     routine_name_(name),
     queue_(queue),
@@ -30,7 +31,7 @@ Routine::Routine(Queue &queue, EventPointer event, const std::string &name,
     context_(queue_.GetContext()),
     device_(queue_.GetDevice()),
     device_name_(device_.Name()),
-    db_(queue_, routines, precision_) {
+    db_(queue_, routines, precision_, userDatabase) {
 }
 
 // =================================================================================================
diff --git a/src/routine.hpp b/src/routine.hpp
index 54b5779f..21506e7b 100644
--- a/src/routine.hpp
+++ b/src/routine.hpp
@@ -34,7 +34,8 @@ class Routine {
 
   // Base class constructor
   explicit Routine(Queue &queue, EventPointer event, const std::string &name,
-                   const std::vector<std::string> &routines, const Precision precision);
+                   const std::vector<std::string> &routines, const Precision precision,
+                   const std::vector<Database::DatabaseEntry> &userDatabase = {});
 
   // Set-up phase of the kernel
   StatusCode SetUp();
diff --git a/src/routines/common.cpp b/src/routines/common.cpp
index 2e82e04d..3969cf9f 100644
--- a/src/routines/common.cpp
+++ b/src/routines/common.cpp
@@ -22,23 +22,25 @@ namespace clblast {
 // Enqueues a kernel, waits for completion, and checks for errors
 StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device,
                      std::vector<size_t> global, const std::vector<size_t> &local,
-                     EventPointer event, std::vector<Event>& waitForEvents) {
+                     EventPointer event, const std::vector<Event> &waitForEvents) {
 
-  // Tests for validity of the local thread sizes
-  if (local.size() > device.MaxWorkItemDimensions()) {
-    return StatusCode::kInvalidLocalNumDimensions; 
-  }
-  const auto max_work_item_sizes = device.MaxWorkItemSizes();
-  for (auto i=size_t{0}; i<local.size(); ++i) {
-    if (local[i] > max_work_item_sizes[i]) { return StatusCode::kInvalidLocalThreadsDim; }
-  }
-  auto local_size = size_t{1};
-  for (auto &item: local) { local_size *= item; }
-  if (local_size > device.MaxWorkGroupSize()) { return StatusCode::kInvalidLocalThreadsTotal; }
+  if (!local.empty()) {
+    // Tests for validity of the local thread sizes
+    if (local.size() > device.MaxWorkItemDimensions()) {
+      return StatusCode::kInvalidLocalNumDimensions;
+    }
+    const auto max_work_item_sizes = device.MaxWorkItemSizes();
+    for (auto i=size_t{0}; i<local.size(); ++i) {
+      if (local[i] > max_work_item_sizes[i]) { return StatusCode::kInvalidLocalThreadsDim; }
+    }
+    auto local_size = size_t{1};
+    for (auto &item: local) { local_size *= item; }
+    if (local_size > device.MaxWorkGroupSize()) { return StatusCode::kInvalidLocalThreadsTotal; }
 
-  // Make sure the global thread sizes are at least equal to the local sizes
-  for (auto i=size_t{0}; i<global.size(); ++i) {
-    if (global[i] < local[i]) { global[i] = local[i]; }
+    // Make sure the global thread sizes are at least equal to the local sizes
+    for (auto i=size_t{0}; i<global.size(); ++i) {
+      if (global[i] < local[i]) { global[i] = local[i]; }
+    }
   }
 
   // Tests for local memory usage
@@ -69,13 +71,5 @@ StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device,
   return StatusCode::kSuccess;
 }
 
-// As above, but without an event waiting list
-StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device,
-                     std::vector<size_t> global, const std::vector<size_t> &local,
-                     EventPointer event) {
-  auto emptyWaitingList = std::vector<Event>();
-  return RunKernel(kernel, queue, device, global, local, event, emptyWaitingList);
-}
-
 // =================================================================================================
 } // namespace clblast
diff --git a/src/routines/common.hpp b/src/routines/common.hpp
index d53bdc25..9d8849c3 100644
--- a/src/routines/common.hpp
+++ b/src/routines/common.hpp
@@ -29,12 +29,7 @@ namespace clblast {
 // Enqueues a kernel, waits for completion, and checks for errors
 StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device,
                      std::vector<size_t> global, const std::vector<size_t> &local,
-                     EventPointer event, std::vector<Event>& waitForEvents);
-
-// As above, but without an event waiting list
-StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device,
-                     std::vector<size_t> global, const std::vector<size_t> &local,
-                     EventPointer event);
+                     EventPointer event, const std::vector<Event> &waitForEvents = {});
 
 // =================================================================================================
 
@@ -43,7 +38,7 @@ StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device,
 template <typename T>
 StatusCode PadCopyTransposeMatrix(Queue &queue, const Device &device,
                                   const Database &db,
-                                  EventPointer event, std::vector<Event>& waitForEvents,
+                                  EventPointer event, const std::vector<Event> &waitForEvents,
                                   const size_t src_one, const size_t src_two,
                                   const size_t src_ld, const size_t src_offset,
                                   const Buffer<T> &src,
diff --git a/src/routines/level3/xgemm.cpp b/src/routines/level3/xgemm.cpp
index 0db28537..fce59622 100644
--- a/src/routines/level3/xgemm.cpp
+++ b/src/routines/level3/xgemm.cpp
@@ -63,9 +63,12 @@ StatusCode Xgemm<T>::DoGemm(const Layout layout,
   const auto b_rotated = (layout == Layout::kColMajor && b_transpose != Transpose::kNo) ||
                          (layout == Layout::kRowMajor && b_transpose == Transpose::kNo);
   const auto c_rotated = (layout == Layout::kRowMajor);
-  const auto a_do_transpose =  a_rotated;
-  const auto b_do_transpose = !b_rotated;
-  const auto c_do_transpose =  c_rotated;
+  static const auto a_want_rotated = false;
+  static const auto b_want_rotated = true;
+  static const auto c_want_rotated = false;
+  const auto a_do_transpose = a_rotated != a_want_rotated;
+  const auto b_do_transpose = b_rotated != b_want_rotated;
+  const auto c_do_transpose = c_rotated != c_want_rotated;
 
   // In case of complex data-types, the transpose can also become a conjugate transpose
   const auto a_conjugate = (a_transpose == Transpose::kConjugate);
@@ -99,6 +102,15 @@ StatusCode Xgemm<T>::DoGemm(const Layout layout,
   const auto n_ceiled = Ceil(n, db_["NWG"]);
   const auto k_ceiled = Ceil(k, db_["KWG"]);
 
+  // Computes the first and second "internal" (ceiled) dimensions of the 3 matrices taking into account
+  // whether the matrices need to be rotated or not for the kernel.
+  const auto a_one_i = (a_want_rotated) ? k_ceiled : m_ceiled;
+  const auto a_two_i = (a_want_rotated) ? m_ceiled : k_ceiled;
+  const auto b_one_i = (b_want_rotated) ? n_ceiled : k_ceiled;
+  const auto b_two_i = (b_want_rotated) ? k_ceiled : n_ceiled;
+  const auto c_one_i = (c_want_rotated) ? n_ceiled : m_ceiled;
+  const auto c_two_i = (c_want_rotated) ? m_ceiled : n_ceiled;
+
   // The padded/transposed input/output matrices: if memory allocation fails, throw an exception
   try {
 
@@ -106,17 +118,17 @@ StatusCode Xgemm<T>::DoGemm(const Layout layout,
     const auto program = GetProgramFromCache(context_, PrecisionValue<T>(), routine_name_);
 
     // Determines whether or not temporary matrices are needed
-    auto a_no_temp = a_one == m_ceiled && a_two == k_ceiled && a_ld == m_ceiled && a_offset == 0 &&
+    auto a_no_temp = a_one == a_one_i && a_two == a_two_i && a_ld == a_one && a_offset == 0 &&
                      a_do_transpose == false && a_conjugate == false;
-    auto b_no_temp = b_one == n_ceiled && b_two == k_ceiled && b_ld == n_ceiled && b_offset == 0 &&
+    auto b_no_temp = b_one == b_one_i && b_two == b_two_i && b_ld == b_one && b_offset == 0 &&
                      b_do_transpose == false && b_conjugate == false;
-    auto c_no_temp = c_one == m_ceiled && c_two == n_ceiled && c_ld == m_ceiled && c_offset == 0 &&
+    auto c_no_temp = c_one == c_one_i && c_two == c_two_i && c_ld == c_one && c_offset == 0 &&
                      c_do_transpose == false;
 
     // Creates the temporary matrices
-    const auto a_temp = (a_no_temp) ? a_buffer : Buffer<T>(context_, k_ceiled*m_ceiled);
-    const auto b_temp = (b_no_temp) ? b_buffer : Buffer<T>(context_, k_ceiled*n_ceiled);
-    const auto c_temp = (c_no_temp) ? c_buffer : Buffer<T>(context_, m_ceiled*n_ceiled);
+    const auto a_temp = (a_no_temp) ? a_buffer : Buffer<T>(context_, a_one_i*a_two_i);
+    const auto b_temp = (b_no_temp) ? b_buffer : Buffer<T>(context_, b_one_i*b_two_i);
+    const auto c_temp = (c_no_temp) ? c_buffer : Buffer<T>(context_, c_one_i*c_two_i);
 
     // Events of all kernels (including pre/post processing kernels)
     auto eventWaitList = std::vector<Event>();
@@ -129,7 +141,7 @@ StatusCode Xgemm<T>::DoGemm(const Layout layout,
       auto eventProcessA = Event();
       status = PadCopyTransposeMatrix(queue_, device_, db_, eventProcessA.pointer(), emptyEventList,
                                       a_one, a_two, a_ld, a_offset, a_buffer,
-                                      m_ceiled, k_ceiled, m_ceiled, 0, a_temp,
+                                      a_one_i, a_two_i, a_one_i, 0, a_temp,
                                       ConstantOne<T>(), program,
                                       true, a_do_transpose, a_conjugate);
       if (ErrorIn(status)) { return status; }
@@ -141,7 +153,7 @@ StatusCode Xgemm<T>::DoGemm(const Layout layout,
       auto eventProcessB = Event();
       status = PadCopyTransposeMatrix(queue_, device_, db_, eventProcessB.pointer(), emptyEventList,
                                       b_one, b_two, b_ld, b_offset, b_buffer,
-                                      n_ceiled, k_ceiled, n_ceiled, 0, b_temp,
+                                      b_one_i, b_two_i, b_one_i, 0, b_temp,
                                       ConstantOne<T>(), program,
                                       true, b_do_transpose, b_conjugate);
       if (ErrorIn(status)) { return status; }
@@ -153,7 +165,7 @@ StatusCode Xgemm<T>::DoGemm(const Layout layout,
       auto eventProcessC = Event();
       status = PadCopyTransposeMatrix(queue_, device_, db_, eventProcessC.pointer(), emptyEventList,
                                       c_one, c_two, c_ld, c_offset, c_buffer,
-                                      m_ceiled, n_ceiled, m_ceiled, 0, c_temp,
+                                      c_one_i, c_two_i, c_one_i, 0, c_temp,
                                       ConstantOne<T>(), program,
                                       true, c_do_transpose, false);
       if (ErrorIn(status)) { return status; }
@@ -176,8 +188,8 @@ StatusCode Xgemm<T>::DoGemm(const Layout layout,
 
       // Computes the global and local thread sizes
       const auto global = std::vector<size_t>{
-        (m_ceiled * db_["MDIMC"]) / db_["MWG"],
-        (n_ceiled * db_["NDIMC"]) / db_["NWG"]
+        (c_one_i * db_["MDIMC"]) / db_["MWG"],
+        (c_two_i * db_["NDIMC"]) / db_["NWG"]
       };
       const auto local = std::vector<size_t>{db_["MDIMC"], db_["NDIMC"]};
 
@@ -191,7 +203,7 @@ StatusCode Xgemm<T>::DoGemm(const Layout layout,
       if (!c_no_temp) {
         eventWaitList.push_back(eventKernel);
         status = PadCopyTransposeMatrix(queue_, device_, db_, event_, eventWaitList,
-                                        m_ceiled, n_ceiled, m_ceiled, 0, c_temp,
+                                        c_one_i, c_two_i, c_one_i, 0, c_temp,
                                         c_one, c_two, c_ld, c_offset, c_buffer,
                                         ConstantOne<T>(), program,
                                         false, c_do_transpose, false);