From a1d80e7402d25a4db73d028f4acba660e3b6f19e Mon Sep 17 00:00:00 2001 From: Ivan Shapovalov Date: Thu, 7 Jul 2016 05:07:59 +0300 Subject: CMakeLists.txt: use ${clblast_SOURCE_DIR} instead of ${CMAKE_SOURCE_DIR} --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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) -- cgit v1.2.3 From 1ae71614ac227e9b80678505fcb5ecc18a3d3383 Mon Sep 17 00:00:00 2001 From: Ivan Shapovalov Date: Sun, 10 Jul 2016 23:11:27 +0300 Subject: xgemm: do not hardcode kernel requirements for internal matrix layout Do not hardcode the knowledge about "A and C col-major, B row-major". This allows for easier reuse of the DoGemm() routine with different kernels. --- src/routines/level3/xgemm.cpp | 42 +++++++++++++++++++++++++++--------------- 1 file changed, 27 insertions(+), 15 deletions(-) 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::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::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::DoGemm(const Layout layout, const auto program = GetProgramFromCache(context_, PrecisionValue(), 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(context_, k_ceiled*m_ceiled); - const auto b_temp = (b_no_temp) ? b_buffer : Buffer(context_, k_ceiled*n_ceiled); - const auto c_temp = (c_no_temp) ? c_buffer : Buffer(context_, m_ceiled*n_ceiled); + const auto a_temp = (a_no_temp) ? a_buffer : Buffer(context_, a_one_i*a_two_i); + const auto b_temp = (b_no_temp) ? b_buffer : Buffer(context_, b_one_i*b_two_i); + const auto c_temp = (c_no_temp) ? c_buffer : Buffer(context_, c_one_i*c_two_i); // Events of all kernels (including pre/post processing kernels) auto eventWaitList = std::vector(); @@ -129,7 +141,7 @@ StatusCode Xgemm::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(), program, true, a_do_transpose, a_conjugate); if (ErrorIn(status)) { return status; } @@ -141,7 +153,7 @@ StatusCode Xgemm::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(), program, true, b_do_transpose, b_conjugate); if (ErrorIn(status)) { return status; } @@ -153,7 +165,7 @@ StatusCode Xgemm::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(), program, true, c_do_transpose, false); if (ErrorIn(status)) { return status; } @@ -176,8 +188,8 @@ StatusCode Xgemm::DoGemm(const Layout layout, // Computes the global and local thread sizes const auto global = std::vector{ - (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{db_["MDIMC"], db_["NDIMC"]}; @@ -191,7 +203,7 @@ StatusCode Xgemm::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(), program, false, c_do_transpose, false); -- cgit v1.2.3 From 2dd5ee3f758f817238cf74c068c5cad6eb3d46dd Mon Sep 17 00:00:00 2001 From: Ivan Shapovalov Date: Mon, 11 Jul 2016 22:36:39 +0300 Subject: clblast::RunKernel, cl::Kernel: take const vector as waitForEvents --- src/clpp11.hpp | 4 +++- src/routines/common.cpp | 2 +- src/routines/common.hpp | 4 ++-- 3 files changed, 6 insertions(+), 4 deletions(-) diff --git a/src/clpp11.hpp b/src/clpp11.hpp index af9d2ea4..9fa683ac 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 event_; }; @@ -686,7 +688,7 @@ class Kernel { // As above, but with an event waiting list void Launch(const Queue &queue, const std::vector &global, const std::vector &local, EventPointer event, - std::vector& waitForEvents) { + const std::vector &waitForEvents) { if (waitForEvents.size() == 0) { return Launch(queue, global, local, event); } // Builds a plain version of the events waiting list diff --git a/src/routines/common.cpp b/src/routines/common.cpp index 2e82e04d..21e16954 100644 --- a/src/routines/common.cpp +++ b/src/routines/common.cpp @@ -22,7 +22,7 @@ namespace clblast { // Enqueues a kernel, waits for completion, and checks for errors StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device, std::vector global, const std::vector &local, - EventPointer event, std::vector& waitForEvents) { + EventPointer event, const std::vector &waitForEvents) { // Tests for validity of the local thread sizes if (local.size() > device.MaxWorkItemDimensions()) { diff --git a/src/routines/common.hpp b/src/routines/common.hpp index d53bdc25..d0bbc707 100644 --- a/src/routines/common.hpp +++ b/src/routines/common.hpp @@ -29,7 +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 global, const std::vector &local, - EventPointer event, std::vector& waitForEvents); + EventPointer event, const std::vector &waitForEvents); // As above, but without an event waiting list StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device, @@ -43,7 +43,7 @@ StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device, template StatusCode PadCopyTransposeMatrix(Queue &queue, const Device &device, const Database &db, - EventPointer event, std::vector& waitForEvents, + EventPointer event, const std::vector &waitForEvents, const size_t src_one, const size_t src_two, const size_t src_ld, const size_t src_offset, const Buffer &src, -- cgit v1.2.3 From 5502c5eec4c7b56c3b44ef04046d1621d58be47f Mon Sep 17 00:00:00 2001 From: Ivan Shapovalov Date: Mon, 11 Jul 2016 22:55:58 +0300 Subject: cl::Kernel: skip NULL entries in waitForEvents --- src/clpp11.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/clpp11.hpp b/src/clpp11.hpp index 9fa683ac..debfff09 100644 --- a/src/clpp11.hpp +++ b/src/clpp11.hpp @@ -689,14 +689,14 @@ class Kernel { void Launch(const Queue &queue, const std::vector &global, const std::vector &local, EventPointer event, const std::vector &waitForEvents) { - if (waitForEvents.size() == 0) { return Launch(queue, global, local, event); } - // Builds a plain version of the events waiting list auto waitForEventsPlain = std::vector(); for (auto &waitEvent : waitForEvents) { - waitForEventsPlain.push_back(waitEvent()); + if (waitEvent()) { waitForEventsPlain.push_back(waitEvent()); } } + if (waitForEvents.size() == 0) { return Launch(queue, global, local, event); } + // Launches the kernel while waiting for other events CheckError(clEnqueueNDRangeKernel(queue(), *kernel_, static_cast(global.size()), nullptr, global.data(), local.data(), -- cgit v1.2.3 From ae3299da302ba2c26f1e3490a1a7ee389d91feee Mon Sep 17 00:00:00 2001 From: Ivan Shapovalov Date: Mon, 11 Jul 2016 23:14:43 +0300 Subject: clblast::RunKernel, cl::Kernel: unify variants with/without waitForEvents, support empty LWS --- src/clpp11.hpp | 13 ++----------- src/routines/common.cpp | 38 ++++++++++++++++---------------------- src/routines/common.hpp | 7 +------ 3 files changed, 19 insertions(+), 39 deletions(-) diff --git a/src/clpp11.hpp b/src/clpp11.hpp index debfff09..d57223dd 100644 --- a/src/clpp11.hpp +++ b/src/clpp11.hpp @@ -695,23 +695,14 @@ class Kernel { if (waitEvent()) { waitForEventsPlain.push_back(waitEvent()); } } - if (waitForEvents.size() == 0) { return Launch(queue, global, local, event); } - // Launches the kernel while waiting for other events CheckError(clEnqueueNDRangeKernel(queue(), *kernel_, static_cast(global.size()), - nullptr, global.data(), local.data(), + nullptr, global.data(), !local.empty() ? local.data() : nullptr, static_cast(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 &global, EventPointer event) { - CheckError(clEnqueueNDRangeKernel(queue(), *kernel_, static_cast(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/routines/common.cpp b/src/routines/common.cpp index 21e16954..3969cf9f 100644 --- a/src/routines/common.cpp +++ b/src/routines/common.cpp @@ -24,21 +24,23 @@ StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device, std::vector global, const std::vector &local, EventPointer event, const std::vector &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 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 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, const std::vector &local, - EventPointer event) { - auto emptyWaitingList = std::vector(); - return RunKernel(kernel, queue, device, global, local, event, emptyWaitingList); -} - // ================================================================================================= } // namespace clblast diff --git a/src/routines/common.hpp b/src/routines/common.hpp index d0bbc707..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 global, const std::vector &local, - EventPointer event, const std::vector &waitForEvents); - -// As above, but without an event waiting list -StatusCode RunKernel(Kernel &kernel, Queue &queue, const Device &device, - std::vector global, const std::vector &local, - EventPointer event); + EventPointer event, const std::vector &waitForEvents = {}); // ================================================================================================= -- cgit v1.2.3 From e4e1f05079273f60f4f15280b3f103810c7eb31f Mon Sep 17 00:00:00 2001 From: Ivan Shapovalov Date: Tue, 12 Jul 2016 13:33:25 +0300 Subject: clblast::Database, clblast::Routine: implement "database overlays" provided by routine implementation --- src/database/database.cpp | 45 ++++++++++++++++++++++++++------------------- src/database/database.hpp | 10 ++++++++-- src/routine.cpp | 5 +++-- src/routine.hpp | 3 ++- 4 files changed, 39 insertions(+), 24 deletions(-) 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::database = { // Constructor, computing device properties and populating the parameter-vector from the database Database::Database(const Queue &queue, const std::vector &kernels, - const Precision precision): + const Precision precision, const std::vector &overlay): parameters_{} { // Finds information of the current device @@ -53,10 +53,23 @@ Database::Database(const Queue &queue, const std::vector &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 &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; + 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 database; - // The constructor + // The constructor with a user-provided database overlay explicit Database(const Queue &queue, const std::vector &routines, - const Precision precision); + const Precision precision, const std::vector &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 &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 &routines, const Precision precision): + const std::vector &routines, const Precision precision, + const std::vector &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 &routines, const Precision precision); + const std::vector &routines, const Precision precision, + const std::vector &userDatabase = {}); // Set-up phase of the kernel StatusCode SetUp(); -- cgit v1.2.3 From 622682ffe30ae4e250d53a88eaacd899c905e20b Mon Sep 17 00:00:00 2001 From: Cedric Nugteren Date: Sun, 24 Jul 2016 16:41:01 +0200 Subject: Refactored the Python database script: separated functionality in modules, now complies to the PEP8 style, added proper command-line argument parsing, and cleaned-up --- scripts/database/database.py | 356 +++++++--------------------------- scripts/database/database/__init__.py | 0 scripts/database/database/bests.py | 20 ++ scripts/database/database/clblast.py | 132 +++++++++++++ scripts/database/database/db.py | 50 +++++ scripts/database/database/defaults.py | 58 ++++++ scripts/database/database/io.py | 58 ++++++ 7 files changed, 389 insertions(+), 285 deletions(-) mode change 100644 => 100755 scripts/database/database.py create mode 100644 scripts/database/database/__init__.py create mode 100644 scripts/database/database/bests.py create mode 100644 scripts/database/database/clblast.py create mode 100644 scripts/database/database/db.py create mode 100644 scripts/database/database/defaults.py create mode 100644 scripts/database/database/io.py diff --git a/scripts/database/database.py b/scripts/database/database.py old mode 100644 new mode 100755 index a70b9fc1..e115d68c --- a/scripts/database/database.py +++ b/scripts/database/database.py @@ -1,326 +1,112 @@ #!/usr/bin/env python -# ================================================================================================== -# This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This -# project loosely follows the Google C++ styleguide and uses a max-width of 100 characters per line. +# This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This file follows the +# PEP8 Python style guide and uses a max-width of 120 characters per line. # # Author(s): # Cedric Nugteren -# -# ================================================================================================== -# System modules import sys import os.path import glob -import re -import json -try: - from urllib.request import urlopen # Python 3 -except ImportError: - from urllib2 import urlopen # Python 2 +import argparse -# Additional modules import pandas as pd -print("## Using pandas version "+pd.__version__+", requires at least 0.17.0") + +import database.io as io +import database.db as db +import database.clblast as clblast +import database.bests as bests +import database.defaults as defaults # Server storing a copy of the database DATABASE_SERVER_URL = "http://www.cedricnugteren.nl/tuning/clblast.db" -# Constants -VENDOR_DEFAULT = "default" -DEVICETYPE_DEFAULT = "All" -DEVICENAME_DEFAULT = "default" - -# Attributes -DEVICETYPE_ATTRIBUTES = ["device_vendor", "device_type"] -DEVICE_ATTRIBUTES = ["device", "device_core_clock", "device_compute_units"] -KERNEL_ATTRIBUTES = ["precision", "kernel_family"] -ARGUMENT_ATTRIBUTES = ["arg_m", "arg_n", "arg_k", "arg_alpha", "arg_beta"] -ATTRIBUTES = DEVICE_ATTRIBUTES + DEVICETYPE_ATTRIBUTES + KERNEL_ATTRIBUTES + ARGUMENT_ATTRIBUTES - # OpenCL vendor names and their short name -VENDOR_NAMES = { "device_vendor": { +VENDOR_TRANSLATION_TABLE = {"device_vendor": { "GenuineIntel": "Intel", "Intel(R) Corporation": "Intel", "Advanced Micro Devices, Inc.": "AMD", "NVIDIA Corporation": "NVIDIA", }} -# Pandas options -pd.set_option('display.width', 1000) - -# ================================================================================================== -# Database operations -# ================================================================================================== - -# Downloads the database and save it to disk -def DownloadDatabase(filename): - print("## Downloading database from '"+DATABASE_SERVER_URL+"'...") - df = urlopen(DATABASE_SERVER_URL) - output = open(file_db,'wb') - output.write(df.read()) - output.close() - -# Loads the database from disk -def LoadDatabase(filename): - return pd.read_pickle(filename) - -# Saves the database to disk -def SaveDatabase(df, filename): - df.to_pickle(filename) - -# Loads JSON data from file -def ImportDataFromFile(filename): - with open(filename) as f: - data = json.load(f) - json_data = pd.DataFrame(data) - df = pd.io.json.json_normalize(json_data["results"]) - for attribute in ATTRIBUTES: - if attribute == "kernel_family": - df[attribute] = re.sub(r'_\d+', '', data[attribute]) - elif attribute in data: - df[attribute] = data[attribute] - else: - df[attribute] = 0 - return df - -# Returns the row-wise concatenation of two dataframes -def ConcatenateData(df1, df2): - return pd.concat([df1, df2]) - -# Removes duplicates from a dataframe -def RemoveDuplicates(df): - return df.drop_duplicates() - -# database = database[(database["device"] != "AMD Radeon R9 M370X Compute Engine") | (database["kernel_family"] != "xgemm") | (database["precision"] != "32")] -def RemoveEntriesByDevice(df, devicename): - return df[df["device"] != devicename] - -def RemoveEntriesByKernelFamily(df, familyname): - return df[df["kernel_family"] != familyname] - -def GetEntriesByField(df, field, value): - return df[df[field] == value] - -# Example usage: -# df = UpdateDatabase(df, (df["kernel_family"] == "xdot") & (df["arg_n"] == "67108864"), "arg_n", "2097152") -def UpdateDatabase(df, condition, field, value): - df.loc[condition, field] = value - return df - -# Fixes the problem that some vendors use multiple different names -def SanitizeVendorNames(df): - df = df.replace(VENDOR_NAMES) - return df - -# Retrieves the results with the lowest execution times -def GetBestResults(df): - dfbest = pd.DataFrame() - grouped = df.groupby(ATTRIBUTES+["kernel"]) - for name, dfgroup in grouped: - besttime = dfgroup["time"].min() - bestcase = dfgroup[dfgroup["time"] == besttime].iloc[0] - dfbest = dfbest.append(bestcase, ignore_index=True) - return dfbest - -# Sets defaults for devices of the same type/vendor based on the smallest values of all know -# entries. The average might be better for performance but some parameters might not be supported -# on other devices. -def CalculateDefaults(df): - dfdefault = pd.DataFrame() - - # Defaults per type/vendor - groups = df.groupby(DEVICETYPE_ATTRIBUTES+KERNEL_ATTRIBUTES+ARGUMENT_ATTRIBUTES+["kernel"]) - for name, dfgroup in groups: - default_values = dfgroup.min(axis=0) - default_values["device"] = DEVICENAME_DEFAULT - default_values["device_compute_units"] = 0 - default_values["device_core_clock"] = 0 - default_values["time"] = 0.0 - dfdefault = dfdefault.append(default_values, ignore_index=True) - - # Checks for mis-matched arguments - groups = dfdefault.groupby(DEVICETYPE_ATTRIBUTES+KERNEL_ATTRIBUTES+["kernel"]) - for name, dfgroup in groups: - if len(dfgroup) != 1: - description = dfgroup["kernel"].min() + " " + dfgroup["device_vendor"].min() - print("[WARNING] Entries for a single kernel with multiple argument values: " + description) - - # Defaults in general - groups = df.groupby(KERNEL_ATTRIBUTES+ARGUMENT_ATTRIBUTES+["kernel"]) - for name, dfgroup in groups: - default_values = dfgroup.min(axis=0) - default_values["device_vendor"] = VENDOR_DEFAULT - default_values["device_type"] = DEVICETYPE_DEFAULT - default_values["device"] = DEVICENAME_DEFAULT - default_values["device_compute_units"] = 0 - default_values["device_core_clock"] = 0 - default_values["time"] = 0.0 - dfdefault = dfdefault.append(default_values, ignore_index=True) - - # Database with both types of defaults only - return dfdefault - -# ================================================================================================== -# C++ header generation -# ================================================================================================== - -# The C++ header -def GetHeader(family): - return(""" -// ================================================================================================= -// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This -// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- -// width of 100 characters per line. -// -// Author(s): -// Database generator -// -// This file populates the database with best-found tuning parameters for the '%s' kernels. -// -// ================================================================================================= - -namespace clblast { -// =================================================================================================""" - % family.title()) - -# The C++ footer -def GetFooter(): - return("\n} // namespace clblast\n") - -# The start of a new C++ precision entry -def GetPrecision(family, precision): - precisionstring = "" - if precision == "16": - precisionstring = "Half" - elif precision == "32": - precisionstring = "Single" - elif precision == "64": - precisionstring = "Double" - elif precision == "3232": - precisionstring = "ComplexSingle" - elif precision == "6464": - precisionstring = "ComplexDouble" - else: - print("[ERROR] Unknown precision") - sys.exit() - return("\n\nconst Database::DatabaseEntry Database::%s%s = {\n \"%s\", Precision::k%s, {\n" - % (family.title(), precisionstring, family.title(), precisionstring)) - -# The C++ device type and vendor -def GetDeviceVendor(vendor, devtype): - if vendor == VENDOR_DEFAULT and devtype == DEVICETYPE_DEFAULT: - return(" { // Default\n kDeviceType%s, \"%s\", {\n" % (devtype, vendor)) - return(" { // %s %ss\n kDeviceType%s, \"%s\", {\n" % (vendor, devtype, devtype[0].upper() + devtype[1:], vendor)) - -# Prints the data to a C++ database -def PrintData(df, outputdir): - - # Iterates over the kernel families: creates a new file per family - for family, dffamily in df.groupby(["kernel_family"]): - dffamily = dffamily.dropna(axis=1, how='all') - f = open(os.path.join(outputdir, family+'.hpp'), 'w+') - f.write(GetHeader(family)) - - # Loops over the different entries for this family and prints their headers - for precision, dfprecision in dffamily.groupby(["precision"]): - f.write(GetPrecision(family, precision)) - for vendor, dfvendor in dfprecision.groupby(["device_vendor"]): - for devtype, dfdevtype in dfvendor.groupby(["device_type"]): - f.write(GetDeviceVendor(vendor, devtype)) - for device, dfdevice in dfdevtype.groupby(["device"]): - devicename = "\"%s\"," % device - f.write(" { %-50s { " % devicename) - # Collects the paramaters for this case and prints them - parameters = [] - for kernel, dfkernel in dfdevice.groupby(["kernel"]): - dfkernel = dfkernel.dropna(axis=1) - col_names = [col for col in list(dfkernel) if col.startswith('parameters.') and col != "parameters.PRECISION"] - parameters += ["{\"%s\",%d}" % (p.replace("parameters.",""), dfkernel[p].iloc[0]) for p in col_names] - f.write(", ".join(parameters)) - f.write(" } },\n") +def main(argv): - # Prints the footers - f.write(" }\n },\n") - f.write(" }\n};\n\n// =================================================================================================") - f.write(GetFooter()) + # Parses the command-line arguments + parser = argparse.ArgumentParser() + parser.add_argument("source_folder", help="The folder with JSON files to parse to add to the database") + parser.add_argument("clblast_root", help="Root of the CLBlast sources") + parser.add_argument("-v", "--verbose", action="store_true", help="Increase verbosity of the script") + cl_args = parser.parse_args(argv) -# ================================================================================================== -# Command-line arguments parsing and verification -# ================================================================================================== + # Parses the path arguments + database_filename = os.path.join(cl_args.clblast_root, "scripts", "database", "database.db") + json_files = os.path.join(cl_args.source_folder, "*.json") + cpp_database_path = os.path.join(cl_args.clblast_root, "src", "database", "kernels") -# Checks for the number of command-line arguments -if len(sys.argv) != 3: - print("[ERROR] Usage: database.py ") - sys.exit() + # Checks whether the command-line arguments are valid + clblast_header = os.path.join(cl_args.clblast_root, "include", "clblast.h") # Not used but just for validation + if not os.path.isfile(clblast_header): + raise RuntimeError("The path '" + cl_args.clblast_root + "' does not point to the root of the CLBlast library") + if len(glob.glob(json_files)) < 1: + print("[database] The path '" + cl_args.source_folder + "' does not contain any JSON files") -# Parses the command-line arguments -path_json = sys.argv[1] -path_clblast = sys.argv[2] -file_db = os.path.join(path_clblast, "scripts", "database", "database.db") -glob_json = os.path.join(path_json, "*.json") + # Pandas options + pd.set_option('display.width', 1000) + if cl_args.verbose: + print("[database] Using pandas version " + pd.__version__) -# Checks whether the command-line arguments are valid; exists otherwise -clblast_h = os.path.join(path_clblast, "include", "clblast.h") # Not used but just for validation -if not os.path.isfile(clblast_h): - print("[ERROR] The path '"+path_clblast+"' does not point to the root of the CLBlast library") - sys.exit() -if len(glob.glob(glob_json)) < 1: - print("## The path '"+path_json+"' does not contain any JSON files") + # Downloads the database if a local copy is not present + if not os.path.isfile(database_filename): + io.download_database(database_filename, DATABASE_SERVER_URL) -# ================================================================================================== -# The main body of the script -# ================================================================================================== + # Loads the database from disk + database = io.load_database(database_filename) -# Downloads the database if a local copy is not present -db_exists = os.path.isfile(file_db) -if not db_exists: - DownloadDatabase(file_db) + # Loops over all JSON files in the supplied folder + for file_json in glob.glob(json_files): -# Loads the database from disk -print("## Loading the database from disk...") -database = LoadDatabase(file_db) + # Loads the newly imported data + sys.stdout.write("[database] Processing '"+file_json+"' ") # No newline printed + imported_data = io.load_json_to_pandas(file_json) -# Loops over all JSON files in the supplied folder -for file_json in glob.glob(glob_json): + # Fixes the problem that some vendors use multiple different names + imported_data = db.find_and_replace(imported_data, VENDOR_TRANSLATION_TABLE) - # Loads the newly imported data - sys.stdout.write("## Processing '"+file_json+"' ") - imported_data = ImportDataFromFile(file_json) - imported_data = SanitizeVendorNames(imported_data) + # Adds the new data to the database + old_size = len(database.index) + database = db.concatenate_database(database, imported_data) + database = db.remove_duplicates(database) + new_size = len(database.index) + print("with " + str(new_size - old_size) + " new items") # Newline printed here - # Adds the new data to the database - old_size = len(database.index) - database = ConcatenateData(database, imported_data) - database = RemoveDuplicates(database) - new_size = len(database.index) - print("with "+str(new_size-old_size)+" new items") + # Stores the modified database back to disk + if len(glob.glob(json_files)) >= 1: + io.save_database(database, database_filename) -# Stores the modified database back to disk -if len(glob.glob(glob_json)) >= 1: - print("## Storing the database to disk...") - SaveDatabase(database, file_db) + # Optional: update the database here. Default is disabled, code below is just an example + if False: # TODO: Use command-line arguments to enable updates in a flexible way + database = db.update_database(database, + ((database["kernel"] == "CopyMatrixFast") & + (database["precision"] == "3232")), + "arg_alpha", "2+0.5i") + io.save_database(database, database_filename) -# Optional: update the database here. Default is disabled, code below is just an example -if False: - database = UpdateDatabase(database, ((database["kernel"] == "CopyMatrixFast") & (database["precision"] == "3232")), "arg_alpha", "2+0.5i") - SaveDatabase(database, file_db) + # Retrieves the best performing results + print("[database] Calculating the best results per device/kernel...") + database_best_results = bests.get_best_results(database) -# Retrieves the best performing results -print("## Calculating the best results per device/kernel...") -bests = GetBestResults(database) + # Determines the defaults for other vendors and per vendor + database_defaults = defaults.calculate_defaults(database_best_results) + database_best_results = db.concatenate_database(database_best_results, database_defaults) -# Determines the defaults for other vendors and per vendor -defaults = CalculateDefaults(bests) -bests = ConcatenateData(bests, defaults) + # Outputs the database as a C++ database + print("[database] Producing a C++ database in '" + cpp_database_path + "'...") + clblast.print_cpp_database(database_best_results, cpp_database_path) -# Outputs the data as a C++ database -path_cpp_database = os.path.join(path_clblast, "src", "database", "kernels") -print("## Producing a C++ database in '"+path_cpp_database+"'...") -PrintData(bests, path_cpp_database) + print("[database] All done") -print("## All done") -# ================================================================================================== +if __name__ == '__main__': + main(sys.argv[1:]) diff --git a/scripts/database/database/__init__.py b/scripts/database/database/__init__.py new file mode 100644 index 00000000..e69de29b diff --git a/scripts/database/database/bests.py b/scripts/database/database/bests.py new file mode 100644 index 00000000..edb81733 --- /dev/null +++ b/scripts/database/database/bests.py @@ -0,0 +1,20 @@ + +# This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This file follows the +# PEP8 Python style guide and uses a max-width of 120 characters per line. +# +# Author(s): +# Cedric Nugteren + +import pandas as pd +import clblast + + +def get_best_results(df): + """Retrieves the results with the lowests execution times""" + database_bests = pd.DataFrame() + database_entries = df.groupby(clblast.ATTRIBUTES + ["kernel"]) + for name, database_entry in database_entries: + best_time = database_entry["time"].min() + best_parameters = database_entry[database_entry["time"] == best_time].iloc[0] + database_bests = database_bests.append(best_parameters, ignore_index=True) + return database_bests diff --git a/scripts/database/database/clblast.py b/scripts/database/database/clblast.py new file mode 100644 index 00000000..9c9f7eb4 --- /dev/null +++ b/scripts/database/database/clblast.py @@ -0,0 +1,132 @@ + +# This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This file follows the +# PEP8 Python style guide and uses a max-width of 120 characters per line. +# +# Author(s): +# Cedric Nugteren + +import os + +# Constants from the C++ code +VENDOR_DEFAULT = "default" +DEVICE_TYPE_DEFAULT = "All" +DEVICE_NAME_DEFAULT = "default" + +# List of attributes +DEVICE_TYPE_ATTRIBUTES = ["device_vendor", "device_type"] +DEVICE_ATTRIBUTES = ["device", "device_core_clock", "device_compute_units"] +KERNEL_ATTRIBUTES = ["precision", "kernel_family"] +ARGUMENT_ATTRIBUTES = ["arg_m", "arg_n", "arg_k", "arg_alpha", "arg_beta"] +ATTRIBUTES = DEVICE_ATTRIBUTES + DEVICE_TYPE_ATTRIBUTES + KERNEL_ATTRIBUTES + ARGUMENT_ATTRIBUTES + + +def precision_to_string(precision): + """Translates a precision number (represented as Python string) into a descriptive string""" + if precision == "16": + return "Half" + elif precision == "32": + return "Single" + elif precision == "64": + return "Double" + elif precision == "3232": + return "ComplexSingle" + elif precision == "6464": + return "ComplexDouble" + else: + raise("Unknown precision: " + precision) + + +def get_cpp_separator(): + """Retrieves a C++ comment separator""" + return "// =================================================================================================" + + +def get_cpp_header(family): + """Retrieves the C++ header""" + return ("\n" + get_cpp_separator() + """ +// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This +// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max- +// width of 100 characters per line. +// +// Author(s): +// Database generator +// +// This file populates the database with best-found tuning parameters for the '%s' kernels. +//\n""" + % family.title() + get_cpp_separator() + "\n\nnamespace clblast {\n" + get_cpp_separator()) + + +def get_cpp_footer(): + """Retrieves the C++ footer""" + return "\n} // namespace clblast\n" + + +def get_cpp_precision(family, precision): + """Retrieves the C++ code for the start of a new precision""" + precision_string = precision_to_string(precision) + return("\n\nconst Database::DatabaseEntry Database::%s%s = {\n \"%s\", Precision::k%s, {\n" + % (family.title(), precision_string, family.title(), precision_string)) + + +def get_cpp_device_vendor(vendor, device_type): + """Retrieves the C++ code for the (default) vendor and device type""" + if vendor == VENDOR_DEFAULT and device_type == DEVICE_TYPE_DEFAULT: + return " { // Default\n kDeviceType%s, \"%s\", {\n" % (device_type, vendor) + device_type_caps = device_type[0].upper() + device_type[1:] + return " { // %s %ss\n kDeviceType%s, \"%s\", {\n" % (vendor, device_type, device_type_caps, vendor) + + +def print_cpp_database(database, output_dir): + """Outputs the database as C++ code""" + + # Iterates over the kernel families + for family_name, family_database in database.groupby(["kernel_family"]): + family_database = family_database.dropna(axis=1, how='all') + + # Opens a new file for each kernel family + full_path = os.path.join(output_dir, family_name+'.hpp') + with open(full_path, 'w+') as f: + f.write(get_cpp_header(family_name)) + + # Loops over the different precision (e.g. 16, 32, 3232, 64, 6464) + for precision, precision_database in family_database.groupby(["precision"]): + f.write(get_cpp_precision(family_name, precision)) + + # Loops over a combination of device vendors and device types (e.g. AMD GPU) + for vendor, vendor_database in precision_database.groupby(["device_vendor"]): + for device_type, device_type_database in vendor_database.groupby(["device_type"]): + f.write(get_cpp_device_vendor(vendor, device_type)) + + # Loops over every device of this vendor-type combination + for device_name, device_database in device_type_database.groupby(["device"]): + device_name_quoted = "\"%s\"," % device_name + device_name_cpp = " { %-50s { " % device_name_quoted + f.write(device_name_cpp) + + # Collects the parameters for this entry + parameters = [] + for kernel, kernel_database in device_database.groupby(["kernel"]): + kernel_database = kernel_database.dropna(axis=1) + + # Only consider the actual parameters, not the precision + def is_parameter(column): + return column.startswith('parameters.') and column != "parameters.PRECISION" + column_names = [col for col in list(kernel_database) if is_parameter(col)] + + for p in column_names: + parameter_name = p.replace("parameters.", "") + parameter_value = int(kernel_database[p].iloc[0]) + parameters.append("{\"" + parameter_name + "\"," + str(parameter_value) + "}") + + # Prints the entry + f.write(", ".join(parameters)) + f.write(" } },\n") + + # Prints the vendor-type combination footer + f.write(" }\n },\n") + + # Prints the precision footer + f.write(" }\n};\n\n" + get_cpp_separator()) + + # Prints the file footer + f.write(get_cpp_footer()) diff --git a/scripts/database/database/db.py b/scripts/database/database/db.py new file mode 100644 index 00000000..60cfbcfa --- /dev/null +++ b/scripts/database/database/db.py @@ -0,0 +1,50 @@ + +# This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This file follows the +# PEP8 Python style guide and uses a max-width of 120 characters per line. +# +# Author(s): +# Cedric Nugteren + +import pandas as pd + + +def get_entries_by_field(database, field, value): + """Retrieves entries from the database with a specific value for a given field""" + return database[database[field] == value] + + +def concatenate_database(database1, database2): + """Concatenates two databases row-wise and returns the result""" + return pd.concat([database1, database2]) + + +def remove_duplicates(database): + """Removes duplicates from a database""" + return database.drop_duplicates() + + +def find_and_replace(database, dictionary): + """Finds and replaces entries in a database based on a dictionary. Example: + dictionary = { "key_to_edit": { find1: replace1, find2, replace2 } }""" + return database.replace(dictionary) + + +def remove_entries_by_key_value(database, key, value): + """Removes entries in the databased which have a specific value for a given key""" + return database[database[key] != value] + + +def remove_entries_by_device(database, device_name): + """Shorthand for the above, specifically removes entries for a given device""" + return remove_entries_by_key_value(database, "device", device_name) + + +def remove_entries_by_kernel_family(database, kernel_family_name): + """Shorthand for the above, specifically removes entries for a given kernel family""" + return remove_entries_by_key_value(database, "kernel_family", kernel_family_name) + + +def update_database(database, condition, field, value): + """Updates the database by writing a specific value to a given field, given certain conditions""" + database.loc[condition, field] = value + return database diff --git a/scripts/database/database/defaults.py b/scripts/database/database/defaults.py new file mode 100644 index 00000000..357c3a3a --- /dev/null +++ b/scripts/database/database/defaults.py @@ -0,0 +1,58 @@ + +# This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This file follows the +# PEP8 Python style guide and uses a max-width of 120 characters per line. +# +# Author(s): +# Cedric Nugteren + +import pandas as pd +import clblast + + +def set_default_device(database_entry): + """Sets the device name and parameters to some default values""" + database_entry["device"] = clblast.DEVICE_NAME_DEFAULT + database_entry["device_compute_units"] = 0 + database_entry["device_core_clock"] = 0 + return database_entry + + +def set_default_time(database_entry): + """Sets the execution time to some default value""" + database_entry["time"] = 0.0 + return database_entry + + +def calculate_defaults(df): + """# Sets defaults for devices of the same type/vendor based on the smallest values of all known entries. The average + might be better for performance but some parameters might not be supported on other devices.""" + database_defaults = pd.DataFrame() + + # Defaults per combination of device vendors and device types (e.g. AMD GPU) + database_type_vendor = df.groupby(clblast.DEVICE_TYPE_ATTRIBUTES + clblast.KERNEL_ATTRIBUTES + ["kernel"] + + clblast.ARGUMENT_ATTRIBUTES) + for group_name, database_group in database_type_vendor: + default_values = database_group.min(axis=0) + default_values = set_default_device(default_values) + default_values = set_default_time(default_values) + database_defaults = database_defaults.append(default_values, ignore_index=True) + + # Checks for mis-matched arguments + groups = database_defaults.groupby(clblast.DEVICE_TYPE_ATTRIBUTES + clblast.KERNEL_ATTRIBUTES + ["kernel"]) + for group_name, database_group in groups: + if len(database_group) != 1: + description = database_group["kernel"].min() + " " + database_group["device_vendor"].min() + print("[WARNING] Entries for a single kernel with multiple argument values: " + description) + + # Defaults over all device types and vendors + groups = df.groupby(clblast.KERNEL_ATTRIBUTES + ["kernel"] + clblast.ARGUMENT_ATTRIBUTES) + for group_name, database_group in groups: + default_values = database_group.min(axis=0) + default_values["device_vendor"] = clblast.VENDOR_DEFAULT + default_values["device_type"] = clblast.DEVICE_TYPE_DEFAULT + default_values = set_default_device(default_values) + default_values = set_default_time(default_values) + database_defaults = database_defaults.append(default_values, ignore_index=True) + + # Database with both types of defaults only + return database_defaults diff --git a/scripts/database/database/io.py b/scripts/database/database/io.py new file mode 100644 index 00000000..ad2f7ae9 --- /dev/null +++ b/scripts/database/database/io.py @@ -0,0 +1,58 @@ + +# This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This file follows the +# PEP8 Python style guide and uses a max-width of 120 characters per line. +# +# Author(s): +# Cedric Nugteren + +import re +import json + +try: + from urllib.request import urlopen # Python 3 +except ImportError: + from urllib2 import urlopen # Python 2 + +import pandas as pd + +import clblast + + +def download_database(filename, database_url): + """Downloads a database and saves it to disk""" + print("[database] Downloading database from '" + database_url + "'...") + database = urlopen(database_url) + with open(filename, 'wb') as f: + f.write(database.read()) + + +def load_database(filename): + """Loads a database from disk""" + print("[database] Loading database from '" + filename + "'") + return pd.read_pickle(filename) + + +def save_database(database, filename): + """Saves a database to disk""" + print("[database] Saving database to '" + filename + "'") + database.to_pickle(filename) + + +def load_json_to_pandas(filename): + """Loads JSON data from file and converts it to a pandas database""" + with open(filename) as f: + json_data = json.load(f) + + # Gathers all results and stores them in a new database + json_database = pd.DataFrame(json_data) + new_database = pd.io.json.json_normalize(json_database["results"]) + + # Sets the common attributes to each entry in the results + for attribute in clblast.ATTRIBUTES: + if attribute == "kernel_family": + new_database[attribute] = re.sub(r'_\d+', '', json_data[attribute]) + elif attribute in json_data: + new_database[attribute] = json_data[attribute] + else: + new_database[attribute] = 0 # For example a parameters that was not used by this kernel + return new_database -- cgit v1.2.3 From ffa35c623af4b0916f625f3a41000e75a1df7e1f Mon Sep 17 00:00:00 2001 From: Cedric Nugteren Date: Sun, 24 Jul 2016 17:00:21 +0200 Subject: Minor improvements after merging in groundwork for custom tuning parameters and kernels --- src/database/database.cpp | 12 ++++++++---- src/database/database.hpp | 8 ++------ src/routine.hpp | 3 ++- 3 files changed, 12 insertions(+), 11 deletions(-) diff --git a/src/database/database.cpp b/src/database/database.cpp index ea1557b9..47f1da16 100644 --- a/src/database/database.cpp +++ b/src/database/database.cpp @@ -42,7 +42,8 @@ const std::vector Database::database = { // ================================================================================================= -// Constructor, computing device properties and populating the parameter-vector from the database +// Constructor, computing device properties and populating the parameter-vector from the database. +// This takes an optional overlay database in case of custom tuning or custom kernels. Database::Database(const Queue &queue, const std::vector &kernels, const Precision precision, const std::vector &overlay): parameters_{} { @@ -66,7 +67,10 @@ Database::Database(const Queue &queue, const std::vector &kernels, 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) { + parameters_.insert(search_result->begin(), search_result->end()); + break; + } } if (!search_result) { throw std::runtime_error("Database error, could not find a suitable entry"); } @@ -86,7 +90,7 @@ std::string Database::GetDefines() const { // ================================================================================================= -// Searches the database for the right kernel and precision +// Searches a particular database for the right kernel and precision Database::ParametersPtr Database::Search(const std::string &this_kernel, const std::string &this_type, const std::string &this_vendor, @@ -119,7 +123,7 @@ Database::ParametersPtr Database::Search(const std::string &this_kernel, } } - // If we reached this point, something is wrong + // If we reached this point, the entry was not found in this database return nullptr; } diff --git a/src/database/database.hpp b/src/database/database.hpp index 5a61fad9..e84357dc 100644 --- a/src/database/database.hpp +++ b/src/database/database.hpp @@ -79,7 +79,7 @@ class Database { static const DatabaseEntry PadtransposeHalf, PadtransposeSingle, PadtransposeDouble, PadtransposeComplexSingle, PadtransposeComplexDouble; static const std::vector database; - // The constructor with a user-provided database overlay + // The constructor with a user-provided database overlay (potentially an empty vector) explicit Database(const Queue &queue, const std::vector &routines, const Precision precision, const std::vector &overlay); @@ -90,11 +90,7 @@ class Database { std::string GetDefines() const; private: - Parameters 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; - - // Alternate search method in a specified database, returning pointer (possibly NULL) + // Search method for a specified database, returning pointer (possibly a nullptr) 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 &db) const; diff --git a/src/routine.hpp b/src/routine.hpp index 21506e7b..f5c607af 100644 --- a/src/routine.hpp +++ b/src/routine.hpp @@ -32,7 +32,8 @@ namespace clblast { class Routine { public: - // Base class constructor + // Base class constructor. The user database is an optional extra database to override the + // built-in database. explicit Routine(Queue &queue, EventPointer event, const std::string &name, const std::vector &routines, const Precision precision, const std::vector &userDatabase = {}); -- cgit v1.2.3