diff options
-rw-r--r-- | src/kernel_preprocessor.cpp | 43 | ||||
-rw-r--r-- | src/kernels/level3/transpose_fast.opencl | 66 | ||||
-rw-r--r-- | src/kernels/level3/xgemm_direct_part1.opencl | 56 | ||||
-rw-r--r-- | src/kernels/level3/xgemm_direct_part3.opencl | 44 | ||||
-rw-r--r-- | src/kernels/level3/xgemm_part3.opencl | 10 | ||||
-rw-r--r-- | test/correctness/misc/preprocessor.cpp | 3 |
6 files changed, 117 insertions, 105 deletions
diff --git a/src/kernel_preprocessor.cpp b/src/kernel_preprocessor.cpp index 493c009c..8738a837 100644 --- a/src/kernel_preprocessor.cpp +++ b/src/kernel_preprocessor.cpp @@ -27,6 +27,7 @@ #include <sstream> #include <algorithm> #include <unordered_map> +#include <map> #include <vector> #include "kernel_preprocessor.hpp" @@ -34,6 +35,17 @@ namespace clblast { // ================================================================================================= +struct compare_longer_string { + bool operator() (const std::string &lhs, const std::string &rhs) const { + if (lhs.size() > rhs.size()) { return true; } + if (lhs.size() < rhs.size()) { return false; } + return lhs < rhs; + } +}; + +using DefinesIntMap = std::map<std::string, int, compare_longer_string>; +using DefinesStringMap = std::map<std::string, std::string, std::greater<std::string>>; + void RaiseError(const std::string& source_line, const std::string& exception_message) { printf("[OpenCL pre-processor] Error in source line: %s\n", source_line.c_str()); throw Error<std::runtime_error>(exception_message); @@ -122,7 +134,7 @@ void FindReplace(std::string &subject, const std::string &search, const std::str } } -void SubstituteDefines(const std::unordered_map<std::string, int>& defines, +void SubstituteDefines(const DefinesIntMap& defines, std::string& source_string) { for (const auto &define : defines) { FindReplace(source_string, define.first, std::to_string(define.second)); @@ -130,8 +142,8 @@ void SubstituteDefines(const std::unordered_map<std::string, int>& defines, } bool EvaluateCondition(std::string condition, - const std::unordered_map<std::string, int> &defines, - const std::unordered_map<std::string, std::string> &defines_string) { + const DefinesIntMap &defines, + const DefinesStringMap &defines_string) { // Replace macros in the string SubstituteDefines(defines, condition); @@ -177,7 +189,7 @@ bool EvaluateCondition(std::string condition, // ================================================================================================= // Array to register promotion, e.g. arr[w] to {arr_0, arr_1} -void ArrayToRegister(std::string &source_line, const std::unordered_map<std::string, int>& defines, +void ArrayToRegister(std::string &source_line, const DefinesIntMap& defines, const std::unordered_map<std::string, size_t>& arrays_to_registers, const size_t num_brackets) { @@ -265,9 +277,9 @@ void ArrayToRegister(std::string &source_line, const std::unordered_map<std::str // First pass: detect defines and comments std::vector<std::string> PreprocessDefinesAndComments(const std::string& source, - std::unordered_map<std::string, int>& defines_int) { + DefinesIntMap& defines_int) { auto lines = std::vector<std::string>(); - auto defines_string = std::unordered_map<std::string, std::string>(); + auto defines_string = DefinesStringMap(); // Parse the input string into a vector of lines const auto max_depth_defines = 30; @@ -276,17 +288,16 @@ std::vector<std::string> PreprocessDefinesAndComments(const std::string& source, auto source_stream = std::stringstream(source); auto line = std::string{""}; while (std::getline(source_stream, line)) { - //printf("[@%d] disabled=%d '%s'\n", depth, disabled[depth], line.c_str()); + //printf("[@%zu] disabled=%d '%s'\n", depth, disabled[depth], line.c_str()); // Decide whether or not to remain in 'disabled' mode + // {0 => enabled, 1 => disabled, but could become enabled again later, 2 => disabled until #endif if (line.find("#endif") != std::string::npos) { disabled[depth] = 0; } - if (line.find("#elif") != std::string::npos) { - disabled[depth] = 0; - } - if (line.find("#else") != std::string::npos) { - disabled[depth] = (disabled[depth] == 0) ? 1 : 0; + if (line.find("#elif") != std::string::npos || line.find("#else") != std::string::npos) { + if (disabled[depth] == 0) { disabled[depth] = 2; } // was enabled, now disabled until #endif + if (disabled[depth] == 1) { disabled[depth] = 0; } // was disabled, now potentially enabled again } // Measures the depth of pre-processor defines @@ -304,7 +315,7 @@ std::vector<std::string> PreprocessDefinesAndComments(const std::string& source, // Verifies whether this level or any level below is disabled auto is_disabled = false; for (auto d = size_t{0}; d <= depth; ++d) { - if (disabled[d] == 1) { is_disabled = true; } + if (disabled[d] >= 1) { is_disabled = true; } } // Not in a disabled-block @@ -382,7 +393,7 @@ std::vector<std::string> PreprocessDefinesAndComments(const std::string& source, // Second pass: detect array-to-register promotion pragma's and replace declarations & function calls std::vector<std::string> PreprocessUnrollLoops(const std::vector<std::string>& source_lines, - const std::unordered_map<std::string, int>& defines, + const DefinesIntMap& defines, std::unordered_map<std::string, size_t>& arrays_to_registers) { auto lines = std::vector<std::string>(); @@ -435,7 +446,7 @@ std::vector<std::string> PreprocessUnrollLoops(const std::vector<std::string>& s // Third pass: unroll loops and perform actual array-to-register promotion std::vector<std::string> PreprocessUnrollLoops(const std::vector<std::string>& source_lines, - const std::unordered_map<std::string, int>& defines, + const DefinesIntMap& defines, std::unordered_map<std::string, size_t>& arrays_to_registers, const bool array_to_register_promotion) { auto lines = std::vector<std::string>(); @@ -538,7 +549,7 @@ std::vector<std::string> PreprocessUnrollLoops(const std::vector<std::string>& s std::string PreprocessKernelSource(const std::string& kernel_source) { // Retrieves the defines and removes comments from the source lines - auto defines = std::unordered_map<std::string, int>(); + auto defines = DefinesIntMap(); auto lines = PreprocessDefinesAndComments(kernel_source, defines); // Unrolls loops (single level each call) diff --git a/src/kernels/level3/transpose_fast.opencl b/src/kernels/level3/transpose_fast.opencl index 8fa7405c..1b9fca45 100644 --- a/src/kernels/level3/transpose_fast.opencl +++ b/src/kernels/level3/transpose_fast.opencl @@ -74,51 +74,51 @@ void TransposeMatrixFast(const int ld, // Loads transposed data from the local memory #pragma promote_to_registers - realT v[TRA_WPT]; + realT vpm[TRA_WPT]; #pragma unroll for (int _w_one = 0; _w_one < TRA_WPT; _w_one += 1) { - v[_w_one] = tile[get_local_id(1)*TRA_WPT + _w_one][get_local_id(0)]; + vpm[_w_one] = tile[get_local_id(1)*TRA_WPT + _w_one][get_local_id(0)]; } // Performs the register-level transpose of the vectorized data #pragma promote_to_registers realT results[TRA_WPT]; #if TRA_WPT == 1 - results[0] = v[0]; + results[0] = vpm[0]; #elif TRA_WPT == 2 - results[0].x = v[0].x; results[0].y = v[1].x; - results[1].x = v[0].y; results[1].y = v[1].y; + results[0].x = vpm[0].x; results[0].y = vpm[1].x; + results[1].x = vpm[0].y; results[1].y = vpm[1].y; #elif TRA_WPT == 4 - results[0].x = v[0].x; results[0].y = v[1].x; results[0].z = v[2].x; results[0].w = v[3].x; - results[1].x = v[0].y; results[1].y = v[1].y; results[1].z = v[2].y; results[1].w = v[3].y; - results[2].x = v[0].z; results[2].y = v[1].z; results[2].z = v[2].z; results[2].w = v[3].z; - results[3].x = v[0].w; results[3].y = v[1].w; results[3].z = v[2].w; results[3].w = v[3].w; + results[0].x = vpm[0].x; results[0].y = vpm[1].x; results[0].z = vpm[2].x; results[0].w = vpm[3].x; + results[1].x = vpm[0].y; results[1].y = vpm[1].y; results[1].z = vpm[2].y; results[1].w = vpm[3].y; + results[2].x = vpm[0].z; results[2].y = vpm[1].z; results[2].z = vpm[2].z; results[2].w = vpm[3].z; + results[3].x = vpm[0].w; results[3].y = vpm[1].w; results[3].z = vpm[2].w; results[3].w = vpm[3].w; #elif TRA_WPT == 8 - results[0].s0 = v[0].s0; results[0].s1 = v[1].s0; results[0].s2 = v[2].s0; results[0].s3 = v[3].s0; results[0].s4 = v[4].s0; results[0].s5 = v[5].s0; results[0].s6 = v[6].s0; results[0].s7 = v[7].s0; - results[1].s0 = v[0].s1; results[1].s1 = v[1].s1; results[1].s2 = v[2].s1; results[1].s3 = v[3].s1; results[1].s4 = v[4].s1; results[1].s5 = v[5].s1; results[1].s6 = v[6].s1; results[1].s7 = v[7].s1; - results[2].s0 = v[0].s2; results[2].s1 = v[1].s2; results[2].s2 = v[2].s2; results[2].s3 = v[3].s2; results[2].s4 = v[4].s2; results[2].s5 = v[5].s2; results[2].s6 = v[6].s2; results[2].s7 = v[7].s2; - results[3].s0 = v[0].s3; results[3].s1 = v[1].s3; results[3].s2 = v[2].s3; results[3].s3 = v[3].s3; results[3].s4 = v[4].s3; results[3].s5 = v[5].s3; results[3].s6 = v[6].s3; results[3].s7 = v[7].s3; - results[4].s0 = v[0].s4; results[4].s1 = v[1].s4; results[4].s2 = v[2].s4; results[4].s3 = v[3].s4; results[4].s4 = v[4].s4; results[4].s5 = v[5].s4; results[4].s6 = v[6].s4; results[4].s7 = v[7].s4; - results[5].s0 = v[0].s5; results[5].s1 = v[1].s5; results[5].s2 = v[2].s5; results[5].s3 = v[3].s5; results[5].s4 = v[4].s5; results[5].s5 = v[5].s5; results[5].s6 = v[6].s5; results[5].s7 = v[7].s5; - results[6].s0 = v[0].s6; results[6].s1 = v[1].s6; results[6].s2 = v[2].s6; results[6].s3 = v[3].s6; results[6].s4 = v[4].s6; results[6].s5 = v[5].s6; results[6].s6 = v[6].s6; results[6].s7 = v[7].s6; - results[7].s0 = v[0].s7; results[7].s1 = v[1].s7; results[7].s2 = v[2].s7; results[7].s3 = v[3].s7; results[7].s4 = v[4].s7; results[7].s5 = v[5].s7; results[7].s6 = v[6].s7; results[7].s7 = v[7].s7; + results[0].s0 = vpm[0].s0; results[0].s1 = vpm[1].s0; results[0].s2 = vpm[2].s0; results[0].s3 = vpm[3].s0; results[0].s4 = vpm[4].s0; results[0].s5 = vpm[5].s0; results[0].s6 = vpm[6].s0; results[0].s7 = vpm[7].s0; + results[1].s0 = vpm[0].s1; results[1].s1 = vpm[1].s1; results[1].s2 = vpm[2].s1; results[1].s3 = vpm[3].s1; results[1].s4 = vpm[4].s1; results[1].s5 = vpm[5].s1; results[1].s6 = vpm[6].s1; results[1].s7 = vpm[7].s1; + results[2].s0 = vpm[0].s2; results[2].s1 = vpm[1].s2; results[2].s2 = vpm[2].s2; results[2].s3 = vpm[3].s2; results[2].s4 = vpm[4].s2; results[2].s5 = vpm[5].s2; results[2].s6 = vpm[6].s2; results[2].s7 = vpm[7].s2; + results[3].s0 = vpm[0].s3; results[3].s1 = vpm[1].s3; results[3].s2 = vpm[2].s3; results[3].s3 = vpm[3].s3; results[3].s4 = vpm[4].s3; results[3].s5 = vpm[5].s3; results[3].s6 = vpm[6].s3; results[3].s7 = vpm[7].s3; + results[4].s0 = vpm[0].s4; results[4].s1 = vpm[1].s4; results[4].s2 = vpm[2].s4; results[4].s3 = vpm[3].s4; results[4].s4 = vpm[4].s4; results[4].s5 = vpm[5].s4; results[4].s6 = vpm[6].s4; results[4].s7 = vpm[7].s4; + results[5].s0 = vpm[0].s5; results[5].s1 = vpm[1].s5; results[5].s2 = vpm[2].s5; results[5].s3 = vpm[3].s5; results[5].s4 = vpm[4].s5; results[5].s5 = vpm[5].s5; results[5].s6 = vpm[6].s5; results[5].s7 = vpm[7].s5; + results[6].s0 = vpm[0].s6; results[6].s1 = vpm[1].s6; results[6].s2 = vpm[2].s6; results[6].s3 = vpm[3].s6; results[6].s4 = vpm[4].s6; results[6].s5 = vpm[5].s6; results[6].s6 = vpm[6].s6; results[6].s7 = vpm[7].s6; + results[7].s0 = vpm[0].s7; results[7].s1 = vpm[1].s7; results[7].s2 = vpm[2].s7; results[7].s3 = vpm[3].s7; results[7].s4 = vpm[4].s7; results[7].s5 = vpm[5].s7; results[7].s6 = vpm[6].s7; results[7].s7 = vpm[7].s7; #elif TRA_WPT == 16 - results[ 0].s0 = v[0].s0; results[ 0].s1 = v[1].s0; results[ 0].s2 = v[2].s0; results[ 0].s3 = v[3].s0; results[ 0].s4 = v[4].s0; results[ 0].s5 = v[5].s0; results[ 0].s6 = v[6].s0; results[ 0].s7 = v[7].s0; results[ 0].s8 = v[8].s0; results[ 0].s9 = v[9].s0; results[ 0].sA = v[10].s0; results[ 0].sB = v[11].s0; results[ 0].sC = v[12].s0; results[ 0].sD = v[13].s0; results[ 0].sE = v[14].s0; results[ 0].sF = v[15].s0; - results[ 1].s0 = v[0].s1; results[ 1].s1 = v[1].s1; results[ 1].s2 = v[2].s1; results[ 1].s3 = v[3].s1; results[ 1].s4 = v[4].s1; results[ 1].s5 = v[5].s1; results[ 1].s6 = v[6].s1; results[ 1].s7 = v[7].s1; results[ 1].s8 = v[8].s1; results[ 1].s9 = v[9].s1; results[ 1].sA = v[10].s1; results[ 1].sB = v[11].s1; results[ 1].sC = v[12].s1; results[ 1].sD = v[13].s1; results[ 1].sE = v[14].s1; results[ 1].sF = v[15].s1; - results[ 2].s0 = v[0].s2; results[ 2].s1 = v[1].s2; results[ 2].s2 = v[2].s2; results[ 2].s3 = v[3].s2; results[ 2].s4 = v[4].s2; results[ 2].s5 = v[5].s2; results[ 2].s6 = v[6].s2; results[ 2].s7 = v[7].s2; results[ 2].s8 = v[8].s2; results[ 2].s9 = v[9].s2; results[ 2].sA = v[10].s2; results[ 2].sB = v[11].s2; results[ 2].sC = v[12].s2; results[ 2].sD = v[13].s2; results[ 2].sE = v[14].s2; results[ 2].sF = v[15].s2; - results[ 3].s0 = v[0].s3; results[ 3].s1 = v[1].s3; results[ 3].s2 = v[2].s3; results[ 3].s3 = v[3].s3; results[ 3].s4 = v[4].s3; results[ 3].s5 = v[5].s3; results[ 3].s6 = v[6].s3; results[ 3].s7 = v[7].s3; results[ 3].s8 = v[8].s3; results[ 3].s9 = v[9].s3; results[ 3].sA = v[10].s3; results[ 3].sB = v[11].s3; results[ 3].sC = v[12].s3; results[ 3].sD = v[13].s3; results[ 3].sE = v[14].s3; results[ 3].sF = v[15].s3; - results[ 4].s0 = v[0].s4; results[ 4].s1 = v[1].s4; results[ 4].s2 = v[2].s4; results[ 4].s3 = v[3].s4; results[ 4].s4 = v[4].s4; results[ 4].s5 = v[5].s4; results[ 4].s6 = v[6].s4; results[ 4].s7 = v[7].s4; results[ 4].s8 = v[8].s4; results[ 4].s9 = v[9].s4; results[ 4].sA = v[10].s4; results[ 4].sB = v[11].s4; results[ 4].sC = v[12].s4; results[ 4].sD = v[13].s4; results[ 4].sE = v[14].s4; results[ 4].sF = v[15].s4; - results[ 5].s0 = v[0].s5; results[ 5].s1 = v[1].s5; results[ 5].s2 = v[2].s5; results[ 5].s3 = v[3].s5; results[ 5].s4 = v[4].s5; results[ 5].s5 = v[5].s5; results[ 5].s6 = v[6].s5; results[ 5].s7 = v[7].s5; results[ 5].s8 = v[8].s5; results[ 5].s9 = v[9].s5; results[ 5].sA = v[10].s5; results[ 5].sB = v[11].s5; results[ 5].sC = v[12].s5; results[ 5].sD = v[13].s5; results[ 5].sE = v[14].s5; results[ 5].sF = v[15].s5; - results[ 6].s0 = v[0].s6; results[ 6].s1 = v[1].s6; results[ 6].s2 = v[2].s6; results[ 6].s3 = v[3].s6; results[ 6].s4 = v[4].s6; results[ 6].s5 = v[5].s6; results[ 6].s6 = v[6].s6; results[ 6].s7 = v[7].s6; results[ 6].s8 = v[8].s6; results[ 6].s9 = v[9].s6; results[ 6].sA = v[10].s6; results[ 6].sB = v[11].s6; results[ 6].sC = v[12].s6; results[ 6].sD = v[13].s6; results[ 6].sE = v[14].s6; results[ 6].sF = v[15].s6; - results[ 7].s0 = v[0].s7; results[ 7].s1 = v[1].s7; results[ 7].s2 = v[2].s7; results[ 7].s3 = v[3].s7; results[ 7].s4 = v[4].s7; results[ 7].s5 = v[5].s7; results[ 7].s6 = v[6].s7; results[ 7].s7 = v[7].s7; results[ 7].s8 = v[8].s7; results[ 7].s9 = v[9].s7; results[ 7].sA = v[10].s7; results[ 7].sB = v[11].s7; results[ 7].sC = v[12].s7; results[ 7].sD = v[13].s7; results[ 7].sE = v[14].s7; results[ 7].sF = v[15].s7; - results[ 8].s0 = v[0].s8; results[ 8].s1 = v[1].s8; results[ 8].s2 = v[2].s8; results[ 8].s3 = v[3].s8; results[ 8].s4 = v[4].s8; results[ 8].s5 = v[5].s8; results[ 8].s6 = v[6].s8; results[ 8].s7 = v[7].s8; results[ 8].s8 = v[8].s8; results[ 8].s9 = v[9].s8; results[ 8].sA = v[10].s8; results[ 8].sB = v[11].s8; results[ 8].sC = v[12].s8; results[ 8].sD = v[13].s8; results[ 8].sE = v[14].s8; results[ 8].sF = v[15].s8; - results[ 9].s0 = v[0].s9; results[ 9].s1 = v[1].s9; results[ 9].s2 = v[2].s9; results[ 9].s3 = v[3].s9; results[ 9].s4 = v[4].s9; results[ 9].s5 = v[5].s9; results[ 9].s6 = v[6].s9; results[ 9].s7 = v[7].s9; results[ 9].s8 = v[8].s9; results[ 9].s9 = v[9].s9; results[ 9].sA = v[10].s9; results[ 9].sB = v[11].s9; results[ 9].sC = v[12].s9; results[ 9].sD = v[13].s9; results[ 9].sE = v[14].s9; results[ 9].sF = v[15].s9; - results[10].s0 = v[0].sA; results[10].s1 = v[1].sA; results[10].s2 = v[2].sA; results[10].s3 = v[3].sA; results[10].s4 = v[4].sA; results[10].s5 = v[5].sA; results[10].s6 = v[6].sA; results[10].s7 = v[7].sA; results[10].s8 = v[8].sA; results[10].s9 = v[9].sA; results[10].sA = v[10].sA; results[10].sB = v[11].sA; results[10].sC = v[12].sA; results[10].sD = v[13].sA; results[10].sE = v[14].sA; results[10].sF = v[15].sA; - results[11].s0 = v[0].sB; results[11].s1 = v[1].sB; results[11].s2 = v[2].sB; results[11].s3 = v[3].sB; results[11].s4 = v[4].sB; results[11].s5 = v[5].sB; results[11].s6 = v[6].sB; results[11].s7 = v[7].sB; results[11].s8 = v[8].sB; results[11].s9 = v[9].sB; results[11].sA = v[10].sB; results[11].sB = v[11].sB; results[11].sC = v[12].sB; results[11].sD = v[13].sB; results[11].sE = v[14].sB; results[11].sF = v[15].sB; - results[12].s0 = v[0].sC; results[12].s1 = v[1].sC; results[12].s2 = v[2].sC; results[12].s3 = v[3].sC; results[12].s4 = v[4].sC; results[12].s5 = v[5].sC; results[12].s6 = v[6].sC; results[12].s7 = v[7].sC; results[12].s8 = v[8].sC; results[12].s9 = v[9].sC; results[12].sA = v[10].sC; results[12].sB = v[11].sC; results[12].sC = v[12].sC; results[12].sD = v[13].sC; results[12].sE = v[14].sC; results[12].sF = v[15].sC; - results[13].s0 = v[0].sD; results[13].s1 = v[1].sD; results[13].s2 = v[2].sD; results[13].s3 = v[3].sD; results[13].s4 = v[4].sD; results[13].s5 = v[5].sD; results[13].s6 = v[6].sD; results[13].s7 = v[7].sD; results[13].s8 = v[8].sD; results[13].s9 = v[9].sD; results[13].sA = v[10].sD; results[13].sB = v[11].sD; results[13].sC = v[12].sD; results[13].sD = v[13].sD; results[13].sE = v[14].sD; results[13].sF = v[15].sD; - results[14].s0 = v[0].sE; results[14].s1 = v[1].sE; results[14].s2 = v[2].sE; results[14].s3 = v[3].sE; results[14].s4 = v[4].sE; results[14].s5 = v[5].sE; results[14].s6 = v[6].sE; results[14].s7 = v[7].sE; results[14].s8 = v[8].sE; results[14].s9 = v[9].sE; results[14].sA = v[10].sE; results[14].sB = v[11].sE; results[14].sC = v[12].sE; results[14].sD = v[13].sE; results[14].sE = v[14].sE; results[14].sF = v[15].sE; - results[15].s0 = v[0].sF; results[15].s1 = v[1].sF; results[15].s2 = v[2].sF; results[15].s3 = v[3].sF; results[15].s4 = v[4].sF; results[15].s5 = v[5].sF; results[15].s6 = v[6].sF; results[15].s7 = v[7].sF; results[15].s8 = v[8].sF; results[15].s9 = v[9].sF; results[15].sA = v[10].sF; results[15].sB = v[11].sF; results[15].sC = v[12].sF; results[15].sD = v[13].sF; results[15].sE = v[14].sF; results[15].sF = v[15].sF; + results[ 0].s0 = vpm[0].s0; results[ 0].s1 = vpm[1].s0; results[ 0].s2 = vpm[2].s0; results[ 0].s3 = vpm[3].s0; results[ 0].s4 = vpm[4].s0; results[ 0].s5 = vpm[5].s0; results[ 0].s6 = vpm[6].s0; results[ 0].s7 = vpm[7].s0; results[ 0].s8 = vpm[8].s0; results[ 0].s9 = vpm[9].s0; results[ 0].sA = vpm[10].s0; results[ 0].sB = vpm[11].s0; results[ 0].sC = vpm[12].s0; results[ 0].sD = vpm[13].s0; results[ 0].sE = vpm[14].s0; results[ 0].sF = vpm[15].s0; + results[ 1].s0 = vpm[0].s1; results[ 1].s1 = vpm[1].s1; results[ 1].s2 = vpm[2].s1; results[ 1].s3 = vpm[3].s1; results[ 1].s4 = vpm[4].s1; results[ 1].s5 = vpm[5].s1; results[ 1].s6 = vpm[6].s1; results[ 1].s7 = vpm[7].s1; results[ 1].s8 = vpm[8].s1; results[ 1].s9 = vpm[9].s1; results[ 1].sA = vpm[10].s1; results[ 1].sB = vpm[11].s1; results[ 1].sC = vpm[12].s1; results[ 1].sD = vpm[13].s1; results[ 1].sE = vpm[14].s1; results[ 1].sF = vpm[15].s1; + results[ 2].s0 = vpm[0].s2; results[ 2].s1 = vpm[1].s2; results[ 2].s2 = vpm[2].s2; results[ 2].s3 = vpm[3].s2; results[ 2].s4 = vpm[4].s2; results[ 2].s5 = vpm[5].s2; results[ 2].s6 = vpm[6].s2; results[ 2].s7 = vpm[7].s2; results[ 2].s8 = vpm[8].s2; results[ 2].s9 = vpm[9].s2; results[ 2].sA = vpm[10].s2; results[ 2].sB = vpm[11].s2; results[ 2].sC = vpm[12].s2; results[ 2].sD = vpm[13].s2; results[ 2].sE = vpm[14].s2; results[ 2].sF = vpm[15].s2; + results[ 3].s0 = vpm[0].s3; results[ 3].s1 = vpm[1].s3; results[ 3].s2 = vpm[2].s3; results[ 3].s3 = vpm[3].s3; results[ 3].s4 = vpm[4].s3; results[ 3].s5 = vpm[5].s3; results[ 3].s6 = vpm[6].s3; results[ 3].s7 = vpm[7].s3; results[ 3].s8 = vpm[8].s3; results[ 3].s9 = vpm[9].s3; results[ 3].sA = vpm[10].s3; results[ 3].sB = vpm[11].s3; results[ 3].sC = vpm[12].s3; results[ 3].sD = vpm[13].s3; results[ 3].sE = vpm[14].s3; results[ 3].sF = vpm[15].s3; + results[ 4].s0 = vpm[0].s4; results[ 4].s1 = vpm[1].s4; results[ 4].s2 = vpm[2].s4; results[ 4].s3 = vpm[3].s4; results[ 4].s4 = vpm[4].s4; results[ 4].s5 = vpm[5].s4; results[ 4].s6 = vpm[6].s4; results[ 4].s7 = vpm[7].s4; results[ 4].s8 = vpm[8].s4; results[ 4].s9 = vpm[9].s4; results[ 4].sA = vpm[10].s4; results[ 4].sB = vpm[11].s4; results[ 4].sC = vpm[12].s4; results[ 4].sD = vpm[13].s4; results[ 4].sE = vpm[14].s4; results[ 4].sF = vpm[15].s4; + results[ 5].s0 = vpm[0].s5; results[ 5].s1 = vpm[1].s5; results[ 5].s2 = vpm[2].s5; results[ 5].s3 = vpm[3].s5; results[ 5].s4 = vpm[4].s5; results[ 5].s5 = vpm[5].s5; results[ 5].s6 = vpm[6].s5; results[ 5].s7 = vpm[7].s5; results[ 5].s8 = vpm[8].s5; results[ 5].s9 = vpm[9].s5; results[ 5].sA = vpm[10].s5; results[ 5].sB = vpm[11].s5; results[ 5].sC = vpm[12].s5; results[ 5].sD = vpm[13].s5; results[ 5].sE = vpm[14].s5; results[ 5].sF = vpm[15].s5; + results[ 6].s0 = vpm[0].s6; results[ 6].s1 = vpm[1].s6; results[ 6].s2 = vpm[2].s6; results[ 6].s3 = vpm[3].s6; results[ 6].s4 = vpm[4].s6; results[ 6].s5 = vpm[5].s6; results[ 6].s6 = vpm[6].s6; results[ 6].s7 = vpm[7].s6; results[ 6].s8 = vpm[8].s6; results[ 6].s9 = vpm[9].s6; results[ 6].sA = vpm[10].s6; results[ 6].sB = vpm[11].s6; results[ 6].sC = vpm[12].s6; results[ 6].sD = vpm[13].s6; results[ 6].sE = vpm[14].s6; results[ 6].sF = vpm[15].s6; + results[ 7].s0 = vpm[0].s7; results[ 7].s1 = vpm[1].s7; results[ 7].s2 = vpm[2].s7; results[ 7].s3 = vpm[3].s7; results[ 7].s4 = vpm[4].s7; results[ 7].s5 = vpm[5].s7; results[ 7].s6 = vpm[6].s7; results[ 7].s7 = vpm[7].s7; results[ 7].s8 = vpm[8].s7; results[ 7].s9 = vpm[9].s7; results[ 7].sA = vpm[10].s7; results[ 7].sB = vpm[11].s7; results[ 7].sC = vpm[12].s7; results[ 7].sD = vpm[13].s7; results[ 7].sE = vpm[14].s7; results[ 7].sF = vpm[15].s7; + results[ 8].s0 = vpm[0].s8; results[ 8].s1 = vpm[1].s8; results[ 8].s2 = vpm[2].s8; results[ 8].s3 = vpm[3].s8; results[ 8].s4 = vpm[4].s8; results[ 8].s5 = vpm[5].s8; results[ 8].s6 = vpm[6].s8; results[ 8].s7 = vpm[7].s8; results[ 8].s8 = vpm[8].s8; results[ 8].s9 = vpm[9].s8; results[ 8].sA = vpm[10].s8; results[ 8].sB = vpm[11].s8; results[ 8].sC = vpm[12].s8; results[ 8].sD = vpm[13].s8; results[ 8].sE = vpm[14].s8; results[ 8].sF = vpm[15].s8; + results[ 9].s0 = vpm[0].s9; results[ 9].s1 = vpm[1].s9; results[ 9].s2 = vpm[2].s9; results[ 9].s3 = vpm[3].s9; results[ 9].s4 = vpm[4].s9; results[ 9].s5 = vpm[5].s9; results[ 9].s6 = vpm[6].s9; results[ 9].s7 = vpm[7].s9; results[ 9].s8 = vpm[8].s9; results[ 9].s9 = vpm[9].s9; results[ 9].sA = vpm[10].s9; results[ 9].sB = vpm[11].s9; results[ 9].sC = vpm[12].s9; results[ 9].sD = vpm[13].s9; results[ 9].sE = vpm[14].s9; results[ 9].sF = vpm[15].s9; + results[10].s0 = vpm[0].sA; results[10].s1 = vpm[1].sA; results[10].s2 = vpm[2].sA; results[10].s3 = vpm[3].sA; results[10].s4 = vpm[4].sA; results[10].s5 = vpm[5].sA; results[10].s6 = vpm[6].sA; results[10].s7 = vpm[7].sA; results[10].s8 = vpm[8].sA; results[10].s9 = vpm[9].sA; results[10].sA = vpm[10].sA; results[10].sB = vpm[11].sA; results[10].sC = vpm[12].sA; results[10].sD = vpm[13].sA; results[10].sE = vpm[14].sA; results[10].sF = vpm[15].sA; + results[11].s0 = vpm[0].sB; results[11].s1 = vpm[1].sB; results[11].s2 = vpm[2].sB; results[11].s3 = vpm[3].sB; results[11].s4 = vpm[4].sB; results[11].s5 = vpm[5].sB; results[11].s6 = vpm[6].sB; results[11].s7 = vpm[7].sB; results[11].s8 = vpm[8].sB; results[11].s9 = vpm[9].sB; results[11].sA = vpm[10].sB; results[11].sB = vpm[11].sB; results[11].sC = vpm[12].sB; results[11].sD = vpm[13].sB; results[11].sE = vpm[14].sB; results[11].sF = vpm[15].sB; + results[12].s0 = vpm[0].sC; results[12].s1 = vpm[1].sC; results[12].s2 = vpm[2].sC; results[12].s3 = vpm[3].sC; results[12].s4 = vpm[4].sC; results[12].s5 = vpm[5].sC; results[12].s6 = vpm[6].sC; results[12].s7 = vpm[7].sC; results[12].s8 = vpm[8].sC; results[12].s9 = vpm[9].sC; results[12].sA = vpm[10].sC; results[12].sB = vpm[11].sC; results[12].sC = vpm[12].sC; results[12].sD = vpm[13].sC; results[12].sE = vpm[14].sC; results[12].sF = vpm[15].sC; + results[13].s0 = vpm[0].sD; results[13].s1 = vpm[1].sD; results[13].s2 = vpm[2].sD; results[13].s3 = vpm[3].sD; results[13].s4 = vpm[4].sD; results[13].s5 = vpm[5].sD; results[13].s6 = vpm[6].sD; results[13].s7 = vpm[7].sD; results[13].s8 = vpm[8].sD; results[13].s9 = vpm[9].sD; results[13].sA = vpm[10].sD; results[13].sB = vpm[11].sD; results[13].sC = vpm[12].sD; results[13].sD = vpm[13].sD; results[13].sE = vpm[14].sD; results[13].sF = vpm[15].sD; + results[14].s0 = vpm[0].sE; results[14].s1 = vpm[1].sE; results[14].s2 = vpm[2].sE; results[14].s3 = vpm[3].sE; results[14].s4 = vpm[4].sE; results[14].s5 = vpm[5].sE; results[14].s6 = vpm[6].sE; results[14].s7 = vpm[7].sE; results[14].s8 = vpm[8].sE; results[14].s9 = vpm[9].sE; results[14].sA = vpm[10].sE; results[14].sB = vpm[11].sE; results[14].sC = vpm[12].sE; results[14].sD = vpm[13].sE; results[14].sE = vpm[14].sE; results[14].sF = vpm[15].sE; + results[15].s0 = vpm[0].sF; results[15].s1 = vpm[1].sF; results[15].s2 = vpm[2].sF; results[15].s3 = vpm[3].sF; results[15].s4 = vpm[4].sF; results[15].s5 = vpm[5].sF; results[15].s6 = vpm[6].sF; results[15].s7 = vpm[7].sF; results[15].s8 = vpm[8].sF; results[15].s9 = vpm[9].sF; results[15].sA = vpm[10].sF; results[15].sB = vpm[11].sF; results[15].sC = vpm[12].sF; results[15].sD = vpm[13].sF; results[15].sE = vpm[14].sF; results[15].sF = vpm[15].sF; #endif // Multiplies by alpha and then stores the results into the destination matrix diff --git a/src/kernels/level3/xgemm_direct_part1.opencl b/src/kernels/level3/xgemm_direct_part1.opencl index e2f9c6a8..80d877cc 100644 --- a/src/kernels/level3/xgemm_direct_part1.opencl +++ b/src/kernels/level3/xgemm_direct_part1.opencl @@ -93,12 +93,12 @@ R"( // ================================================================================================= // Initializes the accumulation registers to zero -INLINE_FUNC void InitAccRegistersDirect(real cpm[NWID][MWID]) { +INLINE_FUNC void InitAccRegistersDirect(real cpd[NWID * MWID]) { #pragma unroll for (int _mi = 0; _mi < MWID; _mi += 1) { #pragma unroll for (int _ni = 0; _ni < NWID; _ni += 1) { - SetToZero(cpm[_ni][_mi]); + SetToZero(cpd[_ni * MWID + _mi]); } } } @@ -106,12 +106,12 @@ INLINE_FUNC void InitAccRegistersDirect(real cpm[NWID][MWID]) { // ================================================================================================= // Performs the actual computation: Cpm += Apm * Bpm -INLINE_FUNC void MultiplyAccumulateDirect(real cpm[NWID][MWID], real apm[MWID], real bpm[NWID]) { +INLINE_FUNC void MultiplyAccumulateDirect(real cpd[NWID * MWID], real apd[MWID], real bpd[NWID]) { #pragma unroll for (int _ni = 0; _ni < NWID; _ni += 1) { #pragma unroll for (int _mi = 0; _mi < MWID; _mi += 1) { - MultiplyAdd(cpm[_ni][_mi], apm[_mi], bpm[_ni]); + MultiplyAdd(cpd[_ni * MWID + _mi], apd[_mi], bpd[_ni]); } } } @@ -120,32 +120,32 @@ INLINE_FUNC void MultiplyAccumulateDirect(real cpm[NWID][MWID], real apm[MWID], // Loads global off-chip memory into thread-private register files. This function is specific for // loading the A input matrix. -INLINE_FUNC void GlobalToPrivateDirectA(const __global real* restrict agms, real apm[MWID], +INLINE_FUNC void GlobalToPrivateDirectA(const __global real* restrict agms, real apd[MWID], const int a_ld, const int a_offset, const int idm, const int idk, const int a_transpose, const int a_conjugate) { #pragma unroll for (int _mi = 0; _mi < MWID; _mi += 1) { const int a_index = (a_transpose) ? (idm + _mi)*a_ld + idk : idk*a_ld + (idm + _mi); - apm[_mi] = agms[a_index + a_offset]; - if (a_conjugate) { COMPLEX_CONJUGATE(apm[_mi]); } + apd[_mi] = agms[a_index + a_offset]; + if (a_conjugate) { COMPLEX_CONJUGATE(apd[_mi]); } } } // Same as above, but now for the B input matrix -INLINE_FUNC void GlobalToPrivateDirectB(const __global real* restrict bgms, real bpm[NWID], +INLINE_FUNC void GlobalToPrivateDirectB(const __global real* restrict bgms, real bpd[NWID], const int b_ld, const int b_offset, const int idn, const int idk, const int b_transpose, const int b_conjugate) { #pragma unroll for (int _ni = 0; _ni < NWID; _ni += 1) { const int b_index = (b_transpose) ? (idn + _ni)*b_ld + idk : idk*b_ld + (idn + _ni); - bpm[_ni] = bgms[b_index + b_offset]; - if (b_conjugate) { COMPLEX_CONJUGATE(bpm[_ni]); } + bpd[_ni] = bgms[b_index + b_offset]; + if (b_conjugate) { COMPLEX_CONJUGATE(bpd[_ni]); } } } // Loads global off-chip memory into thread-private register files. This function is specific for // loading the A input matrix. This is the same as above but now includes a bounds check. -INLINE_FUNC void GlobalToPrivateCheckedA(const __global real* restrict agms, real apm[MWID], +INLINE_FUNC void GlobalToPrivateCheckedA(const __global real* restrict agms, real apd[MWID], const int a_ld, const int a_offset, const int idm, const int idk, const int a_transpose, const int a_conjugate, const int kSizeM) { @@ -153,17 +153,17 @@ INLINE_FUNC void GlobalToPrivateCheckedA(const __global real* restrict agms, rea for (int _mi = 0; _mi < MWID; _mi += 1) { if (idm + _mi < kSizeM) { const int a_index = (a_transpose) ? (idm + _mi)*a_ld + idk : idk*a_ld + (idm + _mi); - apm[_mi] = agms[a_index + a_offset]; - if (a_conjugate) { COMPLEX_CONJUGATE(apm[_mi]); } + apd[_mi] = agms[a_index + a_offset]; + if (a_conjugate) { COMPLEX_CONJUGATE(apd[_mi]); } } else { - SetToZero(apm[_mi]); + SetToZero(apd[_mi]); } } } // Same as above, but now for the B input matrix -INLINE_FUNC void GlobalToPrivateCheckedB(const __global real* restrict bgms, real bpm[NWID], +INLINE_FUNC void GlobalToPrivateCheckedB(const __global real* restrict bgms, real bpd[NWID], const int b_ld, const int b_offset, const int idn, const int idk, const int b_transpose, const int b_conjugate, const int kSizeN) { @@ -171,11 +171,11 @@ INLINE_FUNC void GlobalToPrivateCheckedB(const __global real* restrict bgms, rea for (int _ni = 0; _ni < NWID; _ni += 1) { if (idn + _ni < kSizeN) { const int b_index = (b_transpose) ? (idn + _ni)*b_ld + idk : idk*b_ld + (idn + _ni); - bpm[_ni] = bgms[b_index + b_offset]; - if (b_conjugate) { COMPLEX_CONJUGATE(bpm[_ni]); } + bpd[_ni] = bgms[b_index + b_offset]; + if (b_conjugate) { COMPLEX_CONJUGATE(bpd[_ni]); } } else { - SetToZero(bpm[_ni]); + SetToZero(bpd[_ni]); } } } @@ -184,24 +184,24 @@ INLINE_FUNC void GlobalToPrivateCheckedB(const __global real* restrict bgms, rea // Caches on-chip local memory into per-thread private memory (registers). This function is specific // for caching the A input matrix. -INLINE_FUNC void LocalToPrivateDirectA(LOCAL_PTR real* alm, real apm[MWID], const int kg, +INLINE_FUNC void LocalToPrivateDirectA(LOCAL_PTR real* alm, real apd[MWID], const int kg, const int a_transpose) { #pragma unroll for (int _mi = 0; _mi < MWID; _mi += 1) { const int mg = _mi + get_local_id(0)*MWID; const int index = (a_transpose) ? mg*(WGD + PADA) + kg : kg*(WGD + PADA) + mg; - apm[_mi] = alm[index]; + apd[_mi] = alm[index]; } } // Same as above, but now for the B input matrix -INLINE_FUNC void LocalToPrivateDirectB(LOCAL_PTR real* blm, real bpm[NWID], const int kg, +INLINE_FUNC void LocalToPrivateDirectB(LOCAL_PTR real* blm, real bpd[NWID], const int kg, const int b_transpose) { #pragma unroll for (int _ni = 0; _ni < NWID; _ni += 1) { const int ng = _ni + get_local_id(1)*NWID; const int index = (b_transpose) ? ng*(WGD + PADB) + kg : kg*(WGD + PADB) + ng; - bpm[_ni] = blm[index]; + bpd[_ni] = blm[index]; } } @@ -209,7 +209,7 @@ INLINE_FUNC void LocalToPrivateDirectB(LOCAL_PTR real* blm, real bpm[NWID], cons // Merges the results in Cpm with the global array in Cgm. This also performs the multiplication // with the constants: Cgm = alpha*A*B + beta*Cgm = alpha*Cpm + beta*Cgm -INLINE_FUNC void StoreResultsDirect(__global real* cgm, real cpm[NWID][MWID], +INLINE_FUNC void StoreResultsDirect(__global real* cgm, real cpd[NWID * MWID], const int idm, const int idn, const real alpha, const real beta, const int c_ld, const int c_offset, const int c_transpose) { @@ -224,11 +224,11 @@ INLINE_FUNC void StoreResultsDirect(__global real* cgm, real cpm[NWID][MWID], // The final multiplication with alpha (in case beta == 0) real result; if (IsZero(beta)) { - Multiply(result, alpha, cpm[_ni][_mi]); + Multiply(result, alpha, cpd[_ni * MWID + _mi]); } // The final multiplication with alpha and the addition with beta*C else { - AXPBY(result, alpha, cpm[_ni][_mi], beta, cgm[c_index + c_offset]); + AXPBY(result, alpha, cpd[_ni * MWID + _mi], beta, cgm[c_index + c_offset]); } cgm[c_index + c_offset] = result; } @@ -237,7 +237,7 @@ INLINE_FUNC void StoreResultsDirect(__global real* cgm, real cpm[NWID][MWID], // Merges the results in Cpm with the global array in Cgm. This also performs the multiplication // with the constants: Cgm = alpha*A*B + beta*Cgm = alpha*Cpm + beta*Cgm -INLINE_FUNC void StoreResultsChecked(__global real* cgm, real cpm[NWID][MWID], +INLINE_FUNC void StoreResultsChecked(__global real* cgm, real cpd[NWID * MWID], const int idm, const int idn, const int kSizeM, const int kSizeN, const real alpha, const real beta, const int c_ld, const int c_offset, const int c_transpose) { @@ -253,11 +253,11 @@ INLINE_FUNC void StoreResultsChecked(__global real* cgm, real cpm[NWID][MWID], // The final multiplication with alpha (in case beta == 0) real result; if (IsZero(beta)) { - Multiply(result, alpha, cpm[_ni][_mi]); + Multiply(result, alpha, cpd[_ni * MWID + _mi]); } // The final multiplication with alpha and the addition with beta*C else { - AXPBY(result, alpha, cpm[_ni][_mi], beta, cgm[c_index + c_offset]); + AXPBY(result, alpha, cpd[_ni * MWID + _mi], beta, cgm[c_index + c_offset]); } cgm[c_index + c_offset] = result; } diff --git a/src/kernels/level3/xgemm_direct_part3.opencl b/src/kernels/level3/xgemm_direct_part3.opencl index b24695a1..f9af7a41 100644 --- a/src/kernels/level3/xgemm_direct_part3.opencl +++ b/src/kernels/level3/xgemm_direct_part3.opencl @@ -35,12 +35,12 @@ INLINE_FUNC void XgemmDirect(const int kSizeM, const int kSizeN, const int kSize const __global real* restrict bgms = (const __global real* restrict) bgm; // Allocates workitem-private memory (registers) - real apm[MWID]; - real bpm[NWID]; - real cpm[NWID][MWID]; + real apd[MWID]; + real bpd[NWID]; + real cpd[NWID * MWID]; // Initializes the accumulation registers - InitAccRegistersDirect(cpm); + InitAccRegistersDirect(cpd); // The faster version of GEMM is not allowed on the (incomplete) borders. Therefore, this section // processes only the main parts: output blocks of WGD by WGD. @@ -74,11 +74,11 @@ INLINE_FUNC void XgemmDirect(const int kSizeM, const int kSizeN, const int kSize int kg = pwi + _pit; // Loads data: local --> private (matrix A and B) - LocalToPrivateDirectA(alm, apm, kg, a_transpose); - LocalToPrivateDirectB(blm, bpm, kg, b_transpose); + LocalToPrivateDirectA(alm, apd, kg, a_transpose); + LocalToPrivateDirectB(blm, bpd, kg, b_transpose); - // Performs the accumulation (Cpm += Apm * Bpm) - MultiplyAccumulateDirect(cpm, apm, bpm); + // Performs the accumulation (Cpmd += Apmd * Bpmd) + MultiplyAccumulateDirect(cpd, apd, bpd); } } barrier(CLK_LOCAL_MEM_FENCE); @@ -88,15 +88,15 @@ INLINE_FUNC void XgemmDirect(const int kSizeM, const int kSizeN, const int kSize for (; kwg < kSizeK; ++kwg) { // Loads data: off-chip --> private (matrix A and B) - GlobalToPrivateDirectA(agms, apm, a_ld, a_offset, idm, kwg, a_transpose, a_conjugate); - GlobalToPrivateDirectB(bgms, bpm, b_ld, b_offset, idn, kwg, b_transpose, b_conjugate); + GlobalToPrivateDirectA(agms, apd, a_ld, a_offset, idm, kwg, a_transpose, a_conjugate); + GlobalToPrivateDirectB(bgms, bpd, b_ld, b_offset, idn, kwg, b_transpose, b_conjugate); - // Performs the accumulation (Cpm += Apm * Bpm) - MultiplyAccumulateDirect(cpm, apm, bpm); + // Performs the accumulation (Cpmd += Apmd * Bpmd) + MultiplyAccumulateDirect(cpd, apd, bpd); } // Stores a tile of results and performs the multiplication with alpha and beta - StoreResultsDirect(cgm, cpm, idm, idn, alpha, beta, c_ld, c_offset, c_transpose); + StoreResultsDirect(cgm, cpd, idm, idn, alpha, beta, c_ld, c_offset, c_transpose); } // Simple but slower version for the parts on the edge (incomplete tiles in M and N-dimensions) @@ -118,11 +118,11 @@ INLINE_FUNC void XgemmDirect(const int kSizeM, const int kSizeN, const int kSize int kg = pwi + _pit; // Loads data: local --> private (matrix A and B) - LocalToPrivateDirectA(alm, apm, kg, a_transpose); - LocalToPrivateDirectB(blm, bpm, kg, b_transpose); + LocalToPrivateDirectA(alm, apd, kg, a_transpose); + LocalToPrivateDirectB(blm, bpd, kg, b_transpose); - // Performs the accumulation (Cpm += Apm * Bpm) - MultiplyAccumulateDirect(cpm, apm, bpm); + // Performs the accumulation (Cpmd += Apmd * Bpmd) + MultiplyAccumulateDirect(cpd, apd, bpd); } } barrier(CLK_LOCAL_MEM_FENCE); @@ -132,15 +132,15 @@ INLINE_FUNC void XgemmDirect(const int kSizeM, const int kSizeN, const int kSize for (; kwg < kSizeK; ++kwg) { // Loads data: off-chip --> private (matrix A and B) - GlobalToPrivateCheckedA(agms, apm, a_ld, a_offset, idm, kwg, a_transpose, a_conjugate, kSizeM); - GlobalToPrivateCheckedB(bgms, bpm, b_ld, b_offset, idn, kwg, b_transpose, b_conjugate, kSizeN); + GlobalToPrivateCheckedA(agms, apd, a_ld, a_offset, idm, kwg, a_transpose, a_conjugate, kSizeM); + GlobalToPrivateCheckedB(bgms, bpd, b_ld, b_offset, idn, kwg, b_transpose, b_conjugate, kSizeN); - // Performs the accumulation (Cpm += Apm * Bpm) - MultiplyAccumulateDirect(cpm, apm, bpm); + // Performs the accumulation (Cpmd += Apmd * Bpmd) + MultiplyAccumulateDirect(cpd, apd, bpd); } // Stores a tile of results and performs the multiplication with alpha and beta - StoreResultsChecked(cgm, cpm, idm, idn, kSizeM, kSizeN, alpha, beta, c_ld, c_offset, c_transpose); + StoreResultsChecked(cgm, cpd, idm, idn, kSizeM, kSizeN, alpha, beta, c_ld, c_offset, c_transpose); } } diff --git a/src/kernels/level3/xgemm_part3.opencl b/src/kernels/level3/xgemm_part3.opencl index 7e46cef5..f12fb304 100644 --- a/src/kernels/level3/xgemm_part3.opencl +++ b/src/kernels/level3/xgemm_part3.opencl @@ -31,9 +31,9 @@ INLINE_FUNC void XgemmBody(const int kSizeM, const int kSizeN, const int kSizeK, ) { // Allocates workitem-private memory (registers) - #pragma promote_to_registers + //#pragma promote_to_registers realM apm[MWI/VWM]; - #pragma promote_to_registers + //#pragma promote_to_registers realN bpm[NWI/VWN]; // Combined thread identifier (volatile to disable caching) @@ -128,7 +128,7 @@ void XgemmUpper(const int kSizeN, const int kSizeK, #endif // Computes the matrix-multiplication and stores the result in register memory - #pragma promote_to_registers + //#pragma promote_to_registers realM cpm[NWI*(MWI/VWM)]; #if SA == 1 && SB == 1 XgemmBody(kSizeN, kSizeN, kSizeK, agm, bgm, cgm, cpm, alm, blm); @@ -169,7 +169,7 @@ void XgemmLower(const int kSizeN, const int kSizeK, #endif // Computes the matrix-multiplication and stores the result in register memory - #pragma promote_to_registers + //#pragma promote_to_registers realM cpm[NWI*(MWI/VWM)]; #if SA == 1 && SB == 1 XgemmBody(kSizeN, kSizeN, kSizeK, agm, bgm, cgm, cpm, alm, blm); @@ -214,7 +214,7 @@ void Xgemm(const int kSizeM, const int kSizeN, const int kSizeK, #endif // Computes the matrix-multiplication and stores the result in register memory - #pragma promote_to_registers + //#pragma promote_to_registers realM cpm[NWI*(MWI/VWM)]; #if SA == 1 && SB == 1 XgemmBody(kSizeM, kSizeN, kSizeK, agm, bgm, cgm, cpm, alm, blm); diff --git a/test/correctness/misc/preprocessor.cpp b/test/correctness/misc/preprocessor.cpp index fa0d2ccc..b6a12a38 100644 --- a/test/correctness/misc/preprocessor.cpp +++ b/test/correctness/misc/preprocessor.cpp @@ -219,13 +219,14 @@ size_t RunPreprocessor(int argc, char *argv[], const bool silent, const Precisio #include "../src/kernels/level3/level3.opencl" #include "../src/kernels/level3/transpose_pad.opencl" ; - if (TestKernel(device, context, "TransposePadMatrix", transpose_pad_sources, precision)) { passed++; } else { errors++; } + //if (TestKernel(device, context, "TransposePadMatrix", transpose_pad_sources, precision)) { passed++; } else { errors++; } // GEMM (in-direct) const auto gemm_sources = "#define KWI 2\n" "#define MWG 16\n" "#define NWG 16\n" + "#define SA 1\n" #include "../src/kernels/level3/xgemm_part1.opencl" #include "../src/kernels/level3/xgemm_part2.opencl" #include "../src/kernels/level3/xgemm_part3.opencl" |