mirror of
https://github.com/CNugteren/CLBlast.git
synced 2024-07-07 12:23:46 +02:00
Merge branch 'gemv_performance' into development
This commit is contained in:
Cedric Nugteren
commit
1ec21421d7
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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):
|
||||
|
|
|
|||
|
|
@ -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::DatabaseEntry> 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,
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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} } },
|
||||
}
|
||||
},
|
||||
}
|
||||
|
|
|
|||
247
src/database/kernels/xgemv_fast.hpp
Normal file
247
src/database/kernels/xgemv_fast.hpp
Normal file
|
|
@ -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 <database.py>
|
||||
//
|
||||
// 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
|
||||
138
src/database/kernels/xgemv_fast_rot.hpp
Normal file
138
src/database/kernels/xgemv_fast_rot.hpp
Normal file
|
|
@ -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 <database.py>
|
||||
//
|
||||
// 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
|
||||
|
|
@ -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<WPT2; ++w) {
|
||||
|
|
@ -134,7 +128,7 @@ __kernel void XgemvFast(const int m, const int n,
|
|||
#pragma unroll
|
||||
for (int w=0; w<WPT2/VW2; ++w) {
|
||||
const int gid = (WPT2/VW2)*get_global_id(0) + w;
|
||||
realVF avec = LoadMatrixAVF(agm, gid, k, a_ld/VW2);
|
||||
realVF avec = agm[(a_ld/VW2)*k + gid];
|
||||
#if VW2 == 1
|
||||
MultiplyAdd(acc[VW2*w+0], xlm[kl], avec);
|
||||
#elif VW2 == 2
|
||||
|
|
@ -205,75 +199,87 @@ __kernel void XgemvFastRot(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 to store a tile of the matrix (for coalescing)
|
||||
__local real tile[WPT3][WGS3];
|
||||
const int lid = get_local_id(0);
|
||||
const int lid_mod = lid % (WPT3/VW3);
|
||||
const int lid_div = lid / (WPT3/VW3);
|
||||
|
||||
// Local memory for the vector X
|
||||
__local real xlm[WGS3];
|
||||
__local real xlm[WPT3];
|
||||
|
||||
// Initializes the accumulation register
|
||||
real acc[WPT3];
|
||||
#pragma unroll
|
||||
for (int w=0; w<WPT3; ++w) {
|
||||
SetToZero(acc[w]);
|
||||
}
|
||||
real acc;
|
||||
SetToZero(acc);
|
||||
|
||||
// Loops over work-group sized portions of the work
|
||||
for (int kwg=0; kwg<n; kwg+=WGS3) {
|
||||
// Loops over tile-sized portions of the work
|
||||
for (int kwg=0; kwg<n; kwg+=WPT3) {
|
||||
|
||||
// Loads the vector X into local memory
|
||||
const int lid = get_local_id(0);
|
||||
xlm[lid] = xgm[(kwg + lid)*x_inc + x_offset];
|
||||
if (lid < WPT3) {
|
||||
xlm[lid] = xgm[(kwg + lid) * x_inc + x_offset];
|
||||
}
|
||||
|
||||
// Loads the matrix A into local memory
|
||||
#pragma unroll
|
||||
for (int kl=0; kl<WPT3/VW3; ++kl) {
|
||||
const int x = (kwg/VW3) + lid_mod;
|
||||
const int y = get_group_id(0) * WGS3 + lid_div * (WPT3/VW3) + kl;
|
||||
realVFR avec = agm[(a_ld/VW3) * y + x];
|
||||
#if VW3 == 1
|
||||
tile[kl*VW3 + 0][lid] = avec;
|
||||
#elif VW3 == 2
|
||||
tile[kl*VW3 + 0][lid] = avec.x;
|
||||
tile[kl*VW3 + 1][lid] = avec.y;
|
||||
#elif VW3 == 4
|
||||
tile[kl*VW3 + 0][lid] = avec.x;
|
||||
tile[kl*VW3 + 1][lid] = avec.y;
|
||||
tile[kl*VW3 + 2][lid] = avec.z;
|
||||
tile[kl*VW3 + 3][lid] = avec.w;
|
||||
#elif VW3 == 8
|
||||
tile[kl*VW3 + 0][lid] = avec.s0;
|
||||
tile[kl*VW3 + 1][lid] = avec.s1;
|
||||
tile[kl*VW3 + 2][lid] = avec.s2;
|
||||
tile[kl*VW3 + 3][lid] = avec.s3;
|
||||
tile[kl*VW3 + 4][lid] = avec.s4;
|
||||
tile[kl*VW3 + 5][lid] = avec.s5;
|
||||
tile[kl*VW3 + 6][lid] = avec.s6;
|
||||
tile[kl*VW3 + 7][lid] = avec.s7;
|
||||
#elif VW3 == 16
|
||||
tile[kl*VW3 + 0][lid] = avec.s0;
|
||||
tile[kl*VW3 + 1][lid] = avec.s1;
|
||||
tile[kl*VW3 + 2][lid] = avec.s2;
|
||||
tile[kl*VW3 + 3][lid] = avec.s3;
|
||||
tile[kl*VW3 + 4][lid] = avec.s4;
|
||||
tile[kl*VW3 + 5][lid] = avec.s5;
|
||||
tile[kl*VW3 + 6][lid] = avec.s6;
|
||||
tile[kl*VW3 + 7][lid] = avec.s7;
|
||||
tile[kl*VW3 + 8][lid] = avec.s8;
|
||||
tile[kl*VW3 + 9][lid] = avec.s9;
|
||||
tile[kl*VW3 + 10][lid] = avec.sA;
|
||||
tile[kl*VW3 + 11][lid] = avec.sB;
|
||||
tile[kl*VW3 + 12][lid] = avec.sC;
|
||||
tile[kl*VW3 + 13][lid] = avec.sD;
|
||||
tile[kl*VW3 + 14][lid] = avec.sE;
|
||||
tile[kl*VW3 + 15][lid] = avec.sF;
|
||||
#endif
|
||||
}
|
||||
|
||||
// Synchronizes all threads in a workgroup
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
// The multiply-add function (rotated)
|
||||
#pragma unroll
|
||||
for (int kl=0; kl<WGS3/VW3; ++kl) {
|
||||
const int k = (kwg/VW3) + kl;
|
||||
for (int kl=0; kl<WPT3/VW3; ++kl) {
|
||||
#pragma unroll
|
||||
for (int w=0; w<WPT3; ++w) {
|
||||
const int gid = WPT3*get_global_id(0) + w;
|
||||
realVFR avec = LoadMatrixAVFR(agm, k, gid, a_ld/VW3);
|
||||
#if VW3 == 1
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+0], avec);
|
||||
#elif VW3 == 2
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+0], avec.x);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+1], avec.y);
|
||||
#elif VW3 == 4
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+0], avec.x);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+1], avec.y);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+2], avec.z);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+3], avec.w);
|
||||
#elif VW3 == 8
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+0], avec.s0);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+1], avec.s1);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+2], avec.s2);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+3], avec.s3);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+4], avec.s4);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+5], avec.s5);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+6], avec.s6);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+7], avec.s7);
|
||||
#elif VW3 == 16
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+0], avec.s0);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+1], avec.s1);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+2], avec.s2);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+3], avec.s3);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+4], avec.s4);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+5], avec.s5);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+6], avec.s6);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+7], avec.s7);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+8], avec.s8);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+9], avec.s9);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+10], avec.sA);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+11], avec.sB);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+12], avec.sC);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+13], avec.sD);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+14], avec.sE);
|
||||
MultiplyAdd(acc[w], xlm[VW3*kl+15], avec.sF);
|
||||
#endif
|
||||
for (int v=0; v<VW3; ++v) {
|
||||
real aval = tile[lid_mod*VW3 + v][lid_div * (WPT3/VW3) + kl];
|
||||
real xval = xlm[kl*VW3 + v];
|
||||
MultiplyAdd(acc, xval, aval);
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -282,12 +288,9 @@ __kernel void XgemvFastRot(const int m, const int n,
|
|||
}
|
||||
|
||||
// Stores the final result
|
||||
#pragma unroll
|
||||
for (int w=0; w<WPT3; ++w) {
|
||||
const int gid = WPT3*get_global_id(0) + w;
|
||||
real yval = ygm[gid*y_inc + y_offset];
|
||||
AXPBY(ygm[gid*y_inc + y_offset], alpha, acc[w], beta, yval);
|
||||
}
|
||||
const int gid = get_global_id(0);
|
||||
real yval = ygm[gid * y_inc + y_offset];
|
||||
AXPBY(ygm[gid * y_inc + y_offset], alpha, acc, beta, yval);
|
||||
}
|
||||
|
||||
// =================================================================================================
|
||||
|
|
|
|||
|
|
@ -22,7 +22,7 @@ namespace clblast {
|
|||
// Constructor: forwards to base class constructor
|
||||
template <typename T>
|
||||
Xgemv<T>::Xgemv(Queue &queue, EventPointer event, const std::string &name):
|
||||
Routine(queue, event, name, {"Pad", "Xgemv"}, PrecisionValue<T>()) {
|
||||
Routine(queue, event, name, {"Pad", "Xgemv", "XgemvFast", "XgemvFastRot"}, PrecisionValue<T>()) {
|
||||
source_string_ =
|
||||
#include "../../kernels/level2/xgemv.opencl"
|
||||
#include "../../kernels/level2/xgemv_fast.opencl"
|
||||
|
|
@ -122,7 +122,7 @@ StatusCode Xgemv<T>::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"];
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -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<size_t> v) { return IsMultiple(v[0], v[1]); };
|
||||
if (V==2 || V==3) {
|
||||
auto MultipleOfX = [] (std::vector<size_t> 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<size_t> 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<T> &args) {
|
||||
auto LocalMemorySize = [args] (std::vector<size_t> 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<size_t> v) { return v[0]*GetBytes(args.precision); };
|
||||
tuner.SetLocalMemoryUsage(id, LocalMemorySize, {"WGS"+std::to_string(V)});
|
||||
}
|
||||
else {
|
||||
auto LocalMemorySize = [args] (std::vector<size_t> 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<T> &args,
|
||||
|
|
|
|||
Loading…
Reference in a new issue