summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
Diffstat (limited to 'src')
-rw-r--r--src/cupp11.hpp4
-rw-r--r--src/database/apple_cpu_fallback.hpp7
-rw-r--r--src/database/database.cpp6
-rw-r--r--src/kernels/level2/xtrsv.opencl2
-rw-r--r--src/kernels/level3/invert_diagonal_blocks_part2.opencl24
-rw-r--r--src/kernels/level3/level3.opencl2
-rw-r--r--src/routines/common.cpp70
-rw-r--r--src/routines/common.hpp10
-rw-r--r--src/routines/level2/xtrsv.cpp11
-rw-r--r--src/routines/level3/xgemm.hpp6
-rw-r--r--src/routines/level3/xtrsm.cpp9
-rw-r--r--src/routines/levelx/xinvert.cpp23
-rw-r--r--src/tuning/routines/xgemm.cpp100
-rw-r--r--src/utilities/utilities.hpp2
14 files changed, 189 insertions, 87 deletions
diff --git a/src/cupp11.hpp b/src/cupp11.hpp
index 509ae3e8..a1cb1614 100644
--- a/src/cupp11.hpp
+++ b/src/cupp11.hpp
@@ -678,8 +678,8 @@ public:
}
// Regular constructor with memory management
- explicit Kernel(const Program &program, const std::string &name): name_(name) {
- CheckError(cuModuleGetFunction(&kernel_, program.GetModule(), name.c_str()));
+ explicit Kernel(const std::shared_ptr<Program> program, const std::string &name): name_(name) {
+ CheckError(cuModuleGetFunction(&kernel_, program->GetModule(), name.c_str()));
}
// Sets a kernel argument at the indicated position. This stores both the value of the argument
diff --git a/src/database/apple_cpu_fallback.hpp b/src/database/apple_cpu_fallback.hpp
index fdd9327d..55bcc220 100644
--- a/src/database/apple_cpu_fallback.hpp
+++ b/src/database/apple_cpu_fallback.hpp
@@ -41,7 +41,7 @@ const DatabaseEntry XgerApple = {
"Xger", Precision::kAny, {"WGS1", "WGS2", "WPT"}, { { kDeviceTypeAll, "default", { { "default", { { kDeviceNameDefault, Params{ 64, 1, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 } } } } } } }
};
const DatabaseEntry XtrsvApple = {
- "Xtrsv", Precision::kAny, {"TRSV_BLOCK_SIZE"}, { { kDeviceTypeAll, "default", { { "default", { { kDeviceNameDefault, Params{ 16, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 } } } } } } }
+ "Xtrsv", Precision::kAny, {"TRSV_BLOCK_SIZE"}, { { kDeviceTypeAll, "default", { { "default", { { kDeviceNameDefault, Params{ 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 } } } } } } }
};
const DatabaseEntry XgemmApple = {
"Xgemm", Precision::kAny, {"GEMMK", "KREG", "KWG", "KWI", "MDIMA", "MDIMC", "MWG", "NDIMB", "NDIMC", "NWG", "SA", "SB", "STRM", "STRN", "VWM", "VWN"}, { { kDeviceTypeAll, "default", { { "default", { { kDeviceNameDefault, Params{ 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 1, 1 } } } } } } }
@@ -62,7 +62,10 @@ const DatabaseEntry PadtransposeApple = {
"Padtranspose", Precision::kAny, {"PADTRA_PAD", "PADTRA_TILE", "PADTRA_WPT"}, { { kDeviceTypeAll, "default", { { "default", { { kDeviceNameDefault, Params{ 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 } } } } } } }
};
const DatabaseEntry InvertApple = {
- "Invert", Precision::kAny, {"INTERNAL_BLOCK_SIZE"}, { { kDeviceTypeAll, "default", { { "default", { { kDeviceNameDefault, Params{ 16, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 } } } } } } }
+ "Invert", Precision::kAny, {"INTERNAL_BLOCK_SIZE"}, { { kDeviceTypeAll, "default", { { "default", { { kDeviceNameDefault, Params{ 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 } } } } } } }
+};
+const DatabaseEntry TrsvRoutineApple = {
+ "TrsvRoutine", Precision::kAny, {"TRSV_BLOCK_SIZE"}, { { kDeviceTypeAll, "default", { { "default", { { kDeviceNameDefault, Params{ 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 } } } } } } }
};
// =================================================================================================
diff --git a/src/database/database.cpp b/src/database/database.cpp
index b2f70e49..fca3102d 100644
--- a/src/database/database.cpp
+++ b/src/database/database.cpp
@@ -45,7 +45,8 @@ const std::vector<database::DatabaseEntry> Database::apple_cpu_fallback = std::v
database::XgemvApple, database::XgemvFastApple, database::XgemvFastRotApple, database::XgerApple, database::XtrsvApple,
database::XgemmApple, database::XgemmDirectApple,
database::CopyApple, database::PadApple, database::TransposeApple, database::PadtransposeApple,
- database::InvertApple
+ database::InvertApple,
+ database::TrsvRoutineApple
};
// The default values
@@ -98,7 +99,8 @@ Database::Database(const Device &device, const std::string &kernel_name,
if (device.Type() == "CPU") {
const auto extensions = device.Capabilities();
const auto is_apple = (extensions.find("cl_APPLE_SetMemObjectDestructor") == std::string::npos) ? false : true;
- if (is_apple) {
+ const auto is_likely_apple = device.MaxWorkGroupSize() <= 32;
+ if (is_apple || is_likely_apple) {
databases.push_front(apple_cpu_fallback);
}
}
diff --git a/src/kernels/level2/xtrsv.opencl b/src/kernels/level2/xtrsv.opencl
index 8777eb77..e7b6ae79 100644
--- a/src/kernels/level2/xtrsv.opencl
+++ b/src/kernels/level2/xtrsv.opencl
@@ -18,7 +18,7 @@ R"(
// =================================================================================================
#if defined(ROUTINE_TRSV)
-__kernel __attribute__((reqd_work_group_size(16, 1, 1)))
+__kernel
void FillVector(const int n, const int inc, const int offset,
__global real* restrict dest, const real_arg arg_value) {
const real value = GetRealArg(arg_value);
diff --git a/src/kernels/level3/invert_diagonal_blocks_part2.opencl b/src/kernels/level3/invert_diagonal_blocks_part2.opencl
index 8736203c..8e9b583e 100644
--- a/src/kernels/level3/invert_diagonal_blocks_part2.opencl
+++ b/src/kernels/level3/invert_diagonal_blocks_part2.opencl
@@ -19,7 +19,7 @@ R"(
#if defined(ROUTINE_INVERT)
// B21 = A21 * B11
-__kernel __attribute__((reqd_work_group_size(1 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul16Part1Lower(int n, __global const real* restrict src, const int a_offset, const int lda,
__global real* restrict dest, int current_size, int num_pages, const int block_size)
{
@@ -28,7 +28,7 @@ void TripleMatMul16Part1Lower(int n, __global const real* restrict src, const in
}
// B21 = -B22 * B21
-__kernel __attribute__((reqd_work_group_size(1 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul16Part2Lower(int n, __global real* restrict dest, int current_size, int num_pages, const int block_size)
{
__local real lm[LOCALY * LOCALX];
@@ -36,7 +36,7 @@ void TripleMatMul16Part2Lower(int n, __global real* restrict dest, int current_s
}
// B21 = A21 * B11
-__kernel __attribute__((reqd_work_group_size(2 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul32Part1Lower(int n, __global const real* restrict src, const int a_offset, const int lda,
__global real* restrict dest, int current_size, int num_pages, const int block_size)
{
@@ -45,7 +45,7 @@ void TripleMatMul32Part1Lower(int n, __global const real* restrict src, const in
}
// B21 = -B22 * B21
-__kernel __attribute__((reqd_work_group_size(2 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul32Part2Lower(int n, __global real* restrict dest, int current_size, int num_pages, const int block_size)
{
__local real lm[LOCALY * LOCALX];
@@ -53,7 +53,7 @@ void TripleMatMul32Part2Lower(int n, __global real* restrict dest, int current_s
}
// B21 = A21 * B11
-__kernel __attribute__((reqd_work_group_size(4 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul64Part1Lower(int n, __global const real* restrict src, const int a_offset, const int lda,
__global real* restrict dest, int current_size, int num_pages, const int block_size)
{
@@ -62,7 +62,7 @@ void TripleMatMul64Part1Lower(int n, __global const real* restrict src, const in
}
// B21 = -B22 * B21
-__kernel __attribute__((reqd_work_group_size(4 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul64Part2Lower(int n, __global real* restrict dest, int current_size, int num_pages, const int block_size)
{
__local real lm[LOCALY * LOCALX];
@@ -72,7 +72,7 @@ void TripleMatMul64Part2Lower(int n, __global real* restrict dest, int current_s
// =================================================================================================
// B12 = A12 * B22
-__kernel __attribute__((reqd_work_group_size(1 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul16Part1Upper(int n, __global const real* restrict src, const int a_offset, const int lda,
__global real* restrict dest, int current_size, int num_pages, const int block_size)
{
@@ -81,7 +81,7 @@ void TripleMatMul16Part1Upper(int n, __global const real* restrict src, const in
}
// B12 = -B11 * B12
-__kernel __attribute__((reqd_work_group_size(1 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul16Part2Upper(int n, __global real* restrict dest, int current_size, int num_pages, const int block_size)
{
__local real lm[LOCALY * LOCALX];
@@ -89,7 +89,7 @@ void TripleMatMul16Part2Upper(int n, __global real* restrict dest, int current_s
}
// B12 = A12 * B22
-__kernel __attribute__((reqd_work_group_size(2 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul32Part1Upper(int n, __global const real* restrict src, const int a_offset, const int lda,
__global real* restrict dest, int current_size, int num_pages, const int block_size)
{
@@ -98,7 +98,7 @@ void TripleMatMul32Part1Upper(int n, __global const real* restrict src, const in
}
// B12 = -B11 * B12
-__kernel __attribute__((reqd_work_group_size(2 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul32Part2Upper(int n, __global real* restrict dest, int current_size, int num_pages, const int block_size)
{
__local real lm[LOCALY * LOCALX];
@@ -106,7 +106,7 @@ void TripleMatMul32Part2Upper(int n, __global real* restrict dest, int current_s
}
// B12 = A12 * B22
-__kernel __attribute__((reqd_work_group_size(4 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul64Part1Upper(int n, __global const real* restrict src, const int a_offset, const int lda,
__global real* restrict dest, int current_size, int num_pages, const int block_size)
{
@@ -115,7 +115,7 @@ void TripleMatMul64Part1Upper(int n, __global const real* restrict src, const in
}
// B12 = -B11 * B12
-__kernel __attribute__((reqd_work_group_size(4 * TMMWGSX, TMMWGSY, 1)))
+__kernel
void TripleMatMul64Part2Upper(int n, __global real* restrict dest, int current_size, int num_pages, const int block_size)
{
__local real lm[LOCALY * LOCALX];
diff --git a/src/kernels/level3/level3.opencl b/src/kernels/level3/level3.opencl
index c67851df..bea73daf 100644
--- a/src/kernels/level3/level3.opencl
+++ b/src/kernels/level3/level3.opencl
@@ -76,7 +76,7 @@ R"(
// =================================================================================================
#if defined(ROUTINE_INVERT) || defined(ROUTINE_TRSM)
-__kernel __attribute__((reqd_work_group_size(16, 1, 1)))
+__kernel
void FillMatrix(const int m, const int n, const int ld, const int offset,
__global real* restrict dest, const real_arg arg_value) {
const real value = GetRealArg(arg_value);
diff --git a/src/routines/common.cpp b/src/routines/common.cpp
index 5b80e3f2..695785c4 100644
--- a/src/routines/common.cpp
+++ b/src/routines/common.cpp
@@ -13,6 +13,7 @@
#include <vector>
#include <chrono>
+#include <iostream>
#include "routines/common.hpp"
@@ -38,13 +39,22 @@ void RunKernel(Kernel &kernel, Queue &queue, const Device &device,
auto local_size = size_t{1};
for (auto &item: local) { local_size *= item; }
if (local_size > device.MaxWorkGroupSize()) {
- throw RuntimeErrorCode(StatusCode::kInvalidLocalThreadsTotal);
+ throw RuntimeErrorCode(StatusCode::kInvalidLocalThreadsTotal,
+ ToString(local_size) + " is larger than " + ToString(device.MaxWorkGroupSize()));
}
// Make sure the global thread sizes are at least equal to the local sizes
for (auto i=size_t{0}; i<global.size(); ++i) {
if (global[i] < local[i]) { global[i] = local[i]; }
}
+
+ // Verify that the global thread sizes are a multiple of the local sizes
+ for (auto i=size_t{0}; i<global.size(); ++i) {
+ if ((global[i] / local[i]) * local[i] != global[i]) {
+ throw RuntimeErrorCode(StatusCode::kInvalidLocalThreadsDim,
+ ToString(global[i]) + " is not divisible by " + ToString(local[i]));
+ }
+ }
}
// Tests for local memory usage
@@ -77,11 +87,10 @@ void RunKernel(Kernel &kernel, Queue &queue, const Device &device,
// Sets all elements of a matrix to a constant value
template <typename T>
void FillMatrix(Queue &queue, const Device &device,
- const std::shared_ptr<Program> program, const Databases &,
+ const std::shared_ptr<Program> program,
EventPointer event, const std::vector<Event> &waitForEvents,
const size_t m, const size_t n, const size_t ld, const size_t offset,
- const Buffer<T> &dest,
- const T constant_value) {
+ const Buffer<T> &dest, const T constant_value, const size_t local_size) {
auto kernel = Kernel(program, "FillMatrix");
kernel.SetArgument(0, static_cast<int>(m));
kernel.SetArgument(1, static_cast<int>(n));
@@ -89,63 +98,62 @@ void FillMatrix(Queue &queue, const Device &device,
kernel.SetArgument(3, static_cast<int>(offset));
kernel.SetArgument(4, dest());
kernel.SetArgument(5, GetRealArg(constant_value));
- auto local = std::vector<size_t>{16, 1};
- auto global = std::vector<size_t>{Ceil(m, 16), n};
+ auto local = std::vector<size_t>{local_size, 1};
+ auto global = std::vector<size_t>{Ceil(m, local_size), n};
RunKernel(kernel, queue, device, global, local, event, waitForEvents);
}
// Compiles the above function
-template void FillMatrix<half>(Queue&, const Device&, const std::shared_ptr<Program>, const Databases&,
+template void FillMatrix<half>(Queue&, const Device&, const std::shared_ptr<Program>,
EventPointer, const std::vector<Event>&, const size_t, const size_t,
- const size_t, const size_t, const Buffer<half>&, const half);
-template void FillMatrix<float>(Queue&, const Device&, const std::shared_ptr<Program>, const Databases&,
+ const size_t, const size_t, const Buffer<half>&, const half, const size_t);
+template void FillMatrix<float>(Queue&, const Device&, const std::shared_ptr<Program>,
EventPointer, const std::vector<Event>&, const size_t, const size_t,
- const size_t, const size_t, const Buffer<float>&, const float);
-template void FillMatrix<double>(Queue&, const Device&, const std::shared_ptr<Program>, const Databases&,
+ const size_t, const size_t, const Buffer<float>&, const float, const size_t);
+template void FillMatrix<double>(Queue&, const Device&, const std::shared_ptr<Program>,
EventPointer, const std::vector<Event>&, const size_t, const size_t,
- const size_t, const size_t, const Buffer<double>&, const double);
-template void FillMatrix<float2>(Queue&, const Device&, const std::shared_ptr<Program>, const Databases&,
+ const size_t, const size_t, const Buffer<double>&, const double, const size_t);
+template void FillMatrix<float2>(Queue&, const Device&, const std::shared_ptr<Program>,
EventPointer, const std::vector<Event>&, const size_t, const size_t,
- const size_t, const size_t, const Buffer<float2>&, const float2);
-template void FillMatrix<double2>(Queue&, const Device&, const std::shared_ptr<Program>, const Databases&,
+ const size_t, const size_t, const Buffer<float2>&, const float2, const size_t);
+template void FillMatrix<double2>(Queue&, const Device&, const std::shared_ptr<Program>,
EventPointer, const std::vector<Event>&, const size_t, const size_t,
- const size_t, const size_t, const Buffer<double2>&, const double2);
+ const size_t, const size_t, const Buffer<double2>&, const double2, const size_t);
// Sets all elements of a vector to a constant value
template <typename T>
void FillVector(Queue &queue, const Device &device,
- const std::shared_ptr<Program> program, const Databases &,
+ const std::shared_ptr<Program> program,
EventPointer event, const std::vector<Event> &waitForEvents,
const size_t n, const size_t inc, const size_t offset,
- const Buffer<T> &dest,
- const T constant_value) {
+ const Buffer<T> &dest, const T constant_value, const size_t local_size) {
auto kernel = Kernel(program, "FillVector");
kernel.SetArgument(0, static_cast<int>(n));
kernel.SetArgument(1, static_cast<int>(inc));
kernel.SetArgument(2, static_cast<int>(offset));
kernel.SetArgument(3, dest());
kernel.SetArgument(4, GetRealArg(constant_value));
- auto local = std::vector<size_t>{16};
- auto global = std::vector<size_t>{Ceil(n, 16)};
+ auto local = std::vector<size_t>{local_size};
+ auto global = std::vector<size_t>{Ceil(n, local_size)};
RunKernel(kernel, queue, device, global, local, event, waitForEvents);
}
// Compiles the above function
-template void FillVector<half>(Queue&, const Device&, const std::shared_ptr<Program>, const Databases&,
+template void FillVector<half>(Queue&, const Device&, const std::shared_ptr<Program>,
EventPointer, const std::vector<Event>&, const size_t, const size_t,
- const size_t, const Buffer<half>&, const half);
-template void FillVector<float>(Queue&, const Device&, const std::shared_ptr<Program>, const Databases&,
+ const size_t, const Buffer<half>&, const half, const size_t);
+template void FillVector<float>(Queue&, const Device&, const std::shared_ptr<Program>,
EventPointer, const std::vector<Event>&, const size_t, const size_t,
- const size_t, const Buffer<float>&, const float);
-template void FillVector<double>(Queue&, const Device&, const std::shared_ptr<Program>, const Databases&,
+ const size_t, const Buffer<float>&, const float, const size_t);
+template void FillVector<double>(Queue&, const Device&, const std::shared_ptr<Program>,
EventPointer, const std::vector<Event>&, const size_t, const size_t,
- const size_t, const Buffer<double>&, const double);
-template void FillVector<float2>(Queue&, const Device&, const std::shared_ptr<Program>, const Databases&,
+ const size_t, const Buffer<double>&, const double, const size_t);
+template void FillVector<float2>(Queue&, const Device&, const std::shared_ptr<Program>,
EventPointer, const std::vector<Event>&, const size_t, const size_t,
- const size_t, const Buffer<float2>&, const float2);
-template void FillVector<double2>(Queue&, const Device&, const std::shared_ptr<Program>, const Databases&,
+ const size_t, const Buffer<float2>&, const float2, const size_t);
+template void FillVector<double2>(Queue&, const Device&, const std::shared_ptr<Program>,
EventPointer, const std::vector<Event>&, const size_t, const size_t,
- const size_t, const Buffer<double2>&, const double2);
+ const size_t, const Buffer<double2>&, const double2, const size_t);
// =================================================================================================
} // namespace clblast
diff --git a/src/routines/common.hpp b/src/routines/common.hpp
index b909243d..c30a2e0e 100644
--- a/src/routines/common.hpp
+++ b/src/routines/common.hpp
@@ -36,20 +36,18 @@ void RunKernel(Kernel &kernel, Queue &queue, const Device &device,
// Sets all elements of a matrix to a constant value
template <typename T>
void FillMatrix(Queue &queue, const Device &device,
- const std::shared_ptr<Program> program, const Databases &,
+ const std::shared_ptr<Program> program,
EventPointer event, const std::vector<Event> &waitForEvents,
const size_t m, const size_t n, const size_t ld, const size_t offset,
- const Buffer<T> &dest,
- const T constant_value);
+ const Buffer<T> &dest, const T constant_value, const size_t local_size);
// Sets all elements of a vector to a constant value
template <typename T>
void FillVector(Queue &queue, const Device &device,
- const std::shared_ptr<Program> program, const Databases &,
+ const std::shared_ptr<Program> program,
EventPointer event, const std::vector<Event> &waitForEvents,
const size_t n, const size_t inc, const size_t offset,
- const Buffer<T> &dest,
- const T constant_value);
+ const Buffer<T> &dest, const T constant_value, const size_t local_size);
// =================================================================================================
diff --git a/src/routines/level2/xtrsv.cpp b/src/routines/level2/xtrsv.cpp
index 36c33a76..76401753 100644
--- a/src/routines/level2/xtrsv.cpp
+++ b/src/routines/level2/xtrsv.cpp
@@ -68,7 +68,7 @@ void Xtrsv<T>::Substitution(const Layout layout, const Triangle triangle,
// Launches the kernel
const auto local = std::vector<size_t>{db_["TRSV_BLOCK_SIZE"]};
- const auto global = std::vector<size_t>{1};
+ const auto global = std::vector<size_t>{Ceil(n, db_["TRSV_BLOCK_SIZE"])};
auto event = Event();
RunKernel(kernel, queue_, device_, global, local, event.pointer());
event.WaitForCompletion();
@@ -87,6 +87,11 @@ void Xtrsv<T>::DoTrsv(const Layout layout, const Triangle triangle,
// Makes sure all dimensions are larger than zero
if (n == 0) { throw BLASError(StatusCode::kInvalidDimension); }
+ // Some parts of this kernel are not tunable and thus require some minimal OpenCL properties
+ if (device_.MaxWorkGroupSize() < 16) { // minimum of total local work size of 16
+ throw RuntimeErrorCode(StatusCode::kNotImplemented);
+ }
+
// Tests the matrix and vector
TestMatrixA(n, n, a_buffer, a_offset, a_ld);
TestVectorX(n, b_buffer, b_offset, b_inc);
@@ -102,8 +107,8 @@ void Xtrsv<T>::DoTrsv(const Layout layout, const Triangle triangle,
// Fills the output buffer with zeros
auto eventWaitList = std::vector<Event>();
auto fill_vector_event = Event();
- FillVector(queue_, device_, program_, db_, fill_vector_event.pointer(), eventWaitList,
- n, x_inc, x_offset, x_buffer, ConstantZero<T>());
+ FillVector(queue_, device_, program_, fill_vector_event.pointer(), eventWaitList,
+ n, x_inc, x_offset, x_buffer, ConstantZero<T>(), 16);
fill_vector_event.WaitForCompletion();
// Derives properties based on the arguments
diff --git a/src/routines/level3/xgemm.hpp b/src/routines/level3/xgemm.hpp
index ec84fbb7..ed8cc69d 100644
--- a/src/routines/level3/xgemm.hpp
+++ b/src/routines/level3/xgemm.hpp
@@ -25,9 +25,9 @@ class Xgemm: public Routine {
public:
// Defines the assumptions of the GEMM kernels
- static const bool a_want_rotated_(const size_t gemm_kernel_id) { return gemm_kernel_id == 1; }
- static const bool b_want_rotated_(const size_t gemm_kernel_id) { return true; }
- static const bool c_want_rotated_(const size_t gemm_kernel_id) { return gemm_kernel_id == 1; }
+ static bool a_want_rotated_(const size_t gemm_kernel_id) { return gemm_kernel_id == 1; }
+ static bool b_want_rotated_(const size_t) { return true; }
+ static bool c_want_rotated_(const size_t gemm_kernel_id) { return gemm_kernel_id == 1; }
// Computes the size of the temporary GEMM buffer based on user-arguments
static size_t GetTempSize(const Layout layout, const Transpose a_transpose, const Transpose b_transpose,
diff --git a/src/routines/level3/xtrsm.cpp b/src/routines/level3/xtrsm.cpp
index d622e3bf..905660ff 100644
--- a/src/routines/level3/xtrsm.cpp
+++ b/src/routines/level3/xtrsm.cpp
@@ -78,6 +78,11 @@ void Xtrsm<T>::TrsmColMajor(const Side side, const Triangle triangle,
// Makes sure all dimensions are larger than zero
if ((m == 0) || (n == 0)) { throw BLASError(StatusCode::kInvalidDimension); }
+ // Some parts of this kernel are not tunable and thus require some minimal OpenCL properties
+ if (device_.MaxWorkGroupSize() < 16) { // minimum of total local work size of 16
+ throw RuntimeErrorCode(StatusCode::kNotImplemented);
+ }
+
// Computes the k dimension. This is based on whether or not matrix is A (on the left)
// or B (on the right) in the Xgemm routine.
const auto k = (side == Side::kLeft) ? m : n;
@@ -105,8 +110,8 @@ void Xtrsm<T>::TrsmColMajor(const Side side, const Triangle triangle,
// Fills the output buffer with zeros
auto eventWaitList = std::vector<Event>();
auto fill_matrix_event = Event();
- FillMatrix(queue_, device_, program_, db_, fill_matrix_event.pointer(), eventWaitList,
- x_one, x_two, x_ld, x_offset, x_buffer, ConstantZero<T>());
+ FillMatrix(queue_, device_, program_, fill_matrix_event.pointer(), eventWaitList,
+ x_one, x_two, x_ld, x_offset, x_buffer, ConstantZero<T>(), 16);
fill_matrix_event.WaitForCompletion();
// Inverts the diagonal blocks
diff --git a/src/routines/levelx/xinvert.cpp b/src/routines/levelx/xinvert.cpp
index a5ef9e10..eea8527a 100644
--- a/src/routines/levelx/xinvert.cpp
+++ b/src/routines/levelx/xinvert.cpp
@@ -49,9 +49,16 @@ void Xinvert<T>::InvertMatrixDiagonalBlocks(const Layout layout, const Triangle
throw BLASError(StatusCode::kInvalidDimension);
}
+ // Some parts of this kernel are not tunable and thus require some minimal OpenCL properties
+ if (device_.MaxWorkGroupSize() < 16) { // minimum of total local work size of 16
+ throw RuntimeErrorCode(StatusCode::kNotImplemented);
+ }
+
// Helper variables
const auto internal_block_size = static_cast<size_t>(db_["INTERNAL_BLOCK_SIZE"]);
- assert(internal_block_size == 16);
+ if (internal_block_size != 16) {
+ throw RuntimeErrorCode(StatusCode::kNotImplemented); // e.g. Apple CPU OpenCL with a WGS of 1
+ } // when barriers are present
const auto num_blocks = CeilDiv(n, block_size);
const auto num_internal_blocks = CeilDiv(n, internal_block_size);
const auto unit_diagonal = (diag == Diagonal::kUnit) ? true : false;
@@ -75,8 +82,9 @@ void Xinvert<T>::InvertMatrixDiagonalBlocks(const Layout layout, const Triangle
// Fills the output buffer with zeros
auto event_wait_list = std::vector<Event>();
auto fill_matrix_event = Event();
- FillMatrix(queue_, device_, program_, db_, fill_matrix_event.pointer(), event_wait_list,
- block_size, num_blocks * block_size, block_size, 0, dest, ConstantZero<T>());
+ FillMatrix(queue_, device_, program_, fill_matrix_event.pointer(), event_wait_list,
+ block_size, num_blocks * block_size, block_size, 0, dest, ConstantZero<T>(),
+ 16);
event_wait_list.push_back(fill_matrix_event);
// Inverts the diagonal IB by IB inner blocks of the matrix: one block per work-group
@@ -89,11 +97,11 @@ void Xinvert<T>::InvertMatrixDiagonalBlocks(const Layout layout, const Triangle
kernel.SetArgument(5, static_cast<int>(block_size));
kernel.SetArgument(6, static_cast<int>(unit_diagonal));
kernel.SetArgument(7, static_cast<int>(is_upper));
- const auto local = std::vector<size_t>{internal_block_size};
- const auto global = std::vector<size_t>{num_internal_blocks * internal_block_size};
+ const auto local_invert = std::vector<size_t>{internal_block_size};
+ const auto global_invert = std::vector<size_t>{num_internal_blocks * internal_block_size};
auto base_kernel_event = Event();
auto base_kernel_event_pointer = (internal_block_size == block_size) ? event_ : base_kernel_event.pointer();
- RunKernel(kernel, queue_, device_, global, local, base_kernel_event_pointer, event_wait_list);
+ RunKernel(kernel, queue_, device_, global_invert, local_invert, base_kernel_event_pointer, event_wait_list);
if (internal_block_size == block_size) { event_wait_list.push_back(base_kernel_event); }
// Builds up block_size x block_size blocks. For example, internal_block_size=16:
@@ -107,7 +115,8 @@ void Xinvert<T>::InvertMatrixDiagonalBlocks(const Layout layout, const Triangle
const auto npages = CeilDiv(n, current_size*2);
const auto local0 = (current_size <= 32) ? current_size/4 : 16;
const auto local = std::vector<size_t>{local0, 4};
- const auto global = std::vector<size_t>{(current_size/local[1]), npages*(current_size/16)*local[1]};
+ const auto global = std::vector<size_t>{Ceil(current_size/local[1], local[0]),
+ Ceil(npages*(current_size/16)*local[1], local[1])};
// Part 1
auto kernel1 = Kernel(program_, "TripleMatMul" + ToString(current_size) + "Part1" + name_postfix);
diff --git a/src/tuning/routines/xgemm.cpp b/src/tuning/routines/xgemm.cpp
index 92aab611..7d886ebf 100644
--- a/src/tuning/routines/xgemm.cpp
+++ b/src/tuning/routines/xgemm.cpp
@@ -25,14 +25,15 @@ namespace clblast {
// =================================================================================================
template <typename T>
-void RunGemmRoutine(const size_t value, const Queue& queue, const std::vector<Buffer<T>>& buffers) {
+void RunGemmRoutineMNK(const size_t m, const size_t n, const size_t k,
+ const Queue& queue, const std::vector<Buffer<T>>& buffers) {
auto queue_plain = queue();
auto event = cl_event{};
auto status = Gemm(Layout::kRowMajor, Transpose::kNo, Transpose::kNo,
- value, value, value, ConstantOne<T>(),
- buffers[0](), 0, value,
- buffers[1](), 0, value, ConstantOne<T>(),
- buffers[2](), 0, value,
+ m, n, k, ConstantOne<T>(),
+ buffers[0](), 0, k,
+ buffers[1](), 0, n, ConstantOne<T>(),
+ buffers[2](), 0, n,
&queue_plain, &event);
if (status != StatusCode::kSuccess) {
throw RuntimeError("Gemm failed with status " + ToString(status));
@@ -40,6 +41,10 @@ void RunGemmRoutine(const size_t value, const Queue& queue, const std::vector<Bu
clWaitForEvents(1, &event);
clReleaseEvent(event);
}
+template <typename T>
+void RunGemmRoutine(const size_t value, const Queue& queue, const std::vector<Buffer<T>>& buffers) {
+ RunGemmRoutineMNK(value, value, value, queue, buffers);
+}
template <typename T, size_t batch_count>
void RunGemmBatchedRoutine(const size_t value, const Queue& queue, const std::vector<Buffer<T>>& buffers) {
@@ -80,6 +85,55 @@ void RunGemmStridedBatchedRoutine(const size_t value, const Queue& queue, const
clWaitForEvents(1, &event);
clReleaseEvent(event);
}
+// =================================================================================================
+
+template <typename T>
+void TuneGemmSingleSize(const Platform& platform, const Device& device, const Context& context, Queue& queue,
+ const size_t m, const size_t n, const size_t k, const size_t num_runs) {
+
+ // Buffers
+ auto buffers = std::vector<Buffer<T>>{
+ Buffer<T>(context, m * k),
+ Buffer<T>(context, k * n),
+ Buffer<T>(context, m * n)
+ };
+ const auto FunctionToTune = [&]() { RunGemmRoutineMNK(m, n, k, queue, buffers); };
+
+ // Collects the timings for two methods
+ auto scores = std::vector<TuningResult>();
+ const auto methods = std::vector<std::string>{"in-direct", "direct"};
+ for (auto& method: methods) {
+
+ printf("* Testing the %s routine\n", method.c_str());
+ const auto limit = (method == "in-direct") ? 0 : std::max(std::max(m, n), k) + 1; // small or large number
+ ForceSelectIndirectFrom<T>(limit, device, "GemmRoutine", "XGEMM_MIN_INDIRECT_SIZE");
+ auto time_ms = -1.0;
+ try {
+ time_ms = TimeFunction(num_runs, FunctionToTune);
+ printf(" --> %9.2lf ms\n", time_ms);
+ }
+ catch (...) {
+ const auto status_code = DispatchExceptionCatchAll(true);
+ printf(" --> error %-5d\n", static_cast<int>(status_code));
+ }
+ auto tuning_results = Configuration();
+ tuning_results["XGEMM_MIN_INDIRECT_SIZE"] = limit;
+ tuning_results["PRECISION"] = static_cast<size_t>(PrecisionValue<T>());
+ scores.push_back(TuningResult{"gemm_kernel_selection_single_size", time_ms, tuning_results});
+ }
+
+ // Outputs the results as JSON to disk, including some meta-data
+ const auto precision_string = std::to_string(static_cast<size_t>(PrecisionValue<T>()));
+ auto metadata = std::vector<std::pair<std::string,std::string>>{
+ {"kernel_family", "gemm_routine_single_size"},
+ {"precision", precision_string},
+ {"arg_m", ToString(m)},
+ {"arg_n", ToString(n)},
+ {"arg_k", ToString(k)},
+ };
+ PrintTimingsToFileAsJSON("clblast_gemm_routine_single_size_" + precision_string + ".json",
+ device, platform, metadata, scores);
+}
// =================================================================================================
@@ -91,6 +145,9 @@ void TuneXgemm(int argc, char* argv[]) {
const auto device_id = GetArgument(command_line_args, help, kArgDevice, ConvertArgument(std::getenv("CLBLAST_DEVICE"), size_t{0}));
const auto precision = GetArgument(command_line_args, help, kArgPrecision, Precision::kSingle);
const auto num_runs = GetArgument(command_line_args, help, kArgNumRuns, size_t{10});
+ const auto arg_m = GetArgument(command_line_args, help, kArgM, -1); // optional
+ const auto arg_n = GetArgument(command_line_args, help, kArgN, -1); // optional
+ const auto arg_k = GetArgument(command_line_args, help, kArgK, -1); // optional
fprintf(stdout, "%s\n", help.c_str());
// OpenCL initialisation
@@ -119,16 +176,29 @@ void TuneXgemm(int argc, char* argv[]) {
}
}
- // Run the tuners for the XGEMM routines
- TuneKernelSelection<T>(platform, device, context, queue, precision, RunGemmRoutine<T>,
- 64, 2048, 64, 1, num_runs,
- "gemm", "GemmRoutine", "gemm_routine", "XGEMM_MIN_INDIRECT_SIZE");
- //TuneKernelSelection<T>(platform, device, context, queue, precision, RunGemmBatchedRoutine<T, 30>,
- // 16, 128, 32, 30, num_runs,
- // "gemmbatched", "GemmRoutine", "gemm_routine_2", "XGEMMBATCHED_MIN_INDIRECT_SIZE");
- //TuneKernelSelection<T>(platform, device, context, queue, precision, RunGemmStridedBatchedRoutine<T, 30>,
- // 16, 128, 32, 30, num_runs,
- // "gemmstridedbatched", "GemmRoutine", "gemm_routine_3", "XGEMMSTRIDEDBATCHED_MIN_INDIRECT_SIZE");
+ // Test for only one m/n/k size
+ if (arg_m != -1 || arg_n != -1 || arg_k != -1) {
+ printf("* Tuning for one specific size: m=%d, n=%d, k=%d\n", arg_m, arg_n, arg_k);
+ if (arg_m == -1 || arg_n == -1 || arg_k == -1) {
+ printf("* Error: If one of m/n/k specified, please specify all three\n");
+ return;
+ }
+ TuneGemmSingleSize<T>(platform, device, context, queue, static_cast<size_t>(arg_m),
+ static_cast<size_t>(arg_n), static_cast<size_t>(arg_k), num_runs);
+ }
+
+ else {
+ // Run the tuners for the XGEMM routines
+ TuneKernelSelection<T>(platform, device, context, queue, precision, RunGemmRoutine<T>,
+ 64, 2048, 64, 1, num_runs,
+ "gemm", "GemmRoutine", "gemm_routine", "XGEMM_MIN_INDIRECT_SIZE");
+ //TuneKernelSelection<T>(platform, device, context, queue, precision, RunGemmBatchedRoutine<T, 30>,
+ // 16, 128, 32, 30, num_runs,
+ // "gemmbatched", "GemmRoutine", "gemm_routine_2", "XGEMMBATCHED_MIN_INDIRECT_SIZE");
+ //TuneKernelSelection<T>(platform, device, context, queue, precision, RunGemmStridedBatchedRoutine<T, 30>,
+ // 16, 128, 32, 30, num_runs,
+ // "gemmstridedbatched", "GemmRoutine", "gemm_routine_3", "XGEMMSTRIDEDBATCHED_MIN_INDIRECT_SIZE");
+ }
printf("* Completed tuning process\n");
printf("\n");
diff --git a/src/utilities/utilities.hpp b/src/utilities/utilities.hpp
index a29e531a..16a241af 100644
--- a/src/utilities/utilities.hpp
+++ b/src/utilities/utilities.hpp
@@ -122,6 +122,7 @@ constexpr auto kArgHelp = "h";
constexpr auto kArgQuiet = "q";
constexpr auto kArgNoAbbreviations = "no_abbrv";
constexpr auto kArgNumRuns = "runs";
+constexpr auto kArgFullStatistics = "full_statistics";
// The buffer names
constexpr auto kBufVecX = "X";
@@ -245,6 +246,7 @@ struct Arguments {
size_t num_steps = 0;
size_t num_runs = 10;
std::vector<std::string> tuner_files = {};
+ bool full_statistics = false;
#ifdef CLBLAST_REF_CUBLAS
void* cublas_handle; // cublasHandle_t
#endif
generated by cgit v1.2.3 (git 2.39.1) at 2024-07-07 08:37:56 +0000