MAGMA  1.2.0
MatrixAlgebraonGPUandMulticoreArchitectures
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups
zlarfb_gpu.cpp
Go to the documentation of this file.
1 /*
2  -- MAGMA (version 1.2.0) --
3  Univ. of Tennessee, Knoxville
4  Univ. of California, Berkeley
5  Univ. of Colorado, Denver
6  May 2012
7 
8  @precisions normal z -> s d c
9 
10 */
11 #include "common_magma.h"
12 
13 // === Define what BLAS to use ============================================
14 #define PRECISION_z
15 #if (defined(PRECISION_s) || defined(PRECISION_d))
16 // #define cublasZgemm magmablas_zgemm
17 #endif
18 // === End defining what BLAS to use =======================================
19 
20 extern "C" magma_int_t
21 magma_zlarfb_gpu( char side, char trans, char direct, char storev,
22  magma_int_t m, magma_int_t n, magma_int_t k,
23  cuDoubleComplex *dV, magma_int_t ldv,
24  cuDoubleComplex *dT, magma_int_t ldt,
25  cuDoubleComplex *dC, magma_int_t ldc,
26  cuDoubleComplex *dwork, magma_int_t ldwork)
27 {
28 /* -- MAGMA (version 1.2.0) --
29  Univ. of Tennessee, Univ. of California Berkeley
30  May 2012
31 
32  Purpose
33  =======
34  ZLARFB applies a complex block reflector H or its transpose H' to a
35  COMPLEX_16 m by n matrix C, from the left.
36 
37  Arguments
38  =========
39  SIDE (input) CHARACTER
40  = 'L': apply H or H' from the Left
41  = 'R': apply H or H' from the Right (Not implemented)
42 
43  TRANS (input) CHARACTER
44  = 'N': apply H (No transpose) (Not implemented)
45  = 'C': apply H' (Conjugate transpose)
46 
47  DIRECT (input) CHARACTER
48  Indicates how H is formed from a product of elementary
49  reflectors
50  = 'F': H = H(1) H(2) . . . H(k) (Forward)
51  = 'B': H = H(k) . . . H(2) H(1) (Backward)
52 
53  STOREV (input) CHARACTER
54  Indicates how the vectors which define the elementary
55  reflectors are stored:
56  = 'C': Columnwise
57  = 'R': Rowwise
58 
59  M (input) INTEGER
60  The number of rows of the matrix C.
61 
62  N (input) INTEGER
63  The number of columns of the matrix C.
64 
65  K (input) INTEGER
66  The order of the matrix T (= the number of elementary
67  reflectors whose product defines the block reflector).
68 
69  DV (input) COMPLEX_16 array, dimension (LDV,K)
70  The matrix V. See further details.
71 
72  LDV (input) INTEGER
73  The leading dimension of the array V. LDV >= max(1,M);
74 
75  DT (input) COMPLEX_16 array, dimension (LDT,K)
76  The triangular k by k matrix T in the representation of the
77  block reflector.
78 
79  LDT (input) INTEGER
80  The leading dimension of the array T. LDT >= K.
81 
82  DC (input/output) COMPLEX_16 array, dimension (LDC,N)
83  On entry, the m by n matrix C.
84  On exit, C is overwritten by H*C.
85 
86  LDC (input) INTEGER
87  The leading dimension of the array C. LDA >= max(1,M).
88 
89  WORK (workspace) COMPLEX_16 array, dimension (LDWORK,K)
90 
91  LDWORK (input) INTEGER
92  The leading dimension of the array WORK. LDWORK >= max(1,N);
93  =================================================================== */
94 
95  cuDoubleComplex c_zero = MAGMA_Z_ZERO;
96  cuDoubleComplex c_one = MAGMA_Z_ONE;
97  cuDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
98 
99  /* Function Body */
100  if (m <= 0 || n <= 0) {
101  return MAGMA_SUCCESS;
102  }
103 
104  char transt;
105  if (trans == 'N' || trans == 'n')
106  transt = MagmaConjTrans;
107  else
108  transt = MagmaNoTrans;
109 
110  if ( ( side == 'r' || side == 'R') ) {
111  fprintf(stderr, "The case (side == right) is not implemented\n");
112  magma_xerbla( __func__, 1 );
113  return MAGMA_ERR_ILLEGAL_VALUE;
114  }
115 
116  if ( storev == 'c' || storev == 'C') {
117  /*
118  if (n==1 && m%32==0){
119  // This is used when we have to apply H on only one vector
120  magmablas_zgemvt(m, k, 1., dv_ref(0,0), ldv, dc_ref(0, 0), dwork);
121  printf("m= %d, n = %d, ldwork = %d\n", m, k, ldwork);
122  }
123  else
124  */
125  cublasZgemm( MagmaConjTrans, MagmaNoTrans,
126  n, k, m,
127  c_one, dC, ldc,
128  dV, ldv,
129  c_zero, dwork, ldwork);
130 
131  if (direct == 'F' || direct =='f')
132  cublasZtrmm( MagmaRight, MagmaUpper, transt, MagmaNonUnit,
133  n, k,
134  c_one, dT, ldt,
135  dwork, ldwork);
136  else
137  cublasZtrmm( MagmaRight, MagmaLower, transt, MagmaNonUnit,
138  n, k,
139  c_one, dT, ldt,
140  dwork, ldwork);
141 
142  cublasZgemm( MagmaNoTrans, MagmaConjTrans,
143  m, n, k,
144  c_neg_one, dV, ldv,
145  dwork, ldwork,
146  c_one, dC, ldc);
147  }
148  else {
149  cublasZgemm( MagmaNoTrans, MagmaConjTrans,
150  m, k, n,
151  c_one, dC, ldc,
152  dV, ldv,
153  c_zero, dwork, ldwork);
154 
155  cublasZtrmm( MagmaRight, MagmaUpper, transt, MagmaNonUnit,
156  m, k,
157  c_one, dT, ldt,
158  dwork, ldwork);
159 
160  cublasZgemm( MagmaNoTrans, MagmaNoTrans,
161  m, n, k,
162  c_neg_one, dwork, ldwork,
163  dV, ldv,
164  c_one, dC, ldc);
165  }
166  return MAGMA_SUCCESS;
167 } /* magma_zlarfb */