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 Xher2 kernels for rank-2 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-2 matrix update kernel (HER2, HPR2, SYR2, SPR2)
21__kernel __attribute__((reqd_work_group_size(WGS1, WGS2, 1)))
22void Xher2(const int n,
23           const real_arg arg_alpha,
24           const __global real* restrict xgm, const int x_offset, const int x_inc,
25           const __global real* restrict 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_upper, 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  real xtvalues[WPT];
34  real ytvalues[WPT];
35
36  // Loads the X-vector
37  #pragma unroll
38  for (int w=0; w<WPT; ++w) {
39    const int id2 = w*get_global_size(1) + get_global_id(1);
40    xvalues[w] = LoadVector(id2, n, xgm, x_offset, x_inc, !is_rowmajor);
41  }
42
43  // Loads the X-transposed-vector
44  #pragma unroll
45  for (int w=0; w<WPT; ++w) {
46    const int id1 = w*get_global_size(0) + get_global_id(0);
47    xtvalues[w] = LoadVector(id1, n, xgm, x_offset, x_inc, is_rowmajor);
48  }
49
50  // Loads the Y-vector
51  #pragma unroll
52  for (int w=0; w<WPT; ++w) {
53    const int id1 = w*get_global_size(0) + get_global_id(0);
54    yvalues[w] = LoadVector(id1, n, ygm, y_offset, y_inc, is_rowmajor);
55  }
56
57  // Loads the Y-transposed-vector
58  #pragma unroll
59  for (int w=0; w<WPT; ++w) {
60    const int id2 = w*get_global_size(1) + get_global_id(1);
61    ytvalues[w] = LoadVector(id2, n, ygm, y_offset, y_inc, !is_rowmajor);
62  }
63
64  // Sets the proper value of alpha in case conjugation is needed
65  real alpha1 = alpha;
66  real alpha2 = alpha;
67  #if defined(ROUTINE_HER2) || defined(ROUTINE_HPR2)
68    if (is_rowmajor) {
69      COMPLEX_CONJUGATE(alpha1);
70    }
71    else {
72      COMPLEX_CONJUGATE(alpha2);
73    }
74  #endif
75
76  // Loops over the work per thread twice
77  #pragma unroll
78  for (int w1=0; w1<WPT; ++w1) {
79    #pragma unroll
80    for (int w2=0; w2<WPT; ++w2) {
81
82      // Global thread IDs
83      const int id1 = w1*get_global_size(0) + get_global_id(0);
84      const int id2 = w2*get_global_size(1) + get_global_id(1);
85
86      // Skip these threads if they do not contain threads contributing to the matrix-triangle
87      if ((is_upper && (id1 > id2)) || (!is_upper && (id2 > id1))) {
88        // Do nothing
89      }
90
91      // Loads A, performs the operation, and stores the result into A
92      else {
93        MatrixUpdate2(id1, id2, n, n, agm, a_offset, a_ld,
94                      alpha1, xvalues[w2], yvalues[w1],
95                      alpha2, xtvalues[w1], ytvalues[w2], is_upper);
96      }
97    }
98  }
99}
100
101// =================================================================================================
102
103// End of the C++11 raw string literal
104)"
105
106// =================================================================================================
107