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 Xher 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// Symmetric version of the rank-1 matrix update kernel (HER, HPR, SYR, SPR)
21__kernel __attribute__((reqd_work_group_size(WGS1, WGS2, 1)))
22void Xher(const int n,
23          const real_arg arg_alpha,
24          const __global real* restrict xgm, const int x_offset, const int x_inc,
25          __global real* restrict agm, const int a_offset, const int a_ld,
26          const int is_upper, const int is_rowmajor) {
27  const real alpha = GetRealArg(arg_alpha);
28
29  // Register storage for X and XT
30  real xvalues[WPT];
31  real xtvalues[WPT];
32
33  // Loads the X-vector
34  #pragma unroll
35  for (int w=0; w<WPT; ++w) {
36    const int id2 = w*get_global_size(1) + get_global_id(1);
37    xvalues[w] = LoadVector(id2, n, xgm, x_offset, x_inc, !is_rowmajor);
38  }
39
40  // Loads the X-transposed-vector
41  #pragma unroll
42  for (int w=0; w<WPT; ++w) {
43    const int id1 = w*get_global_size(0) + get_global_id(0);
44    xtvalues[w] = LoadVector(id1, n, xgm, x_offset, x_inc, is_rowmajor);
45  }
46
47  // Loops over the work per thread twice
48  #pragma unroll
49  for (int w1=0; w1<WPT; ++w1) {
50    #pragma unroll
51    for (int w2=0; w2<WPT; ++w2) {
52
53      // Global thread IDs
54      const int id1 = w1*get_global_size(0) + get_global_id(0);
55      const int id2 = w2*get_global_size(1) + get_global_id(1);
56
57      // Skip these threads if they do not contain threads contributing to the matrix-triangle
58      if ((is_upper && (id1 > id2)) || (!is_upper && (id2 > id1))) {
59        // Do nothing
60      }
61
62      // Loads A, performs the operation, and stores the result into A
63      else {
64        MatrixUpdate(id1, id2, n, n, agm, a_offset, a_ld, alpha, xvalues[w2], xtvalues[w1], is_upper);
65      }
66    }
67  }
68}
69
70// =================================================================================================
71
72// End of the C++11 raw string literal
73)"
74
75// =================================================================================================
76