summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--CHANGELOG1
-rw-r--r--src/database/kernels/xgemv.hpp2
-rw-r--r--src/kernels/level2/xgemv_fast.opencl143
-rw-r--r--src/routines/level2/xgemv.cpp2
-rw-r--r--src/tuning/kernels/xgemv.cpp38
5 files changed, 107 insertions, 79 deletions
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/src/database/kernels/xgemv.hpp b/src/database/kernels/xgemv.hpp
index 65f4b5c8..3aa1863f 100644
--- a/src/database/kernels/xgemv.hpp
+++ b/src/database/kernels/xgemv.hpp
@@ -36,7 +36,7 @@ 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} } },
+ { "AMD Radeon R9 M370X Compute Engine", { {"WGS1",128}, {"WPT1",1}, {"VW2",1}, {"WGS2",128}, {"WPT2",1}, {"VW3",8}, {"WGS3",16}, {"WPT3",16} } },
{ "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} } },
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<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);
}
// =================================================================================================
diff --git a/src/routines/level2/xgemv.cpp b/src/routines/level2/xgemv.cpp
index 2842ef07..e4d407c8 100644
--- a/src/routines/level2/xgemv.cpp
+++ b/src/routines/level2/xgemv.cpp
@@ -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"];
}
diff --git a/src/tuning/kernels/xgemv.cpp b/src/tuning/kernels/xgemv.cpp
index 5c187d33..96d4a5f2 100644
--- a/src/tuning/kernels/xgemv.cpp
+++ b/src/tuning/kernels/xgemv.cpp
@@ -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,
generated by cgit v1.2.3 (git 2.39.1) at 2024-08-02 02:05:18 +0000