1
2// =================================================================================================
3// This file is part of the CLBlast project. The project is licensed under Apache Version 2.0. This
4// project loosely follows the Google C++ styleguide and uses a tab-size of two spaces and a max-
5// width of 100 characters per line.
6//
7// Author(s):
8//   Cedric Nugteren <www.cedricnugteren.nl>
9//
10// This file contains the Xger kernels for rank-1 matrix update.
11//
12// =================================================================================================
13
14// Enables loading of this file using the C++ pre-processor's #include (C++11 standard raw string
15// literal). Comment-out this line for syntax-highlighting when developing.
16R"(
17
18// =================================================================================================
19
20// Regular version of the rank-1 matrix update kernel (GER, GERU, GERC)
21__kernel __attribute__((reqd_work_group_size(WGS1, WGS2, 1)))
22void Xger(const int max1, const int max2,
23          const real_arg arg_alpha,
24          const __global real* restrict xgm, const int x_offset, const int x_inc,
25          const __global real* ygm, const int y_offset, const int y_inc,
26          __global real* restrict agm, const int a_offset, const int a_ld,
27          const int is_rowmajor) {
28  const real alpha = GetRealArg(arg_alpha);
29
30  // Register storage for X and Y
31  real xvalues[WPT];
32  real yvalues[WPT];
33
34  // Row-major version
35  if (is_rowmajor) {
36
37    // Loads the X-vector
38    #pragma unroll
39    for (int w=0; w<WPT; ++w) {
40      const int id2 = w*get_global_size(1) + get_global_id(1);
41      xvalues[w] = LoadVector(id2, max2, xgm, x_offset, x_inc, false);
42    }
43
44    // Loads the Y-vector
45    #pragma unroll
46    for (int w=0; w<WPT; ++w) {
47      const int id1 = w*get_global_size(0) + get_global_id(0);
48      yvalues[w] = LoadVector(id1, max1, ygm, y_offset, y_inc, true);
49    }
50
51    // Loops over the work per thread twice
52    #pragma unroll
53    for (int w1=0; w1<WPT; ++w1) {
54      #pragma unroll
55      for (int w2=0; w2<WPT; ++w2) {
56
57        // Global thread IDs
58        const int id1 = w1*get_global_size(0) + get_global_id(0);
59        const int id2 = w2*get_global_size(1) + get_global_id(1);
60
61        // Loads A, performs the operation, and stores the result into A
62        MatrixUpdate(id1, id2, max1, max2, agm, a_offset, a_ld,
63                     alpha, xvalues[w2], yvalues[w1], false);
64      }
65    }
66  }
67
68  // Col-major version
69  else {
70
71    // Loads the X-vector
72    #pragma unroll
73    for (int w=0; w<WPT; ++w) {
74      const int id1 = w*get_global_size(0) + get_global_id(0);
75      xvalues[w] = LoadVector(id1, max1, xgm, x_offset, x_inc, false);
76    }
77
78    // Loads the Y-vector
79    #pragma unroll
80    for (int w=0; w<WPT; ++w) {
81      const int id2 = w*get_global_size(1) + get_global_id(1);
82      yvalues[w] = LoadVector(id2, max2, ygm, y_offset, y_inc, true);
83    }
84
85    // Loops over the work per thread twice
86    #pragma unroll
87    for (int w1=0; w1<WPT; ++w1) {
88      #pragma unroll
89      for (int w2=0; w2<WPT; ++w2) {
90
91        // Global thread IDs
92        const int id1 = w1*get_global_size(0) + get_global_id(0);
93        const int id2 = w2*get_global_size(1) + get_global_id(1);
94
95        // Loads A, performs the operation, and stores the result into A
96        MatrixUpdate(id1, id2, max1, max2, agm, a_offset, a_ld,
97                     alpha, xvalues[w1], yvalues[w2], false);
98      }
99    }
100  }
101}
102
103// =================================================================================================
104
105// End of the C++11 raw string literal
106)"
107
108// =================================================================================================
109