diff --git a/CHANGELOG b/CHANGELOG index b6e09102..d018e211 100644 --- a/CHANGELOG +++ b/CHANGELOG @@ -7,6 +7,7 @@ Development version (next release) - Fixed a bug related to the cache and retrieval of programs based on the OpenCL context - Fixed a performance issue (caused by fp16 support) by optimizing alpha/beta parameter passing to kernels - Added an option (-warm_up) to do a warm-up run before timing in the performance clients +- Improved performance significantly of rotated GEMV computations - Added tuned parameters for various devices (see README) Version 0.8.0 diff --git a/scripts/database/database/clblast.py b/scripts/database/database/clblast.py index 9c9f7eb4..46b711cc 100644 --- a/scripts/database/database/clblast.py +++ b/scripts/database/database/clblast.py @@ -64,8 +64,9 @@ def get_cpp_footer(): def get_cpp_precision(family, precision): """Retrieves the C++ code for the start of a new precision""" precision_string = precision_to_string(precision) + camelcase_name = family.title().replace("_", "") return("\n\nconst Database::DatabaseEntry Database::%s%s = {\n \"%s\", Precision::k%s, {\n" - % (family.title(), precision_string, family.title(), precision_string)) + % (camelcase_name, precision_string, camelcase_name, precision_string)) def get_cpp_device_vendor(vendor, device_type): diff --git a/src/database/database.cpp b/src/database/database.cpp index 47f1da16..38974b95 100644 --- a/src/database/database.cpp +++ b/src/database/database.cpp @@ -17,6 +17,8 @@ #include "database/kernels/xaxpy.hpp" #include "database/kernels/xdot.hpp" #include "database/kernels/xgemv.hpp" +#include "database/kernels/xgemv_fast.hpp" +#include "database/kernels/xgemv_fast_rot.hpp" #include "database/kernels/xger.hpp" #include "database/kernels/xgemm.hpp" #include "database/kernels/copy.hpp" @@ -32,8 +34,10 @@ const std::vector Database::database = { XaxpyHalf, XaxpySingle, XaxpyDouble, XaxpyComplexSingle, XaxpyComplexDouble, XdotHalf, XdotSingle, XdotDouble, XdotComplexSingle, XdotComplexDouble, XgemvHalf, XgemvSingle, XgemvDouble, XgemvComplexSingle, XgemvComplexDouble, + XgemvFastHalf, XgemvFastSingle, XgemvFastDouble, XgemvFastComplexSingle, XgemvFastComplexDouble, + /* XgemvFastRotHalf, */ XgemvFastRotSingle, XgemvFastRotDouble, XgemvFastRotComplexSingle, XgemvFastRotComplexDouble, XgerHalf, XgerSingle, XgerDouble, XgerComplexSingle, XgerComplexDouble, - XgemmHalf, XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble, + /* XgemmHalf, */ XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble, CopyHalf, CopySingle, CopyDouble, CopyComplexSingle, CopyComplexDouble, PadHalf, PadSingle, PadDouble, PadComplexSingle, PadComplexDouble, TransposeHalf, TransposeSingle, TransposeDouble, TransposeComplexSingle, TransposeComplexDouble, diff --git a/src/database/database.hpp b/src/database/database.hpp index e84357dc..8d6d3863 100644 --- a/src/database/database.hpp +++ b/src/database/database.hpp @@ -71,8 +71,10 @@ class Database { static const DatabaseEntry XaxpyHalf, XaxpySingle, XaxpyDouble, XaxpyComplexSingle, XaxpyComplexDouble; static const DatabaseEntry XdotHalf, XdotSingle, XdotDouble, XdotComplexSingle, XdotComplexDouble; static const DatabaseEntry XgemvHalf, XgemvSingle, XgemvDouble, XgemvComplexSingle, XgemvComplexDouble; + static const DatabaseEntry XgemvFastHalf, XgemvFastSingle, XgemvFastDouble, XgemvFastComplexSingle, XgemvFastComplexDouble; + static const DatabaseEntry /* XgemvFastRotHalf, */ XgemvFastRotSingle, XgemvFastRotDouble, XgemvFastRotComplexSingle, XgemvFastRotComplexDouble; static const DatabaseEntry XgerHalf, XgerSingle, XgerDouble, XgerComplexSingle, XgerComplexDouble; - static const DatabaseEntry XgemmHalf, XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble; + static const DatabaseEntry /* XgemmHalf, */ XgemmSingle, XgemmDouble, XgemmComplexSingle, XgemmComplexDouble; static const DatabaseEntry CopyHalf, CopySingle, CopyDouble, CopyComplexSingle, CopyComplexDouble; static const DatabaseEntry PadHalf, PadSingle, PadDouble, PadComplexSingle, PadComplexDouble; static const DatabaseEntry TransposeHalf, TransposeSingle, TransposeDouble, TransposeComplexSingle, TransposeComplexDouble; diff --git a/src/database/kernels/xgemm.hpp b/src/database/kernels/xgemm.hpp index 736f2695..61b7ff05 100644 --- a/src/database/kernels/xgemm.hpp +++ b/src/database/kernels/xgemm.hpp @@ -14,18 +14,6 @@ namespace clblast { // ================================================================================================= -const Database::DatabaseEntry Database::XgemmHalf = { - "Xgemm", Precision::kHalf, { - { // Default - kDeviceTypeAll, "default", { - { "default", { {"KWG",16}, {"KWI",2}, {"MDIMA",8}, {"MDIMC",8}, {"MWG",32}, {"NDIMB",8}, {"NDIMC",8}, {"NWG",64}, {"SA",0}, {"SB",0}, {"STRM",0}, {"STRN",0}, {"VWM",1}, {"VWN",1} } }, - } - }, - } -}; - -// ================================================================================================= - const Database::DatabaseEntry Database::XgemmSingle = { "Xgemm", Precision::kSingle, { { // AMD GPUs diff --git a/src/database/kernels/xgemv.hpp b/src/database/kernels/xgemv.hpp index 65f4b5c8..6d680b06 100644 --- a/src/database/kernels/xgemv.hpp +++ b/src/database/kernels/xgemv.hpp @@ -18,13 +18,13 @@ const Database::DatabaseEntry Database::XgemvHalf = { "Xgemv", Precision::kHalf, { { // Intel GPUs kDeviceTypeGPU, "Intel", { - { "Intel(R) HD Graphics Skylake ULT GT2", { {"WGS1",128}, {"WPT1",1}, {"VW2",2}, {"WGS2",128}, {"WPT2",2}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "default", { {"WGS1",128}, {"WPT1",1}, {"VW2",2}, {"WGS2",128}, {"WPT2",2}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "Intel(R) HD Graphics Skylake ULT GT2", { {"WGS1",128}, {"WPT1",1} } }, + { "default", { {"WGS1",128}, {"WPT1",1} } }, } }, { // Default kDeviceTypeAll, "default", { - { "default", { {"WGS1",128}, {"WPT1",1}, {"VW2",2}, {"WGS2",128}, {"WPT2",2}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "default", { {"WGS1",128}, {"WPT1",1} } }, } }, } @@ -36,57 +36,57 @@ const Database::DatabaseEntry Database::XgemvSingle = { "Xgemv", Precision::kSingle, { { // AMD GPUs kDeviceTypeGPU, "AMD", { - { "AMD Radeon R9 M370X Compute Engine", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Hawaii", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Oland", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",4}, {"WGS3",256}, {"WPT3",4} } }, - { "Pitcairn", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Tahiti", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "default", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "AMD Radeon R9 M370X Compute Engine", { {"WGS1",128}, {"WPT1",1} } }, + { "Hawaii", { {"WGS1",128}, {"WPT1",1} } }, + { "Oland", { {"WGS1",128}, {"WPT1",1} } }, + { "Pitcairn", { {"WGS1",256}, {"WPT1",1} } }, + { "Tahiti", { {"WGS1",256}, {"WPT1",1} } }, + { "default", { {"WGS1",128}, {"WPT1",1} } }, } }, { // Intel CPUs kDeviceTypeCPU, "Intel", { - { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"WGS1",64}, {"WPT1",1}, {"VW2",4}, {"WGS2",128}, {"WPT2",4}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"WGS1",64}, {"WPT1",4}, {"VW2",1}, {"WGS2",64}, {"WPT2",4}, {"VW3",2}, {"WGS3",64}, {"WPT3",4} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",4}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"WGS1",64}, {"WPT1",1} } }, + { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"WGS1",64}, {"WPT1",4} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Intel GPUs kDeviceTypeGPU, "Intel", { - { "Intel(R) HD Graphics 530", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Intel(R) HD Graphics Haswell Ultrabook GT2 Mobile", { {"WGS1",256}, {"WPT1",1}, {"VW2",4}, {"WGS2",128}, {"WPT2",4}, {"VW3",4}, {"WGS3",256}, {"WPT3",4} } }, - { "Intel(R) HD Graphics Skylake ULT GT2", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",4}, {"WGS3",64}, {"WPT3",4} } }, - { "Iris", { {"WGS1",64}, {"WPT1",2}, {"VW2",1}, {"WGS2",128}, {"WPT2",2}, {"VW3",4}, {"WGS3",64}, {"WPT3",8} } }, - { "Iris Pro", { {"WGS1",256}, {"WPT1",2}, {"VW2",1}, {"WGS2",128}, {"WPT2",2}, {"VW3",4}, {"WGS3",64}, {"WPT3",4} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "Intel(R) HD Graphics 530", { {"WGS1",256}, {"WPT1",1} } }, + { "Intel(R) HD Graphics Haswell Ultrabook GT2 Mobile", { {"WGS1",256}, {"WPT1",1} } }, + { "Intel(R) HD Graphics Skylake ULT GT2", { {"WGS1",64}, {"WPT1",1} } }, + { "Iris", { {"WGS1",64}, {"WPT1",2} } }, + { "Iris Pro", { {"WGS1",256}, {"WPT1",2} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Intel accelerators kDeviceTypeAccelerator, "Intel", { - { "Intel(R) Many Integrated Core Acceleration Card", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "Intel(R) Many Integrated Core Acceleration Card", { {"WGS1",64}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // NVIDIA GPUs kDeviceTypeGPU, "NVIDIA", { - { "GRID K520", { {"WGS1",256}, {"WPT1",1}, {"VW2",2}, {"WGS2",256}, {"WPT2",2}, {"VW3",2}, {"WGS3",64}, {"WPT3",2} } }, - { "GeForce GTX 1070", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",2}, {"WGS3",64}, {"WPT3",2} } }, - { "GeForce GTX 480", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",2}, {"WGS3",64}, {"WPT3",2} } }, - { "GeForce GTX 670", { {"WGS1",64}, {"WPT1",1}, {"VW2",2}, {"WGS2",256}, {"WPT2",2}, {"VW3",2}, {"WGS3",64}, {"WPT3",2} } }, - { "GeForce GTX 680", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",2}, {"WGS3",128}, {"WPT3",2} } }, - { "GeForce GTX 750", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",2}, {"WGS3",128}, {"WPT3",2} } }, - { "GeForce GTX 750 Ti", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",4}, {"WGS3",128}, {"WPT3",4} } }, - { "GeForce GTX 980", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",128}, {"WPT3",1} } }, - { "GeForce GTX TITAN", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",256}, {"WPT3",1} } }, - { "GeForce GTX TITAN X", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",128}, {"WPT3",1} } }, - { "Tesla K20m", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",256}, {"WPT3",1} } }, - { "Tesla K40m", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",128}, {"WPT3",1} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "GRID K520", { {"WGS1",256}, {"WPT1",1} } }, + { "GeForce GTX 1070", { {"WGS1",128}, {"WPT1",1} } }, + { "GeForce GTX 480", { {"WGS1",64}, {"WPT1",1} } }, + { "GeForce GTX 670", { {"WGS1",64}, {"WPT1",1} } }, + { "GeForce GTX 680", { {"WGS1",256}, {"WPT1",1} } }, + { "GeForce GTX 750", { {"WGS1",256}, {"WPT1",1} } }, + { "GeForce GTX 750 Ti", { {"WGS1",256}, {"WPT1",1} } }, + { "GeForce GTX 980", { {"WGS1",128}, {"WPT1",1} } }, + { "GeForce GTX TITAN", { {"WGS1",256}, {"WPT1",1} } }, + { "GeForce GTX TITAN X", { {"WGS1",256}, {"WPT1",1} } }, + { "Tesla K20m", { {"WGS1",128}, {"WPT1",1} } }, + { "Tesla K40m", { {"WGS1",256}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Default kDeviceTypeAll, "default", { - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, } @@ -98,53 +98,53 @@ const Database::DatabaseEntry Database::XgemvComplexSingle = { "Xgemv", Precision::kComplexSingle, { { // AMD GPUs kDeviceTypeGPU, "AMD", { - { "AMD Radeon R9 M370X Compute Engine", { {"WGS1",64}, {"WPT1",1}, {"VW2",2}, {"WGS2",256}, {"WPT2",2}, {"VW3",1}, {"WGS3",128}, {"WPT3",1} } }, - { "Hawaii", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Oland", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",4}, {"WGS3",64}, {"WPT3",4} } }, - { "Pitcairn", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",2}, {"WGS3",64}, {"WPT3",2} } }, - { "Tahiti", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "AMD Radeon R9 M370X Compute Engine", { {"WGS1",64}, {"WPT1",1} } }, + { "Hawaii", { {"WGS1",64}, {"WPT1",1} } }, + { "Oland", { {"WGS1",64}, {"WPT1",1} } }, + { "Pitcairn", { {"WGS1",64}, {"WPT1",1} } }, + { "Tahiti", { {"WGS1",64}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Intel CPUs kDeviceTypeCPU, "Intel", { - { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",2}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"WGS1",64}, {"WPT1",4}, {"VW2",4}, {"WGS2",64}, {"WPT2",4}, {"VW3",2}, {"WGS3",64}, {"WPT3",2} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",2}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"WGS1",128}, {"WPT1",1} } }, + { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"WGS1",64}, {"WPT1",4} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Intel GPUs kDeviceTypeGPU, "Intel", { - { "Intel(R) HD Graphics 530", { {"WGS1",64}, {"WPT1",1}, {"VW2",2}, {"WGS2",128}, {"WPT2",2}, {"VW3",4}, {"WGS3",64}, {"WPT3",4} } }, - { "Intel(R) HD Graphics Haswell Ultrabook GT2 Mobile", { {"WGS1",128}, {"WPT1",1}, {"VW2",2}, {"WGS2",128}, {"WPT2",2}, {"VW3",4}, {"WGS3",128}, {"WPT3",4} } }, - { "Intel(R) HD Graphics Skylake ULT GT2", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",4}, {"WGS3",64}, {"WPT3",4} } }, - { "Iris", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Iris Pro", { {"WGS1",64}, {"WPT1",1}, {"VW2",4}, {"WGS2",128}, {"WPT2",4}, {"VW3",2}, {"WGS3",64}, {"WPT3",2} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "Intel(R) HD Graphics 530", { {"WGS1",64}, {"WPT1",1} } }, + { "Intel(R) HD Graphics Haswell Ultrabook GT2 Mobile", { {"WGS1",128}, {"WPT1",1} } }, + { "Intel(R) HD Graphics Skylake ULT GT2", { {"WGS1",64}, {"WPT1",1} } }, + { "Iris", { {"WGS1",256}, {"WPT1",1} } }, + { "Iris Pro", { {"WGS1",64}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Intel accelerators kDeviceTypeAccelerator, "Intel", { - { "Intel(R) Many Integrated Core Acceleration Card", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "Intel(R) Many Integrated Core Acceleration Card", { {"WGS1",64}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // NVIDIA GPUs kDeviceTypeGPU, "NVIDIA", { - { "GRID K520", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "GeForce GTX 1070", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, - { "GeForce GTX 480", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "GeForce GTX 670", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "GeForce GTX 680", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "GRID K520", { {"WGS1",256}, {"WPT1",1} } }, + { "GeForce GTX 1070", { {"WGS1",64}, {"WPT1",1} } }, + { "GeForce GTX 480", { {"WGS1",64}, {"WPT1",1} } }, + { "GeForce GTX 670", { {"WGS1",64}, {"WPT1",1} } }, + { "GeForce GTX 680", { {"WGS1",64}, {"WPT1",1} } }, { "GeForce GTX 750", { {"WGS1",128}, {"WPT1",1} } }, - { "GeForce GTX 750 Ti", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "GeForce GTX 750 Ti", { {"WGS1",64}, {"WPT1",1} } }, { "GeForce GTX TITAN", { {"WGS1",256}, {"WPT1",1} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Default kDeviceTypeAll, "default", { - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, } @@ -156,47 +156,47 @@ const Database::DatabaseEntry Database::XgemvDouble = { "Xgemv", Precision::kDouble, { { // AMD GPUs kDeviceTypeGPU, "AMD", { - { "AMD Radeon R9 M370X Compute Engine", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",128}, {"WPT3",1} } }, - { "Hawaii", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Oland", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",4}, {"WGS3",256}, {"WPT3",4} } }, - { "Pitcairn", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",2}, {"WGS3",64}, {"WPT3",2} } }, - { "Tahiti", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",2}, {"WGS3",64}, {"WPT3",2} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "AMD Radeon R9 M370X Compute Engine", { {"WGS1",64}, {"WPT1",1} } }, + { "Hawaii", { {"WGS1",128}, {"WPT1",1} } }, + { "Oland", { {"WGS1",256}, {"WPT1",1} } }, + { "Pitcairn", { {"WGS1",256}, {"WPT1",1} } }, + { "Tahiti", { {"WGS1",256}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Intel CPUs kDeviceTypeCPU, "Intel", { - { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"WGS1",64}, {"WPT1",2}, {"VW2",4}, {"WGS2",128}, {"WPT2",4}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"WGS1",64}, {"WPT1",4}, {"VW2",1}, {"WGS2",64}, {"WPT2",4}, {"VW3",1}, {"WGS3",64}, {"WPT3",2} } }, - { "default", { {"WGS1",64}, {"WPT1",2}, {"VW2",1}, {"WGS2",64}, {"WPT2",4}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"WGS1",64}, {"WPT1",2} } }, + { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"WGS1",64}, {"WPT1",4} } }, + { "default", { {"WGS1",64}, {"WPT1",2} } }, } }, { // Intel accelerators kDeviceTypeAccelerator, "Intel", { - { "Intel(R) Many Integrated Core Acceleration Card", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "Intel(R) Many Integrated Core Acceleration Card", { {"WGS1",64}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // NVIDIA GPUs kDeviceTypeGPU, "NVIDIA", { - { "GRID K520", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "GeForce GTX 1070", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",128}, {"WPT3",1} } }, - { "GeForce GTX 480", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "GeForce GTX 670", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "GeForce GTX 680", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",2}, {"WGS3",128}, {"WPT3",2} } }, - { "GeForce GTX 750", { {"WGS1",64}, {"WPT1",1}, {"VW2",2}, {"WGS2",256}, {"WPT2",2}, {"VW3",2}, {"WGS3",64}, {"WPT3",2} } }, - { "GeForce GTX 750 Ti", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",2}, {"WGS3",256}, {"WPT3",2} } }, - { "GeForce GTX 980", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "GeForce GTX TITAN", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",256}, {"WPT3",1} } }, - { "GeForce GTX TITAN X", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",1}, {"WGS3",128}, {"WPT3",1} } }, - { "Tesla K20m", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Tesla K40m", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",256}, {"WPT3",1} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "GRID K520", { {"WGS1",128}, {"WPT1",1} } }, + { "GeForce GTX 1070", { {"WGS1",64}, {"WPT1",1} } }, + { "GeForce GTX 480", { {"WGS1",256}, {"WPT1",1} } }, + { "GeForce GTX 670", { {"WGS1",128}, {"WPT1",1} } }, + { "GeForce GTX 680", { {"WGS1",128}, {"WPT1",1} } }, + { "GeForce GTX 750", { {"WGS1",64}, {"WPT1",1} } }, + { "GeForce GTX 750 Ti", { {"WGS1",64}, {"WPT1",1} } }, + { "GeForce GTX 980", { {"WGS1",64}, {"WPT1",1} } }, + { "GeForce GTX TITAN", { {"WGS1",256}, {"WPT1",1} } }, + { "GeForce GTX TITAN X", { {"WGS1",64}, {"WPT1",1} } }, + { "Tesla K20m", { {"WGS1",256}, {"WPT1",1} } }, + { "Tesla K40m", { {"WGS1",256}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Default kDeviceTypeAll, "default", { - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, } @@ -208,38 +208,38 @@ const Database::DatabaseEntry Database::XgemvComplexDouble = { "Xgemv", Precision::kComplexDouble, { { // AMD GPUs kDeviceTypeGPU, "AMD", { - { "AMD Radeon R9 M370X Compute Engine", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",128}, {"WPT3",1} } }, - { "Hawaii", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Oland", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",256}, {"WPT2",1}, {"VW3",1}, {"WGS3",256}, {"WPT3",1} } }, - { "Pitcairn", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "Tahiti", { {"WGS1",256}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "AMD Radeon R9 M370X Compute Engine", { {"WGS1",64}, {"WPT1",1} } }, + { "Hawaii", { {"WGS1",64}, {"WPT1",1} } }, + { "Oland", { {"WGS1",256}, {"WPT1",1} } }, + { "Pitcairn", { {"WGS1",256}, {"WPT1",1} } }, + { "Tahiti", { {"WGS1",256}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Intel CPUs kDeviceTypeCPU, "Intel", { - { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"WGS1",64}, {"WPT1",1}, {"VW2",2}, {"WGS2",64}, {"WPT2",4}, {"VW3",1}, {"WGS3",256}, {"WPT3",1} } }, - { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"WGS1",64}, {"WPT1",4}, {"VW2",4}, {"WGS2",64}, {"WPT2",4}, {"VW3",2}, {"WGS3",256}, {"WPT3",2} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",2}, {"WGS2",64}, {"WPT2",4}, {"VW3",1}, {"WGS3",256}, {"WPT3",1} } }, + { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"WGS1",64}, {"WPT1",1} } }, + { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"WGS1",64}, {"WPT1",4} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Intel accelerators kDeviceTypeAccelerator, "Intel", { - { "Intel(R) Many Integrated Core Acceleration Card", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "Intel(R) Many Integrated Core Acceleration Card", { {"WGS1",64}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // NVIDIA GPUs kDeviceTypeGPU, "NVIDIA", { - { "GRID K520", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",1}, {"WGS3",256}, {"WPT3",1} } }, - { "GeForce GTX 480", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "GeForce GTX 670", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "GRID K520", { {"WGS1",128}, {"WPT1",1} } }, + { "GeForce GTX 480", { {"WGS1",64}, {"WPT1",1} } }, + { "GeForce GTX 670", { {"WGS1",128}, {"WPT1",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, { // Default kDeviceTypeAll, "default", { - { "default", { {"WGS1",64}, {"WPT1",1}, {"VW2",1}, {"WGS2",64}, {"WPT2",1}, {"VW3",1}, {"WGS3",64}, {"WPT3",1} } }, + { "default", { {"WGS1",64}, {"WPT1",1} } }, } }, } diff --git a/src/database/kernels/xgemv_fast.hpp b/src/database/kernels/xgemv_fast.hpp new file mode 100644 index 00000000..65b15030 --- /dev/null +++ b/src/database/kernels/xgemv_fast.hpp @@ -0,0 +1,247 @@ + +// ================================================================================================= +// 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 'Xgemv_Fast' kernels. +// +// ================================================================================================= + +namespace clblast { +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvFastHalf = { + "XgemvFast", Precision::kHalf, { + { // Intel GPUs + kDeviceTypeGPU, "Intel", { + { "Intel(R) HD Graphics Skylake ULT GT2", { {"VW2",2}, {"WGS2",128}, {"WPT2",2} } }, + { "default", { {"VW2",2}, {"WGS2",128}, {"WPT2",2} } }, + } + }, + { // Default + kDeviceTypeAll, "default", { + { "default", { {"VW2",2}, {"WGS2",128}, {"WPT2",2} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvFastSingle = { + "XgemvFast", Precision::kSingle, { + { // AMD GPUs + kDeviceTypeGPU, "AMD", { + { "AMD Radeon R9 M370X Compute Engine", { {"VW2",1}, {"WGS2",128}, {"WPT2",1} } }, + { "Hawaii", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Oland", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Pitcairn", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Tahiti", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // Intel CPUs + kDeviceTypeCPU, "Intel", { + { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"VW2",4}, {"WGS2",128}, {"WPT2",4} } }, + { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"VW2",1}, {"WGS2",64}, {"WPT2",4} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",4} } }, + } + }, + { // Intel GPUs + kDeviceTypeGPU, "Intel", { + { "Intel(R) HD Graphics 530", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "Intel(R) HD Graphics Haswell Ultrabook GT2 Mobile", { {"VW2",4}, {"WGS2",128}, {"WPT2",4} } }, + { "Intel(R) HD Graphics Skylake ULT GT2", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "Iris", { {"VW2",1}, {"WGS2",128}, {"WPT2",2} } }, + { "Iris Pro", { {"VW2",1}, {"WGS2",128}, {"WPT2",2} } }, + { "default", { {"VW2",1}, {"WGS2",128}, {"WPT2",1} } }, + } + }, + { // Intel accelerators + kDeviceTypeAccelerator, "Intel", { + { "Intel(R) Many Integrated Core Acceleration Card", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // NVIDIA GPUs + kDeviceTypeGPU, "NVIDIA", { + { "GRID K520", { {"VW2",2}, {"WGS2",256}, {"WPT2",2} } }, + { "GeForce GTX 1070", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "GeForce GTX 480", { {"VW2",1}, {"WGS2",128}, {"WPT2",1} } }, + { "GeForce GTX 670", { {"VW2",2}, {"WGS2",256}, {"WPT2",2} } }, + { "GeForce GTX 680", { {"VW2",1}, {"WGS2",128}, {"WPT2",1} } }, + { "GeForce GTX 750", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "GeForce GTX 750 Ti", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "GeForce GTX 980", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "GeForce GTX TITAN", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "GeForce GTX TITAN X", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Tesla K20m", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "Tesla K40m", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // Default + kDeviceTypeAll, "default", { + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvFastComplexSingle = { + "XgemvFast", Precision::kComplexSingle, { + { // AMD GPUs + kDeviceTypeGPU, "AMD", { + { "AMD Radeon R9 M370X Compute Engine", { {"VW2",2}, {"WGS2",256}, {"WPT2",2} } }, + { "Hawaii", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Oland", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Pitcairn", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Tahiti", { {"VW2",1}, {"WGS2",128}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // Intel CPUs + kDeviceTypeCPU, "Intel", { + { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"VW2",1}, {"WGS2",128}, {"WPT2",2} } }, + { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"VW2",4}, {"WGS2",64}, {"WPT2",4} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",2} } }, + } + }, + { // Intel GPUs + kDeviceTypeGPU, "Intel", { + { "Intel(R) HD Graphics 530", { {"VW2",2}, {"WGS2",128}, {"WPT2",2} } }, + { "Intel(R) HD Graphics Haswell Ultrabook GT2 Mobile", { {"VW2",2}, {"WGS2",128}, {"WPT2",2} } }, + { "Intel(R) HD Graphics Skylake ULT GT2", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Iris", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Iris Pro", { {"VW2",4}, {"WGS2",128}, {"WPT2",4} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // Intel accelerators + kDeviceTypeAccelerator, "Intel", { + { "Intel(R) Many Integrated Core Acceleration Card", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // NVIDIA GPUs + kDeviceTypeGPU, "NVIDIA", { + { "GRID K520", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "GeForce GTX 1070", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "GeForce GTX 480", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "GeForce GTX 670", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "GeForce GTX 680", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "GeForce GTX 750 Ti", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // Default + kDeviceTypeAll, "default", { + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvFastDouble = { + "XgemvFast", Precision::kDouble, { + { // AMD GPUs + kDeviceTypeGPU, "AMD", { + { "AMD Radeon R9 M370X Compute Engine", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "Hawaii", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Oland", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Pitcairn", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Tahiti", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // Intel CPUs + kDeviceTypeCPU, "Intel", { + { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"VW2",4}, {"WGS2",128}, {"WPT2",4} } }, + { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"VW2",1}, {"WGS2",64}, {"WPT2",4} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",4} } }, + } + }, + { // Intel accelerators + kDeviceTypeAccelerator, "Intel", { + { "Intel(R) Many Integrated Core Acceleration Card", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // NVIDIA GPUs + kDeviceTypeGPU, "NVIDIA", { + { "GRID K520", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "GeForce GTX 1070", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "GeForce GTX 480", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "GeForce GTX 670", { {"VW2",1}, {"WGS2",128}, {"WPT2",1} } }, + { "GeForce GTX 680", { {"VW2",1}, {"WGS2",128}, {"WPT2",1} } }, + { "GeForce GTX 750", { {"VW2",2}, {"WGS2",256}, {"WPT2",2} } }, + { "GeForce GTX 750 Ti", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "GeForce GTX 980", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "GeForce GTX TITAN", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "GeForce GTX TITAN X", { {"VW2",1}, {"WGS2",128}, {"WPT2",1} } }, + { "Tesla K20m", { {"VW2",1}, {"WGS2",128}, {"WPT2",1} } }, + { "Tesla K40m", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // Default + kDeviceTypeAll, "default", { + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvFastComplexDouble = { + "XgemvFast", Precision::kComplexDouble, { + { // AMD GPUs + kDeviceTypeGPU, "AMD", { + { "AMD Radeon R9 M370X Compute Engine", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "Hawaii", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Oland", { {"VW2",1}, {"WGS2",256}, {"WPT2",1} } }, + { "Pitcairn", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "Tahiti", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // Intel CPUs + kDeviceTypeCPU, "Intel", { + { "Intel(R) Core(TM) i5-6200U CPU @ 2.30GHz", { {"VW2",2}, {"WGS2",64}, {"WPT2",4} } }, + { "Intel(R) Core(TM) i7-3770 CPU @ 3.40GHz", { {"VW2",4}, {"WGS2",64}, {"WPT2",4} } }, + { "default", { {"VW2",2}, {"WGS2",64}, {"WPT2",4} } }, + } + }, + { // Intel accelerators + kDeviceTypeAccelerator, "Intel", { + { "Intel(R) Many Integrated Core Acceleration Card", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // NVIDIA GPUs + kDeviceTypeGPU, "NVIDIA", { + { "GRID K520", { {"VW2",1}, {"WGS2",128}, {"WPT2",1} } }, + { "GeForce GTX 480", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "GeForce GTX 670", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + { // Default + kDeviceTypeAll, "default", { + { "default", { {"VW2",1}, {"WGS2",64}, {"WPT2",1} } }, + } + }, + } +}; + +// ================================================================================================= +} // namespace clblast diff --git a/src/database/kernels/xgemv_fast_rot.hpp b/src/database/kernels/xgemv_fast_rot.hpp new file mode 100644 index 00000000..9822fb20 --- /dev/null +++ b/src/database/kernels/xgemv_fast_rot.hpp @@ -0,0 +1,138 @@ + +// ================================================================================================= +// 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 'Xgemv_Fast_Rot' kernels. +// +// ================================================================================================= + +namespace clblast { +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvFastRotSingle = { + "XgemvFastRot", Precision::kSingle, { + { // AMD GPUs + kDeviceTypeGPU, "AMD", { + { "AMD Radeon R9 M370X Compute Engine", { {"VW3",8}, {"WGS3",64}, {"WPT3",32} } }, + { "default", { {"VW3",8}, {"WGS3",64}, {"WPT3",32} } }, + } + }, + { // Intel CPUs + kDeviceTypeCPU, "Intel", { + { "Intel(R) Core(TM) i7-5930K CPU @ 3.50GHz", { {"VW3",8}, {"WGS3",16}, {"WPT3",8} } }, + { "default", { {"VW3",8}, {"WGS3",16}, {"WPT3",8} } }, + } + }, + { // Intel GPUs + kDeviceTypeGPU, "Intel", { + { "Intel(R) HD Graphics Skylake ULT GT2", { {"VW3",4}, {"WGS3",128}, {"WPT3",16} } }, + { "Iris Pro", { {"VW3",4}, {"WGS3",32}, {"WPT3",16} } }, + { "default", { {"VW3",4}, {"WGS3",32}, {"WPT3",16} } }, + } + }, + { // NVIDIA GPUs + kDeviceTypeGPU, "NVIDIA", { + { "GeForce GTX TITAN", { {"VW3",1}, {"WGS3",16}, {"WPT3",16} } }, + { "default", { {"VW3",1}, {"WGS3",16}, {"WPT3",16} } }, + } + }, + { // Default + kDeviceTypeAll, "default", { + { "default", { {"VW3",1}, {"WGS3",16}, {"WPT3",8} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvFastRotComplexSingle = { + "XgemvFastRot", Precision::kComplexSingle, { + { // AMD GPUs + kDeviceTypeGPU, "AMD", { + { "AMD Radeon R9 M370X Compute Engine", { {"VW3",8}, {"WGS3",16}, {"WPT3",16} } }, + { "default", { {"VW3",8}, {"WGS3",16}, {"WPT3",16} } }, + } + }, + { // Intel CPUs + kDeviceTypeCPU, "Intel", { + { "Intel(R) Core(TM) i7-5930K CPU @ 3.50GHz", { {"VW3",4}, {"WGS3",16}, {"WPT3",16} } }, + { "default", { {"VW3",4}, {"WGS3",16}, {"WPT3",16} } }, + } + }, + { // Intel GPUs + kDeviceTypeGPU, "Intel", { + { "Intel(R) HD Graphics Skylake ULT GT2", { {"VW3",2}, {"WGS3",32}, {"WPT3",16} } }, + { "Iris Pro", { {"VW3",4}, {"WGS3",16}, {"WPT3",16} } }, + { "default", { {"VW3",2}, {"WGS3",16}, {"WPT3",16} } }, + } + }, + { // Default + kDeviceTypeAll, "default", { + { "default", { {"VW3",2}, {"WGS3",16}, {"WPT3",16} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvFastRotDouble = { + "XgemvFastRot", Precision::kDouble, { + { // AMD GPUs + kDeviceTypeGPU, "AMD", { + { "AMD Radeon R9 M370X Compute Engine", { {"VW3",4}, {"WGS3",16}, {"WPT3",16} } }, + { "default", { {"VW3",4}, {"WGS3",16}, {"WPT3",16} } }, + } + }, + { // Intel CPUs + kDeviceTypeCPU, "Intel", { + { "Intel(R) Core(TM) i7-5930K CPU @ 3.50GHz", { {"VW3",8}, {"WGS3",16}, {"WPT3",8} } }, + { "default", { {"VW3",8}, {"WGS3",16}, {"WPT3",8} } }, + } + }, + { // NVIDIA GPUs + kDeviceTypeGPU, "NVIDIA", { + { "GeForce GTX TITAN", { {"VW3",1}, {"WGS3",16}, {"WPT3",16} } }, + { "default", { {"VW3",1}, {"WGS3",16}, {"WPT3",16} } }, + } + }, + { // Default + kDeviceTypeAll, "default", { + { "default", { {"VW3",1}, {"WGS3",16}, {"WPT3",8} } }, + } + }, + } +}; + +// ================================================================================================= + +const Database::DatabaseEntry Database::XgemvFastRotComplexDouble = { + "XgemvFastRot", Precision::kComplexDouble, { + { // AMD GPUs + kDeviceTypeGPU, "AMD", { + { "AMD Radeon R9 M370X Compute Engine", { {"VW3",4}, {"WGS3",32}, {"WPT3",16} } }, + { "default", { {"VW3",4}, {"WGS3",32}, {"WPT3",16} } }, + } + }, + { // Intel CPUs + kDeviceTypeCPU, "Intel", { + { "Intel(R) Core(TM) i7-5930K CPU @ 3.50GHz", { {"VW3",8}, {"WGS3",16}, {"WPT3",16} } }, + { "default", { {"VW3",8}, {"WGS3",16}, {"WPT3",16} } }, + } + }, + { // Default + kDeviceTypeAll, "default", { + { "default", { {"VW3",4}, {"WGS3",16}, {"WPT3",16} } }, + } + }, + } +}; + +// ================================================================================================= +} // namespace clblast diff --git a/src/kernels/level2/xgemv_fast.opencl b/src/kernels/level2/xgemv_fast.opencl index 1127a0b6..210c42c1 100644 --- a/src/kernels/level2/xgemv_fast.opencl +++ b/src/kernels/level2/xgemv_fast.opencl @@ -38,7 +38,7 @@ R"( #define WGS3 64 // The local work-group size #endif #ifndef WPT3 - #define WPT3 1 // The amount of work-per-thread + #define WPT3 1 // The tile-size #endif #ifndef VW3 #define VW3 1 // Vector width of matrix A loads @@ -74,18 +74,12 @@ R"( // ================================================================================================= -// Loads a vector input value (1/2) +// Loads a vector input value inline realVF LoadMatrixAVF(const __global realVF* restrict agm, const int x, const int y, const int a_ld) { return agm[a_ld*y + x]; } -// Loads a vector input value (2/2): as before, but different data-type -inline realVFR LoadMatrixAVFR(const __global realVFR* restrict agm, const int x, const int y, - const int a_ld) { - return agm[a_ld*y + x]; -} - // ================================================================================================= // Faster version of the kernel, assuming that: @@ -103,14 +97,14 @@ __kernel void XgemvFast(const int m, const int n, const __global real* restrict xgm, const int x_offset, const int x_inc, __global real* ygm, const int y_offset, const int y_inc, const int do_conjugate, const int parameter, - const int kl, const int ku) { + const int kl_unused, const int ku_unused) { const real alpha = GetRealArg(arg_alpha); const real beta = GetRealArg(arg_beta); // Local memory for the vector X __local real xlm[WGS2]; - // Initializes the accumulation register + // Initializes the accumulation registers real acc[WPT2]; #pragma unroll for (int w=0; w Xgemv::Xgemv(Queue &queue, EventPointer event, const std::string &name): - Routine(queue, event, name, {"Pad", "Xgemv"}, PrecisionValue()) { + Routine(queue, event, name, {"Pad", "Xgemv", "XgemvFast", "XgemvFastRot"}, PrecisionValue()) { source_string_ = #include "../../kernels/level2/xgemv.opencl" #include "../../kernels/level2/xgemv_fast.opencl" @@ -122,7 +122,7 @@ StatusCode Xgemv::MatVec(const Layout layout, const Transpose a_transpose, } if (fast_kernel_rot) { kernel_name = "XgemvFastRot"; - global_size = m_real / db_["WPT3"]; + global_size = m_real; local_size = db_["WGS3"]; } diff --git a/src/tuning/kernels/xgemv.cpp b/src/tuning/kernels/xgemv.cpp index 5c187d33..7229602d 100644 --- a/src/tuning/kernels/xgemv.cpp +++ b/src/tuning/kernels/xgemv.cpp @@ -29,7 +29,7 @@ class TuneXgemv { public: // The representative kernel and the source code - static std::string KernelFamily() { return "xgemv_"+std::to_string(V); } + static std::string KernelFamily() { return (V==1) ? "xgemv" : ((V==2) ? "xgemv_fast" : "xgemv_fast_rot"); } static std::string KernelName() { return (V==1) ? "Xgemv" : ((V==2) ? "XgemvFast" : "XgemvFastRot"); } static std::string GetSources() { return @@ -61,21 +61,42 @@ class TuneXgemv { // Sets the tuning parameters and their possible values static void SetParameters(cltune::Tuner &tuner, const size_t id) { - tuner.AddParameter(id, "WGS"+std::to_string(V), {64, 128, 256}); - tuner.AddParameter(id, "WPT"+std::to_string(V), {1, 2, 4}); - if (V==2 || V==3) { tuner.AddParameter(id, "VW"+std::to_string(V), {1, 2, 4, 8}); } + if (V==1) { + tuner.AddParameter(id, "WGS"+std::to_string(V), {32, 64, 128, 256}); + tuner.AddParameter(id, "WPT"+std::to_string(V), {1, 2, 4}); + } + if (V==2) { + tuner.AddParameter(id, "WGS"+std::to_string(V), {16, 32, 64, 128, 256}); + tuner.AddParameter(id, "WPT"+std::to_string(V), {1, 2, 4}); + tuner.AddParameter(id, "VW"+std::to_string(V), {1, 2, 4, 8}); + } + if (V==3) { + tuner.AddParameter(id, "WGS"+std::to_string(V), {16, 32, 64, 128}); + tuner.AddParameter(id, "WPT"+std::to_string(V), {1, 2, 4, 8, 16, 32}); + tuner.AddParameter(id, "VW"+std::to_string(V), {1, 2, 4, 8}); + } } // Sets the constraints and local memory size static void SetConstraints(cltune::Tuner &tuner, const size_t id) { - auto MultipleOfX = [] (std::vector v) { return IsMultiple(v[0], v[1]); }; if (V==2 || V==3) { + auto MultipleOfX = [] (std::vector v) { return IsMultiple(v[0], v[1]); }; tuner.AddConstraint(id, MultipleOfX, {"WPT"+std::to_string(V), "VW"+std::to_string(V)}); } + if (V==3) { + auto LargerOrEqual = [] (std::vector v) { return v[0] >= v[1]; }; + tuner.AddConstraint(id, LargerOrEqual, {"WGS"+std::to_string(V), "WPT"+std::to_string(V)}); + } } static void SetLocalMemorySize(cltune::Tuner &tuner, const size_t id, const Arguments &args) { - auto LocalMemorySize = [args] (std::vector v) { return v[0]*GetBytes(args.precision); }; - tuner.SetLocalMemoryUsage(id, LocalMemorySize, {"WGS"+std::to_string(V)}); + if (V==1 || V==2) { + auto LocalMemorySize = [args] (std::vector v) { return v[0]*GetBytes(args.precision); }; + tuner.SetLocalMemoryUsage(id, LocalMemorySize, {"WGS"+std::to_string(V)}); + } + else { + auto LocalMemorySize = [args] (std::vector v) { return (v[0]*v[1] + v[1])*GetBytes(args.precision); }; + tuner.SetLocalMemoryUsage(id, LocalMemorySize, {"WGS"+std::to_string(V), "WPT"+std::to_string(V)}); + } } // Sets the base thread configuration @@ -89,7 +110,10 @@ class TuneXgemv { static TransformVector MulLocal() { return {{"WGS"+std::to_string(V)}}; } static TransformVector DivLocal() { return {}; } static TransformVector MulGlobal() { return {}; } - static TransformVector DivGlobal() { return {{"WPT"+std::to_string(V)}}; } + static TransformVector DivGlobal() { + if (V==1 || V==2) return {{"WPT"+std::to_string(V)}}; + return {}; + } // Sets the kernel's arguments static void SetArguments(cltune::Tuner &tuner, const Arguments &args,