From 66649125b8192c2b4c83064fd4adb9ce6b384987 Mon Sep 17 00:00:00 2001 From: Doc CI Action Date: Mon, 25 Nov 2024 18:53:32 +0000 Subject: [PATCH] Doc: TFactor: Separate larft_gemv kernel and add a faster option (#1219) --- master/blas_2tile_8h.html | 56 +- master/blas_2tile_8h_source.html | 689 ++++++++---------- .../dir_2e3e3bc658385778082583f320919a4c.html | 2 + .../dir_5ca20fc8e29b0f8739133582ef745158.html | 2 + master/factorization_2qr_2api_8h_source.html | 2 +- master/files.html | 4 +- master/gpublas_8h.html | 152 ++++ master/gpublas_8h_source.html | 237 ++++++ master/larft_8h_source.html | 142 ++++ master/qr_8h_source.html | 2 +- master/search/all_6.js | 11 +- master/search/files_6.js | 3 +- master/t__factor__impl_8h_source.html | 660 ++++++++--------- 13 files changed, 1158 insertions(+), 804 deletions(-) create mode 100644 master/gpublas_8h.html create mode 100644 master/gpublas_8h_source.html create mode 100644 master/larft_8h_source.html diff --git a/master/blas_2tile_8h.html b/master/blas_2tile_8h.html index ddf21eabce..2c4bf0c428 100644 --- a/master/blas_2tile_8h.html +++ b/master/blas_2tile_8h.html @@ -74,7 +74,6 @@
-Macros | Functions
tile.h File Reference
@@ -91,44 +90,14 @@ #include <dlaf/sender/transform.h>
#include <dlaf/types.h>
#include <dlaf/util_blas.h>
-#include <whip.hpp>
#include <dlaf/gpu/blas/api.h>
-#include <dlaf/gpu/blas/error.h>
+#include <dlaf/gpu/blas/gpublas.h>
#include <dlaf/util_cublas.h>

Go to the source code of this file.

- - - -

-Macros

#define DLAF_DECLARE_GPUBLAS_OP(Name)
 
- - - - - - - - - - - - - - - - @@ -195,28 +164,7 @@

Functions

- dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_OP (Axpy, axpy)
 
- dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_OP (Gemv, gemv)
 
- dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_OP (Trmv, trmv)
 
- dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_OP (Gemm, gemm)
 
- dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_SYHE_OP (Hemm, mm)
 
- dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_SYHE_OP (Her2k, r2k)
 
- dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_SYHE_OP (Herk, rk)
 
- dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_OP (Trsm, trsm)
 
template<Backend B, class T , Device D>
void dlaf::tile::gemm (const blas::Op op_a, const blas::Op op_b, const T alpha, const Tile< const T, D > &a, const Tile< const T, D > &b, const T beta, const Tile< T, D > &c)
 

Detailed Description

Provides Tile wrappers for BLAS operations.

-

Macro Definition Documentation

- -

◆ DLAF_DECLARE_GPUBLAS_OP

- -
-
- - - - - - - - -
#define DLAF_DECLARE_GPUBLAS_OP( Name)
-
-Value:
template <typename T> \
-
struct Name
-
-
-
-

Function Documentation

+

Function Documentation

◆ gemm() [1/3]

diff --git a/master/blas_2tile_8h_source.html b/master/blas_2tile_8h_source.html index d02fe35bed..24944242af 100644 --- a/master/blas_2tile_8h_source.html +++ b/master/blas_2tile_8h_source.html @@ -109,420 +109,299 @@
28#include <dlaf/util_blas.h>
29
30#ifdef DLAF_WITH_GPU
-
31#include <whip.hpp>
-
32
-
33#include <dlaf/gpu/blas/api.h>
-
34#include <dlaf/gpu/blas/error.h>
-
35#include <dlaf/util_cublas.h>
-
36
-
37#ifdef DLAF_WITH_HIP
-
38
-
39#define DLAF_GET_ROCBLAS_WORKSPACE(f) \
-
40 [&]() { \
-
41 std::size_t workspace_size; \
-
42 DLAF_GPUBLAS_CHECK_ERROR( \
-
43 rocblas_start_device_memory_size_query(static_cast<rocblas_handle>(handle))); \
-
44 DLAF_ROCBLAS_WORKSPACE_CHECK_ERROR(rocblas_##f(handle, std::forward<Args>(args)...)); \
-
45 DLAF_GPUBLAS_CHECK_ERROR(rocblas_stop_device_memory_size_query(static_cast<rocblas_handle>(handle), \
-
46 &workspace_size)); \
-
47 return ::dlaf::memory::MemoryView<std::byte, Device::GPU>(to_int(workspace_size)); \
-
48 }();
-
49
-
50namespace dlaf::tile::internal {
-
51inline void extendROCBlasWorkspace(cublasHandle_t handle,
-
52 ::dlaf::memory::MemoryView<std::byte, Device::GPU>&& workspace) {
-
53 whip::stream_t stream;
-
54 DLAF_GPUBLAS_CHECK_ERROR(cublasGetStream(handle, &stream));
-
55 auto f = [workspace = std::move(workspace)](whip::error_t status) { whip::check_error(status); };
-
56 pika::cuda::experimental::detail::add_event_callback(std::move(f), stream);
-
57}
-
58}
+
31#include <dlaf/gpu/blas/api.h>
+
32#include <dlaf/gpu/blas/gpublas.h>
+
33#include <dlaf/util_cublas.h>
+
34#endif
+
35
+
36namespace dlaf {
+
37namespace tile {
+
38using matrix::Tile;
+
39
+
40#ifdef DLAF_DOXYGEN
+
41
+
45template <Backend B, class T, Device D>
+
46void gemm(const blas::Op op_a, const blas::Op op_b, const T alpha, const Tile<const T, D>& a,
+
47 const Tile<const T, D>& b, const T beta, const Tile<T, D>& c);
+
48
+
51template <Backend B, typename Sender,
+
52 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
+
53auto gemm(const dlaf::internal::Policy<B>& p, Sender&& s);
+
54
+
57template <Backend B>
+
58auto gemm(const dlaf::internal::Policy<B>& p);
59
-
60#define DLAF_DEFINE_GPUBLAS_OP(Name, Type, f) \
-
61 template <> \
-
62 struct Name<Type> { \
-
63 template <typename... Args> \
-
64 static void call(cublasHandle_t handle, Args&&... args) { \
-
65 auto workspace = DLAF_GET_ROCBLAS_WORKSPACE(f); \
-
66 DLAF_GPUBLAS_CHECK_ERROR(rocblas_set_workspace(static_cast<rocblas_handle>(handle), workspace(), \
-
67 to_sizet(workspace.size()))); \
-
68 DLAF_GPUBLAS_CHECK_ERROR(rocblas_##f(handle, std::forward<Args>(args)...)); \
-
69 DLAF_GPUBLAS_CHECK_ERROR(rocblas_set_workspace(static_cast<rocblas_handle>(handle), nullptr, 0)); \
-
70 ::dlaf::tile::internal::extendROCBlasWorkspace(handle, std::move(workspace)); \
-
71 } \
-
72 }
-
73
-
74#elif defined(DLAF_WITH_CUDA)
-
75
-
76#define DLAF_DEFINE_GPUBLAS_OP(Name, Type, f) \
-
77 template <> \
-
78 struct Name<Type> { \
-
79 template <typename... Args> \
-
80 static void call(Args&&... args) { \
-
81 DLAF_GPUBLAS_CHECK_ERROR(cublas##f##_v2(std::forward<Args>(args)...)); \
-
82 } \
-
83 }
+
63template <Backend B, class T, Device D>
+
64void hemm(const blas::Side side, const blas::Uplo uplo, const T alpha, const Tile<const T, D>& a,
+
65 const Tile<const T, D>& b, const T beta, const Tile<T, D>& c);
+
66
+
69template <Backend B, typename Sender,
+
70 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
+
71auto hemm(const dlaf::internal::Policy<B>& p, Sender&& s);
+
72
+
75template <Backend B>
+
76auto hemm(const dlaf::internal::Policy<B>& p);
+
77
+
81template <Backend B, class T, Device D>
+
82void her2k(const blas::Uplo uplo, const blas::Op op, const T alpha, const Tile<const T, D>& a,
+
83 const Tile<const T, D>& b, const BaseType<T> beta, const Tile<T, D>& c);
84
-
85#endif
-
86
-
87#define DLAF_DECLARE_GPUBLAS_OP(Name) \
-
88 template <typename T> \
-
89 struct Name
+
87template <Backend B, typename Sender,
+
88 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
+
89auto her2k(const dlaf::internal::Policy<B>& p, Sender&& s);
90
-
91#ifdef DLAF_WITH_HIP
-
92#define DLAF_MAKE_GPUBLAS_OP(Name, f) \
-
93 DLAF_DECLARE_GPUBLAS_OP(Name); \
-
94 DLAF_DEFINE_GPUBLAS_OP(Name, float, s##f); \
-
95 DLAF_DEFINE_GPUBLAS_OP(Name, double, d##f); \
-
96 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<float>, c##f); \
-
97 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<double>, z##f)
-
98
-
99#define DLAF_MAKE_GPUBLAS_SYHE_OP(Name, f) \
-
100 DLAF_DECLARE_GPUBLAS_OP(Name); \
-
101 DLAF_DEFINE_GPUBLAS_OP(Name, float, ssy##f); \
-
102 DLAF_DEFINE_GPUBLAS_OP(Name, double, dsy##f); \
-
103 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<float>, che##f); \
-
104 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<double>, zhe##f)
-
105
-
106#elif defined(DLAF_WITH_CUDA)
-
107#define DLAF_MAKE_GPUBLAS_OP(Name, f) \
-
108 DLAF_DECLARE_GPUBLAS_OP(Name); \
-
109 DLAF_DEFINE_GPUBLAS_OP(Name, float, S##f); \
-
110 DLAF_DEFINE_GPUBLAS_OP(Name, double, D##f); \
-
111 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<float>, C##f); \
-
112 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<double>, Z##f)
+
93template <Backend B>
+
94auto her2k(const dlaf::internal::Policy<B>& p);
+
95
+
99template <Backend B, class T, Device D>
+
100void herk(const blas::Uplo uplo, const blas::Op op, const BaseType<T> alpha, const Tile<const T, D>& a,
+
101 const BaseType<T> beta, const Tile<T, D>& c);
+
102
+
105template <Backend B, typename Sender,
+
106 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
+
107auto herk(const dlaf::internal::Policy<B>& p, Sender&& s);
+
108
+
111template <Backend B>
+
112auto herk(const dlaf::internal::Policy<B>& p);
113
-
114#define DLAF_MAKE_GPUBLAS_SYHE_OP(Name, f) \
-
115 DLAF_DECLARE_GPUBLAS_OP(Name); \
-
116 DLAF_DEFINE_GPUBLAS_OP(Name, float, Ssy##f); \
-
117 DLAF_DEFINE_GPUBLAS_OP(Name, double, Dsy##f); \
-
118 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<float>, Che##f); \
-
119 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<double>, Zhe##f)
-
120#endif
+
117template <Backend B, class T, Device D>
+
118void trmm(const dlaf::internal::Policy<B>& policy, const blas::Side side, const blas::Uplo uplo,
+
119 const blas::Op op, const blas::Diag diag, const T alpha, const Tile<const T, D>& a,
+
120 const Tile<T, D>& b);
121
-
122namespace dlaf::gpublas::internal {
-
123
-
124// Level 1
-
125DLAF_MAKE_GPUBLAS_OP(Axpy, axpy);
-
126
-
127// Level 2
-
128DLAF_MAKE_GPUBLAS_OP(Gemv, gemv);
-
129
-
130DLAF_MAKE_GPUBLAS_OP(Trmv, trmv);
-
131
-
132// Level 3
-
133DLAF_MAKE_GPUBLAS_OP(Gemm, gemm);
-
134
-
135DLAF_MAKE_GPUBLAS_SYHE_OP(Hemm, mm);
-
136
-
137DLAF_MAKE_GPUBLAS_SYHE_OP(Her2k, r2k);
-
138
-
139DLAF_MAKE_GPUBLAS_SYHE_OP(Herk, rk);
-
140
-
141#if defined(DLAF_WITH_CUDA)
-
142DLAF_MAKE_GPUBLAS_OP(Trmm, trmm);
-
143#elif defined(DLAF_WITH_HIP)
-
144
-
145#if ROCBLAS_VERSION_MAJOR >= 3 && defined(ROCBLAS_V3)
-
146DLAF_MAKE_GPUBLAS_OP(Trmm, trmm);
-
147#else
-
148DLAF_MAKE_GPUBLAS_OP(Trmm, trmm_outofplace);
-
149#endif
-
150
-
151#endif
+
124template <Backend B, typename Sender,
+
125 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
+
126auto trmm(const dlaf::internal::Policy<B>& p, Sender&& s);
+
127
+
130template <Backend B>
+
131auto trmm(const dlaf::internal::Policy<B>& p);
+
132
+
137template <Backend B, class T, Device D>
+
138void trmm3(const dlaf::internal::Policy<B>& policy, const blas::Side side, const blas::Uplo uplo,
+
139 const blas::Op op, const blas::Diag diag, const T alpha, const Tile<const T, D>& a,
+
140 const Tile<const T, D>& b, const Tile<T, D>& c);
+
141
+
144template <Backend B, typename Sender,
+
145 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
+
146auto trmm3(const dlaf::internal::Policy<B>& p, Sender&& s);
+
147
+
150template <Backend B>
+
151auto trmm3(const dlaf::internal::Policy<B>& p);
152
-
153DLAF_MAKE_GPUBLAS_OP(Trsm, trsm);
-
154}
-
155#endif
-
156
-
157namespace dlaf {
-
158namespace tile {
-
159using matrix::Tile;
+
156template <Backend B, class T, Device D>
+
157void trsm(const dlaf::internal::Policy<B>& policy, const blas::Side side, const blas::Uplo uplo,
+
158 const blas::Op op, const blas::Diag diag, const T alpha, const Tile<const T, D>& a,
+
159 const Tile<T, D>& b);
160
-
161#ifdef DLAF_DOXYGEN
-
162
-
166template <Backend B, class T, Device D>
-
167void gemm(const blas::Op op_a, const blas::Op op_b, const T alpha, const Tile<const T, D>& a,
-
168 const Tile<const T, D>& b, const T beta, const Tile<T, D>& c);
-
169
-
172template <Backend B, typename Sender,
-
173 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
-
174auto gemm(const dlaf::internal::Policy<B>& p, Sender&& s);
-
175
-
178template <Backend B>
-
179auto gemm(const dlaf::internal::Policy<B>& p);
-
180
-
184template <Backend B, class T, Device D>
-
185void hemm(const blas::Side side, const blas::Uplo uplo, const T alpha, const Tile<const T, D>& a,
-
186 const Tile<const T, D>& b, const T beta, const Tile<T, D>& c);
-
187
-
190template <Backend B, typename Sender,
-
191 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
-
192auto hemm(const dlaf::internal::Policy<B>& p, Sender&& s);
+
163template <Backend B, typename Sender,
+
164 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
+
165auto trsm(const dlaf::internal::Policy<B>& p, Sender&& s);
+
166
+
169template <Backend B>
+
170auto trsm(const dlaf::internal::Policy<B>& p);
+
171#else
+
172
+
173namespace internal {
+
174
+
175template <class T>
+
176void gemm(const blas::Op op_a, const blas::Op op_b, const T alpha, const Tile<const T, Device::CPU>& a,
+
177 const Tile<const T, Device::CPU>& b, const T beta, const Tile<T, Device::CPU>& c) noexcept {
+
178 auto s = tile::internal::getGemmSizes(op_a, op_b, a, b, c);
+
179 common::internal::SingleThreadedBlasScope single;
+
180 blas::gemm(blas::Layout::ColMajor, op_a, op_b, s.m, s.n, s.k, alpha, a.ptr(), a.ld(), b.ptr(), b.ld(),
+
181 beta, c.ptr(), c.ld());
+
182}
+
183
+
184template <class T>
+
185void hemm(const blas::Side side, const blas::Uplo uplo, const T alpha,
+
186 const Tile<const T, Device::CPU>& a, const Tile<const T, Device::CPU>& b, const T beta,
+
187 const Tile<T, Device::CPU>& c) {
+
188 auto s = tile::internal::getHemmSizes(side, a, b, c);
+
189 common::internal::SingleThreadedBlasScope single;
+
190 blas::hemm(blas::Layout::ColMajor, side, uplo, s.m, s.n, alpha, a.ptr(), a.ld(), b.ptr(), b.ld(), beta,
+
191 c.ptr(), c.ld());
+
192}
193
-
196template <Backend B>
-
197auto hemm(const dlaf::internal::Policy<B>& p);
-
198
-
202template <Backend B, class T, Device D>
-
203void her2k(const blas::Uplo uplo, const blas::Op op, const T alpha, const Tile<const T, D>& a,
-
204 const Tile<const T, D>& b, const BaseType<T> beta, const Tile<T, D>& c);
-
205
-
208template <Backend B, typename Sender,
-
209 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
-
210auto her2k(const dlaf::internal::Policy<B>& p, Sender&& s);
-
211
-
214template <Backend B>
-
215auto her2k(const dlaf::internal::Policy<B>& p);
-
216
-
220template <Backend B, class T, Device D>
-
221void herk(const blas::Uplo uplo, const blas::Op op, const BaseType<T> alpha, const Tile<const T, D>& a,
-
222 const BaseType<T> beta, const Tile<T, D>& c);
-
223
-
226template <Backend B, typename Sender,
-
227 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
-
228auto herk(const dlaf::internal::Policy<B>& p, Sender&& s);
-
229
-
232template <Backend B>
-
233auto herk(const dlaf::internal::Policy<B>& p);
-
234
-
238template <Backend B, class T, Device D>
-
239void trmm(const dlaf::internal::Policy<B>& policy, const blas::Side side, const blas::Uplo uplo,
-
240 const blas::Op op, const blas::Diag diag, const T alpha, const Tile<const T, D>& a,
-
241 const Tile<T, D>& b);
-
242
-
245template <Backend B, typename Sender,
-
246 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
-
247auto trmm(const dlaf::internal::Policy<B>& p, Sender&& s);
+
194template <class T>
+
195void her2k(const blas::Uplo uplo, const blas::Op op, const T alpha, const Tile<const T, Device::CPU>& a,
+
196 const Tile<const T, Device::CPU>& b, const BaseType<T> beta,
+
197 const Tile<T, Device::CPU>& c) noexcept {
+
198 auto s = tile::internal::getHer2kSizes(op, a, b, c);
+
199 common::internal::SingleThreadedBlasScope single;
+
200 blas::her2k(blas::Layout::ColMajor, uplo, op, s.n, s.k, alpha, a.ptr(), a.ld(), b.ptr(), b.ld(), beta,
+
201 c.ptr(), c.ld());
+
202}
+
203
+
204template <class T>
+
205void herk(const blas::Uplo uplo, const blas::Op op, const BaseType<T> alpha,
+
206 const Tile<const T, Device::CPU>& a, const BaseType<T> beta,
+
207 const Tile<T, Device::CPU>& c) noexcept {
+
208 auto s = tile::internal::getHerkSizes(op, a, c);
+
209 common::internal::SingleThreadedBlasScope single;
+
210 blas::herk(blas::Layout::ColMajor, uplo, op, s.n, s.k, alpha, a.ptr(), a.ld(), beta, c.ptr(), c.ld());
+
211}
+
212
+
213// Triangular matrix-matrix multiplication.
+
214template <class T>
+
215void trmm(const blas::Side side, const blas::Uplo uplo, const blas::Op op, const blas::Diag diag,
+
216 const T alpha, const Tile<const T, Device::CPU>& a, const Tile<T, Device::CPU>& b) noexcept {
+
217 auto s = tile::internal::getTrmmSizes(side, a, b);
+
218 common::internal::SingleThreadedBlasScope single;
+
219 blas::trmm(blas::Layout::ColMajor, side, uplo, op, diag, s.m, s.n, alpha, a.ptr(), a.ld(), b.ptr(),
+
220 b.ld());
+
221}
+
222
+
223// Triangular matrix-matrix multiplication.
+
224// Version with 3 tile arguments (different output tile).
+
225template <class T>
+
226void trmm3(const blas::Side side, const blas::Uplo uplo, const blas::Op op, const blas::Diag diag,
+
227 const T alpha, const Tile<const T, Device::CPU>& a, const Tile<const T, Device::CPU>& b,
+
228 const Tile<T, Device::CPU>& c) noexcept {
+
229 auto s = tile::internal::getTrmm3Sizes(side, a, b, c);
+
230 DLAF_ASSERT(b.ptr() == nullptr || b.ptr() != c.ptr(), b.ptr(), c.ptr());
+
231
+
232 matrix::internal::copy(b, c);
+
233 common::internal::SingleThreadedBlasScope single;
+
234 blas::trmm(blas::Layout::ColMajor, side, uplo, op, diag, s.m, s.n, alpha, a.ptr(), a.ld(), c.ptr(),
+
235 c.ld());
+
236}
+
237
+
238template <class T>
+
239void trsm(const blas::Side side, const blas::Uplo uplo, const blas::Op op, const blas::Diag diag,
+
240 const T alpha, const Tile<const T, Device::CPU>& a, const Tile<T, Device::CPU>& b) noexcept {
+
241 auto s = tile::internal::getTrsmSizes(side, a, b);
+
242 common::internal::SingleThreadedBlasScope single;
+
243 blas::trsm(blas::Layout::ColMajor, side, uplo, op, diag, s.m, s.n, alpha, a.ptr(), a.ld(), b.ptr(),
+
244 b.ld());
+
245}
+
246
+
247#ifdef DLAF_WITH_GPU
248
-
251template <Backend B>
-
252auto trmm(const dlaf::internal::Policy<B>& p);
-
253
-
258template <Backend B, class T, Device D>
-
259void trmm3(const dlaf::internal::Policy<B>& policy, const blas::Side side, const blas::Uplo uplo,
-
260 const blas::Op op, const blas::Diag diag, const T alpha, const Tile<const T, D>& a,
-
261 const Tile<const T, D>& b, const Tile<T, D>& c);
+
249template <class T>
+
250void gemm(cublasHandle_t handle, const blas::Op op_a, const blas::Op op_b, const T alpha,
+
251 const matrix::Tile<const T, Device::GPU>& a, const matrix::Tile<const T, Device::GPU>& b,
+
252 const T beta, const matrix::Tile<T, Device::GPU>& c) {
+
253 using util::blasToCublas;
+
254 using util::blasToCublasCast;
+
255 auto s = getGemmSizes(op_a, op_b, a, b, c);
+
256 gpublas::internal::Gemm<T>::call(handle, blasToCublas(op_a), blasToCublas(op_b), to_int(s.m),
+
257 to_int(s.n), to_int(s.k), blasToCublasCast(&alpha),
+
258 blasToCublasCast(a.ptr()), to_int(a.ld()), blasToCublasCast(b.ptr()),
+
259 to_int(b.ld()), blasToCublasCast(&beta), blasToCublasCast(c.ptr()),
+
260 to_int(c.ld()));
+
261}
262
-
265template <Backend B, typename Sender,
-
266 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
-
267auto trmm3(const dlaf::internal::Policy<B>& p, Sender&& s);
-
268
-
271template <Backend B>
-
272auto trmm3(const dlaf::internal::Policy<B>& p);
-
273
-
277template <Backend B, class T, Device D>
-
278void trsm(const dlaf::internal::Policy<B>& policy, const blas::Side side, const blas::Uplo uplo,
-
279 const blas::Op op, const blas::Diag diag, const T alpha, const Tile<const T, D>& a,
-
280 const Tile<T, D>& b);
-
281
-
284template <Backend B, typename Sender,
-
285 typename = std::enable_if_t<pika::execution::experimental::is_sender_v<Sender>>>
-
286auto trsm(const dlaf::internal::Policy<B>& p, Sender&& s);
-
287
-
290template <Backend B>
-
291auto trsm(const dlaf::internal::Policy<B>& p);
-
292#else
-
293
-
294namespace internal {
-
295
-
296template <class T>
-
297void gemm(const blas::Op op_a, const blas::Op op_b, const T alpha, const Tile<const T, Device::CPU>& a,
-
298 const Tile<const T, Device::CPU>& b, const T beta, const Tile<T, Device::CPU>& c) noexcept {
-
299 auto s = tile::internal::getGemmSizes(op_a, op_b, a, b, c);
-
300 common::internal::SingleThreadedBlasScope single;
-
301 blas::gemm(blas::Layout::ColMajor, op_a, op_b, s.m, s.n, s.k, alpha, a.ptr(), a.ld(), b.ptr(), b.ld(),
-
302 beta, c.ptr(), c.ld());
-
303}
-
304
-
305template <class T>
-
306void hemm(const blas::Side side, const blas::Uplo uplo, const T alpha,
-
307 const Tile<const T, Device::CPU>& a, const Tile<const T, Device::CPU>& b, const T beta,
-
308 const Tile<T, Device::CPU>& c) {
-
309 auto s = tile::internal::getHemmSizes(side, a, b, c);
-
310 common::internal::SingleThreadedBlasScope single;
-
311 blas::hemm(blas::Layout::ColMajor, side, uplo, s.m, s.n, alpha, a.ptr(), a.ld(), b.ptr(), b.ld(), beta,
-
312 c.ptr(), c.ld());
-
313}
-
314
-
315template <class T>
-
316void her2k(const blas::Uplo uplo, const blas::Op op, const T alpha, const Tile<const T, Device::CPU>& a,
-
317 const Tile<const T, Device::CPU>& b, const BaseType<T> beta,
-
318 const Tile<T, Device::CPU>& c) noexcept {
-
319 auto s = tile::internal::getHer2kSizes(op, a, b, c);
-
320 common::internal::SingleThreadedBlasScope single;
-
321 blas::her2k(blas::Layout::ColMajor, uplo, op, s.n, s.k, alpha, a.ptr(), a.ld(), b.ptr(), b.ld(), beta,
-
322 c.ptr(), c.ld());
-
323}
-
324
-
325template <class T>
-
326void herk(const blas::Uplo uplo, const blas::Op op, const BaseType<T> alpha,
-
327 const Tile<const T, Device::CPU>& a, const BaseType<T> beta,
-
328 const Tile<T, Device::CPU>& c) noexcept {
-
329 auto s = tile::internal::getHerkSizes(op, a, c);
-
330 common::internal::SingleThreadedBlasScope single;
-
331 blas::herk(blas::Layout::ColMajor, uplo, op, s.n, s.k, alpha, a.ptr(), a.ld(), beta, c.ptr(), c.ld());
-
332}
-
333
-
334// Triangular matrix-matrix multiplication.
-
335template <class T>
-
336void trmm(const blas::Side side, const blas::Uplo uplo, const blas::Op op, const blas::Diag diag,
-
337 const T alpha, const Tile<const T, Device::CPU>& a, const Tile<T, Device::CPU>& b) noexcept {
-
338 auto s = tile::internal::getTrmmSizes(side, a, b);
-
339 common::internal::SingleThreadedBlasScope single;
-
340 blas::trmm(blas::Layout::ColMajor, side, uplo, op, diag, s.m, s.n, alpha, a.ptr(), a.ld(), b.ptr(),
-
341 b.ld());
-
342}
-
343
-
344// Triangular matrix-matrix multiplication.
-
345// Version with 3 tile arguments (different output tile).
-
346template <class T>
-
347void trmm3(const blas::Side side, const blas::Uplo uplo, const blas::Op op, const blas::Diag diag,
-
348 const T alpha, const Tile<const T, Device::CPU>& a, const Tile<const T, Device::CPU>& b,
-
349 const Tile<T, Device::CPU>& c) noexcept {
-
350 auto s = tile::internal::getTrmm3Sizes(side, a, b, c);
-
351 DLAF_ASSERT(b.ptr() == nullptr || b.ptr() != c.ptr(), b.ptr(), c.ptr());
-
352
-
353 matrix::internal::copy(b, c);
-
354 common::internal::SingleThreadedBlasScope single;
-
355 blas::trmm(blas::Layout::ColMajor, side, uplo, op, diag, s.m, s.n, alpha, a.ptr(), a.ld(), c.ptr(),
-
356 c.ld());
-
357}
-
358
-
359template <class T>
-
360void trsm(const blas::Side side, const blas::Uplo uplo, const blas::Op op, const blas::Diag diag,
-
361 const T alpha, const Tile<const T, Device::CPU>& a, const Tile<T, Device::CPU>& b) noexcept {
-
362 auto s = tile::internal::getTrsmSizes(side, a, b);
-
363 common::internal::SingleThreadedBlasScope single;
-
364 blas::trsm(blas::Layout::ColMajor, side, uplo, op, diag, s.m, s.n, alpha, a.ptr(), a.ld(), b.ptr(),
-
365 b.ld());
-
366}
-
367
-
368#ifdef DLAF_WITH_GPU
-
369
-
370template <class T>
-
371void gemm(cublasHandle_t handle, const blas::Op op_a, const blas::Op op_b, const T alpha,
-
372 const matrix::Tile<const T, Device::GPU>& a, const matrix::Tile<const T, Device::GPU>& b,
-
373 const T beta, const matrix::Tile<T, Device::GPU>& c) {
-
374 using util::blasToCublas;
-
375 using util::blasToCublasCast;
-
376 auto s = getGemmSizes(op_a, op_b, a, b, c);
-
377 gpublas::internal::Gemm<T>::call(handle, blasToCublas(op_a), blasToCublas(op_b), to_int(s.m),
-
378 to_int(s.n), to_int(s.k), blasToCublasCast(&alpha),
-
379 blasToCublasCast(a.ptr()), to_int(a.ld()), blasToCublasCast(b.ptr()),
-
380 to_int(b.ld()), blasToCublasCast(&beta), blasToCublasCast(c.ptr()),
-
381 to_int(c.ld()));
-
382}
-
383
-
384template <class T>
-
385void hemm(cublasHandle_t handle, const blas::Side side, const blas::Uplo uplo, const T alpha,
-
386 const Tile<const T, Device::GPU>& a, const Tile<const T, Device::GPU>& b, const T beta,
-
387 const Tile<T, Device::GPU>& c) {
-
388 using util::blasToCublas;
-
389 using util::blasToCublasCast;
-
390 auto s = getHemmSizes(side, a, b, c);
-
391 gpublas::internal::Hemm<T>::call(handle, blasToCublas(side), blasToCublas(uplo), to_int(s.m),
-
392 to_int(s.n), blasToCublasCast(&alpha), blasToCublasCast(a.ptr()),
-
393 to_int(a.ld()), blasToCublasCast(b.ptr()), to_int(b.ld()),
-
394 blasToCublasCast(&beta), blasToCublasCast(c.ptr()), to_int(c.ld()));
-
395}
-
396
-
397template <class T>
-
398void her2k(cublasHandle_t handle, const blas::Uplo uplo, blas::Op op, const T alpha,
-
399 const matrix::Tile<const T, Device::GPU>& a, const Tile<const T, Device::GPU>& b,
-
400 const BaseType<T> beta, const matrix::Tile<T, Device::GPU>& c) {
-
401 using util::blasToCublas;
-
402 using util::blasToCublasCast;
-
403 auto s = getHer2kSizes(op, a, b, c);
-
404#if defined(DLAF_WITH_HIP) && HIP_VERSION < 50200000
-
405 if (!isComplex_v<T> && op == blas::Op::ConjTrans)
-
406 op = blas::Op::Trans;
-
407#endif
-
408 gpublas::internal::Her2k<T>::call(handle, blasToCublas(uplo), blasToCublas(op), to_int(s.n),
-
409 to_int(s.k), blasToCublasCast(&alpha), blasToCublasCast(a.ptr()),
-
410 to_int(a.ld()), blasToCublasCast(b.ptr()), to_int(b.ld()),
-
411 blasToCublasCast(&beta), blasToCublasCast(c.ptr()), to_int(c.ld()));
-
412}
-
413
-
414template <class T>
-
415void herk(cublasHandle_t handle, const blas::Uplo uplo, const blas::Op op, const BaseType<T> alpha,
-
416 const matrix::Tile<const T, Device::GPU>& a, const BaseType<T> beta,
-
417 const matrix::Tile<T, Device::GPU>& c) {
-
418 using util::blasToCublas;
-
419 using util::blasToCublasCast;
-
420 auto s = getHerkSizes(op, a, c);
-
421 gpublas::internal::Herk<T>::call(handle, blasToCublas(uplo), blasToCublas(op), to_int(s.n),
-
422 to_int(s.k), blasToCublasCast(&alpha), blasToCublasCast(a.ptr()),
-
423 to_int(a.ld()), blasToCublasCast(&beta), blasToCublasCast(c.ptr()),
-
424 to_int(c.ld()));
-
425}
-
426
-
427template <class T>
-
428void trmm(cublasHandle_t handle, const blas::Side side, const blas::Uplo uplo, const blas::Op op,
-
429 const blas::Diag diag, const T alpha, const matrix::Tile<const T, Device::GPU>& a,
-
430 const matrix::Tile<T, Device::GPU>& b) {
-
431 using util::blasToCublas;
-
432 using util::blasToCublasCast;
-
433 auto s = tile::internal::getTrmmSizes(side, a, b);
-
434
-
435 gpublas::internal::Trmm<T>::call(handle, blasToCublas(side), blasToCublas(uplo), blasToCublas(op),
-
436 blasToCublas(diag), to_int(s.m), to_int(s.n),
-
437 blasToCublasCast(&alpha), blasToCublasCast(a.ptr()), to_int(a.ld()),
-
438 blasToCublasCast(b.ptr()), to_int(b.ld()), blasToCublasCast(b.ptr()),
-
439 to_int(b.ld()));
-
440}
-
441
-
442template <class T>
-
443void trmm3(cublasHandle_t handle, const blas::Side side, const blas::Uplo uplo, const blas::Op op,
-
444 const blas::Diag diag, const T alpha, const matrix::Tile<const T, Device::GPU>& a,
-
445 const matrix::Tile<const T, Device::GPU>& b, const matrix::Tile<T, Device::GPU>& c) {
-
446 using util::blasToCublas;
-
447 using util::blasToCublasCast;
-
448 auto s = tile::internal::getTrmm3Sizes(side, a, b, c);
-
449 DLAF_ASSERT(b.ptr() == nullptr || b.ptr() != c.ptr(), b.ptr(), c.ptr());
-
450
-
451 gpublas::internal::Trmm<T>::call(handle, blasToCublas(side), blasToCublas(uplo), blasToCublas(op),
-
452 blasToCublas(diag), to_int(s.m), to_int(s.n),
-
453 blasToCublasCast(&alpha), blasToCublasCast(a.ptr()), to_int(a.ld()),
-
454 blasToCublasCast(b.ptr()), to_int(b.ld()), blasToCublasCast(c.ptr()),
-
455 to_int(c.ld()));
-
456}
-
457
-
458template <class T>
-
459void trsm(cublasHandle_t handle, const blas::Side side, const blas::Uplo uplo, const blas::Op op,
-
460 const blas::Diag diag, const T alpha, const matrix::Tile<const T, Device::GPU>& a,
-
461 const matrix::Tile<T, Device::GPU>& b) {
-
462 using util::blasToCublas;
-
463 using util::blasToCublasCast;
-
464 auto s = getTrsmSizes(side, a, b);
-
465 auto a_ptr = blasToCublasCast(a.ptr());
-
466 gpublas::internal::Trsm<T>::call(handle, blasToCublas(side), blasToCublas(uplo), blasToCublas(op),
-
467 blasToCublas(diag), to_int(s.m), to_int(s.n),
-
468 blasToCublasCast(&alpha), a_ptr, to_int(a.ld()),
-
469 blasToCublasCast(b.ptr()), to_int(b.ld()));
-
470}
-
471#endif // defined(DLAF_WITH_GPU)
-
472
-
473DLAF_MAKE_CALLABLE_OBJECT(gemm);
-
474DLAF_MAKE_CALLABLE_OBJECT(hemm);
-
475DLAF_MAKE_CALLABLE_OBJECT(her2k);
-
476DLAF_MAKE_CALLABLE_OBJECT(herk);
-
477DLAF_MAKE_CALLABLE_OBJECT(trmm);
-
478DLAF_MAKE_CALLABLE_OBJECT(trmm3);
-
479DLAF_MAKE_CALLABLE_OBJECT(trsm);
-
480}
-
481
-
482DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, gemm, internal::gemm_o)
-
483DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, hemm, internal::hemm_o)
-
484DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, her2k,
-
485 internal::her2k_o)
-
486DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, herk, internal::herk_o)
-
487DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, trmm, internal::trmm_o)
-
488DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, trmm3,
-
489 internal::trmm3_o)
-
490DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, trsm, internal::trsm_o)
-
491
-
492#endif
-
493}
-
494}
+
263template <class T>
+
264void hemm(cublasHandle_t handle, const blas::Side side, const blas::Uplo uplo, const T alpha,
+
265 const Tile<const T, Device::GPU>& a, const Tile<const T, Device::GPU>& b, const T beta,
+
266 const Tile<T, Device::GPU>& c) {
+
267 using util::blasToCublas;
+
268 using util::blasToCublasCast;
+
269 auto s = getHemmSizes(side, a, b, c);
+
270 gpublas::internal::Hemm<T>::call(handle, blasToCublas(side), blasToCublas(uplo), to_int(s.m),
+
271 to_int(s.n), blasToCublasCast(&alpha), blasToCublasCast(a.ptr()),
+
272 to_int(a.ld()), blasToCublasCast(b.ptr()), to_int(b.ld()),
+
273 blasToCublasCast(&beta), blasToCublasCast(c.ptr()), to_int(c.ld()));
+
274}
+
275
+
276template <class T>
+
277void her2k(cublasHandle_t handle, const blas::Uplo uplo, blas::Op op, const T alpha,
+
278 const matrix::Tile<const T, Device::GPU>& a, const Tile<const T, Device::GPU>& b,
+
279 const BaseType<T> beta, const matrix::Tile<T, Device::GPU>& c) {
+
280 using util::blasToCublas;
+
281 using util::blasToCublasCast;
+
282 auto s = getHer2kSizes(op, a, b, c);
+
283#if defined(DLAF_WITH_HIP) && HIP_VERSION < 50200000
+
284 if (!isComplex_v<T> && op == blas::Op::ConjTrans)
+
285 op = blas::Op::Trans;
+
286#endif
+
287 gpublas::internal::Her2k<T>::call(handle, blasToCublas(uplo), blasToCublas(op), to_int(s.n),
+
288 to_int(s.k), blasToCublasCast(&alpha), blasToCublasCast(a.ptr()),
+
289 to_int(a.ld()), blasToCublasCast(b.ptr()), to_int(b.ld()),
+
290 blasToCublasCast(&beta), blasToCublasCast(c.ptr()), to_int(c.ld()));
+
291}
+
292
+
293template <class T>
+
294void herk(cublasHandle_t handle, const blas::Uplo uplo, const blas::Op op, const BaseType<T> alpha,
+
295 const matrix::Tile<const T, Device::GPU>& a, const BaseType<T> beta,
+
296 const matrix::Tile<T, Device::GPU>& c) {
+
297 using util::blasToCublas;
+
298 using util::blasToCublasCast;
+
299 auto s = getHerkSizes(op, a, c);
+
300 gpublas::internal::Herk<T>::call(handle, blasToCublas(uplo), blasToCublas(op), to_int(s.n),
+
301 to_int(s.k), blasToCublasCast(&alpha), blasToCublasCast(a.ptr()),
+
302 to_int(a.ld()), blasToCublasCast(&beta), blasToCublasCast(c.ptr()),
+
303 to_int(c.ld()));
+
304}
+
305
+
306template <class T>
+
307void trmm(cublasHandle_t handle, const blas::Side side, const blas::Uplo uplo, const blas::Op op,
+
308 const blas::Diag diag, const T alpha, const matrix::Tile<const T, Device::GPU>& a,
+
309 const matrix::Tile<T, Device::GPU>& b) {
+
310 using util::blasToCublas;
+
311 using util::blasToCublasCast;
+
312 auto s = tile::internal::getTrmmSizes(side, a, b);
+
313
+
314 gpublas::internal::Trmm<T>::call(handle, blasToCublas(side), blasToCublas(uplo), blasToCublas(op),
+
315 blasToCublas(diag), to_int(s.m), to_int(s.n),
+
316 blasToCublasCast(&alpha), blasToCublasCast(a.ptr()), to_int(a.ld()),
+
317 blasToCublasCast(b.ptr()), to_int(b.ld()), blasToCublasCast(b.ptr()),
+
318 to_int(b.ld()));
+
319}
+
320
+
321template <class T>
+
322void trmm3(cublasHandle_t handle, const blas::Side side, const blas::Uplo uplo, const blas::Op op,
+
323 const blas::Diag diag, const T alpha, const matrix::Tile<const T, Device::GPU>& a,
+
324 const matrix::Tile<const T, Device::GPU>& b, const matrix::Tile<T, Device::GPU>& c) {
+
325 using util::blasToCublas;
+
326 using util::blasToCublasCast;
+
327 auto s = tile::internal::getTrmm3Sizes(side, a, b, c);
+
328 DLAF_ASSERT(b.ptr() == nullptr || b.ptr() != c.ptr(), b.ptr(), c.ptr());
+
329
+
330 gpublas::internal::Trmm<T>::call(handle, blasToCublas(side), blasToCublas(uplo), blasToCublas(op),
+
331 blasToCublas(diag), to_int(s.m), to_int(s.n),
+
332 blasToCublasCast(&alpha), blasToCublasCast(a.ptr()), to_int(a.ld()),
+
333 blasToCublasCast(b.ptr()), to_int(b.ld()), blasToCublasCast(c.ptr()),
+
334 to_int(c.ld()));
+
335}
+
336
+
337template <class T>
+
338void trsm(cublasHandle_t handle, const blas::Side side, const blas::Uplo uplo, const blas::Op op,
+
339 const blas::Diag diag, const T alpha, const matrix::Tile<const T, Device::GPU>& a,
+
340 const matrix::Tile<T, Device::GPU>& b) {
+
341 using util::blasToCublas;
+
342 using util::blasToCublasCast;
+
343 auto s = getTrsmSizes(side, a, b);
+
344 auto a_ptr = blasToCublasCast(a.ptr());
+
345 gpublas::internal::Trsm<T>::call(handle, blasToCublas(side), blasToCublas(uplo), blasToCublas(op),
+
346 blasToCublas(diag), to_int(s.m), to_int(s.n),
+
347 blasToCublasCast(&alpha), a_ptr, to_int(a.ld()),
+
348 blasToCublasCast(b.ptr()), to_int(b.ld()));
+
349}
+
350#endif // defined(DLAF_WITH_GPU)
+
351
+
352DLAF_MAKE_CALLABLE_OBJECT(gemm);
+
353DLAF_MAKE_CALLABLE_OBJECT(hemm);
+
354DLAF_MAKE_CALLABLE_OBJECT(her2k);
+
355DLAF_MAKE_CALLABLE_OBJECT(herk);
+
356DLAF_MAKE_CALLABLE_OBJECT(trmm);
+
357DLAF_MAKE_CALLABLE_OBJECT(trmm3);
+
358DLAF_MAKE_CALLABLE_OBJECT(trsm);
+
359}
+
360
+
361DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, gemm, internal::gemm_o)
+
362DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, hemm, internal::hemm_o)
+
363DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, her2k,
+
364 internal::her2k_o)
+
365DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, herk, internal::herk_o)
+
366DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, trmm, internal::trmm_o)
+
367DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, trmm3,
+
368 internal::trmm3_o)
+
369DLAF_MAKE_SENDER_ALGORITHM_OVERLOADS(dlaf::internal::TransformDispatchType::Blas, trsm, internal::trsm_o)
+
370
+
371#endif
+
372}
+
373}
dlaf::tile::gemm
void gemm(const blas::Op op_a, const blas::Op op_b, const T alpha, const Tile< const T, D > &a, const Tile< const T, D > &b, const T beta, const Tile< T, D > &c)
dlaf::tile::trsm
void trsm(const dlaf::internal::Policy< B > &policy, const blas::Side side, const blas::Uplo uplo, const blas::Op op, const blas::Diag diag, const T alpha, const Tile< const T, D > &a, const Tile< T, D > &b)
dlaf::tile::her2k
void her2k(const blas::Uplo uplo, const blas::Op op, const T alpha, const Tile< const T, D > &a, const Tile< const T, D > &b, const BaseType< T > beta, const Tile< T, D > &c)
@@ -535,8 +414,8 @@
dlaf::common::internal::SingleThreadedBlasScope
Definition single_threaded_blas.h:20
dlaf::internal::Policy
Definition policy.h:24
dlaf::matrix::Tile
Definition tile.h:332
-
dlaf::memory::MemoryView
Definition memory_view.h:32
copy_tile.h
+
gpublas.h
tile.h
single_threaded_blas.h
types.h
diff --git a/master/dir_2e3e3bc658385778082583f320919a4c.html b/master/dir_2e3e3bc658385778082583f320919a4c.html index 4160fe3484..174779ed7d 100644 --- a/master/dir_2e3e3bc658385778082583f320919a4c.html +++ b/master/dir_2e3e3bc658385778082583f320919a4c.html @@ -83,6 +83,8 @@    lacpy.h   + larft.h +   laset.h   diff --git a/master/dir_5ca20fc8e29b0f8739133582ef745158.html b/master/dir_5ca20fc8e29b0f8739133582ef745158.html index 496865bca8..411ec93637 100644 --- a/master/dir_5ca20fc8e29b0f8739133582ef745158.html +++ b/master/dir_5ca20fc8e29b0f8739133582ef745158.html @@ -83,6 +83,8 @@    error.h   + gpublas.h +  diff --git a/master/factorization_2qr_2api_8h_source.html b/master/factorization_2qr_2api_8h_source.html index b698df830c..2fc65e3e6c 100644 --- a/master/factorization_2qr_2api_8h_source.html +++ b/master/factorization_2qr_2api_8h_source.html @@ -142,7 +142,7 @@
tile.h
panel.h
dlaf::factorization::internal::QR_Tfactor
Definition api.h:27
-
dlaf::factorization::internal::QR_Tfactor::call
static void call(matrix::Panel< Coord::Col, T, device > &panel_view, matrix::ReadOnlyTileSender< T, Device::CPU > taus, matrix::ReadWriteTileSender< T, device > t)
Definition t_factor_impl.h:260
+
dlaf::factorization::internal::QR_Tfactor::call
static void call(matrix::Panel< Coord::Col, T, device > &panel_view, matrix::ReadOnlyTileSender< T, Device::CPU > taus, matrix::ReadWriteTileSender< T, device > t)
Definition t_factor_impl.h:248
dlaf::factorization::internal::QR
Definition api.h:24
types.h
views.h
diff --git a/master/files.html b/master/files.html index 156026ba95..40ec64de49 100644 --- a/master/files.html +++ b/master/files.html @@ -199,6 +199,7 @@  ► blas  api.h  error.h + gpublas.h  ► cublas  error.h  ► cusolver @@ -217,7 +218,8 @@  ► gpu  add.h  lacpy.h - laset.h + larft.h + laset.h  enum_output.h  tile.h  ► matrix diff --git a/master/gpublas_8h.html b/master/gpublas_8h.html new file mode 100644 index 0000000000..620d9e1337 --- /dev/null +++ b/master/gpublas_8h.html @@ -0,0 +1,152 @@ + + + + + + + +DLAF: /home/runner/work/DLA-Future/DLA-Future/include/dlaf/gpu/blas/gpublas.h File Reference + + + + + + + + + +
+
+ + + + + + +
+
DLAF +
+
+
+ + + + + + + + +
+
+ + +
+
+
+
+
+
Loading...
+
Searching...
+
No Matches
+
+
+
+
+ + +
+
+
+Macros | +Functions
+
gpublas.h File Reference
+
+
+
#include <cstddef>
+#include <utility>
+#include <whip.hpp>
+#include <dlaf/gpu/blas/api.h>
+#include <dlaf/gpu/blas/error.h>
+#include <dlaf/util_cublas.h>
+
+

Go to the source code of this file.

+ + + + +

+Macros

#define DLAF_DECLARE_GPUBLAS_OP(Name)
 
+ + + + + + + + + + + + + + + + + +

+Functions

+ dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_OP (Axpy, axpy)
 
+ dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_OP (Gemv, gemv)
 
+ dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_OP (Trmv, trmv)
 
+ dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_OP (Gemm, gemm)
 
+ dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_SYHE_OP (Hemm, mm)
 
+ dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_SYHE_OP (Her2k, r2k)
 
+ dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_SYHE_OP (Herk, rk)
 
+ dlaf::gpublas::internal::DLAF_MAKE_GPUBLAS_OP (Trsm, trsm)
 
+

Detailed Description

+

Provides gpublas wrappers for BLAS operations.

+

Macro Definition Documentation

+ +

◆ DLAF_DECLARE_GPUBLAS_OP

+ +
+
+ + + + + + + + +
#define DLAF_DECLARE_GPUBLAS_OP( Name)
+
+Value:
template <typename T> \
+
struct Name
+
+
+
+
+ +
+Generated by doxygen 1.9.8 +
+ + diff --git a/master/gpublas_8h_source.html b/master/gpublas_8h_source.html new file mode 100644 index 0000000000..42128c0ac5 --- /dev/null +++ b/master/gpublas_8h_source.html @@ -0,0 +1,237 @@ + + + + + + + +DLAF: /home/runner/work/DLA-Future/DLA-Future/include/dlaf/gpu/blas/gpublas.h Source File + + + + + + + + + +
+
+ + + + + + +
+
DLAF +
+
+
+ + + + + + + + + +
+
+ + +
+
+
+
+
+
Loading...
+
Searching...
+
No Matches
+
+
+
+
+ + +
+
+
gpublas.h
+
+
+Go to the documentation of this file.
1//
+
2// Distributed Linear Algebra with Future (DLAF)
+
3//
+
4// Copyright (c) 2018-2024, ETH Zurich
+
5// All rights reserved.
+
6//
+
7// Please, refer to the LICENSE file in the root directory.
+
8// SPDX-License-Identifier: BSD-3-Clause
+
9//
+
10#pragma once
+
11
+
14
+
15#ifdef DLAF_WITH_GPU
+
16#include <cstddef>
+
17#include <utility>
+
18
+
19#include <whip.hpp>
+
20
+
21#include <dlaf/gpu/blas/api.h>
+
22#include <dlaf/gpu/blas/error.h>
+
23#include <dlaf/util_cublas.h>
+
24
+
25#ifdef DLAF_WITH_HIP
+
26
+
27#include <pika/async_cuda/detail/cuda_event_callback.hpp>
+
28
+
29#include <dlaf/memory/memory_view.h>
+
30
+
31#define DLAF_GET_ROCBLAS_WORKSPACE(f) \
+
32 [&]() { \
+
33 std::size_t workspace_size; \
+
34 DLAF_GPUBLAS_CHECK_ERROR( \
+
35 rocblas_start_device_memory_size_query(static_cast<rocblas_handle>(handle))); \
+
36 DLAF_ROCBLAS_WORKSPACE_CHECK_ERROR(rocblas_##f(handle, std::forward<Args>(args)...)); \
+
37 DLAF_GPUBLAS_CHECK_ERROR(rocblas_stop_device_memory_size_query(static_cast<rocblas_handle>(handle), \
+
38 &workspace_size)); \
+
39 return ::dlaf::memory::MemoryView<std::byte, Device::GPU>(to_int(workspace_size)); \
+
40 }();
+
41
+
42namespace dlaf::tile::internal {
+
43inline void extendROCBlasWorkspace(cublasHandle_t handle,
+
44 ::dlaf::memory::MemoryView<std::byte, Device::GPU>&& workspace) {
+
45 whip::stream_t stream;
+
46 DLAF_GPUBLAS_CHECK_ERROR(cublasGetStream(handle, &stream));
+
47 auto f = [workspace = std::move(workspace)](whip::error_t status) { whip::check_error(status); };
+
48 pika::cuda::experimental::detail::add_event_callback(std::move(f), stream);
+
49}
+
50}
+
51
+
52#define DLAF_DEFINE_GPUBLAS_OP(Name, Type, f) \
+
53 template <> \
+
54 struct Name<Type> { \
+
55 template <typename... Args> \
+
56 static void call(cublasHandle_t handle, Args&&... args) { \
+
57 auto workspace = DLAF_GET_ROCBLAS_WORKSPACE(f); \
+
58 DLAF_GPUBLAS_CHECK_ERROR(rocblas_set_workspace(static_cast<rocblas_handle>(handle), workspace(), \
+
59 to_sizet(workspace.size()))); \
+
60 DLAF_GPUBLAS_CHECK_ERROR(rocblas_##f(handle, std::forward<Args>(args)...)); \
+
61 DLAF_GPUBLAS_CHECK_ERROR(rocblas_set_workspace(static_cast<rocblas_handle>(handle), nullptr, 0)); \
+
62 ::dlaf::tile::internal::extendROCBlasWorkspace(handle, std::move(workspace)); \
+
63 } \
+
64 }
+
65
+
66#elif defined(DLAF_WITH_CUDA)
+
67
+
68#define DLAF_DEFINE_GPUBLAS_OP(Name, Type, f) \
+
69 template <> \
+
70 struct Name<Type> { \
+
71 template <typename... Args> \
+
72 static void call(Args&&... args) { \
+
73 DLAF_GPUBLAS_CHECK_ERROR(cublas##f##_v2(std::forward<Args>(args)...)); \
+
74 } \
+
75 }
+
76
+
77#endif
+
78
+
79#define DLAF_DECLARE_GPUBLAS_OP(Name) \
+
80 template <typename T> \
+
81 struct Name
+
82
+
83#ifdef DLAF_WITH_HIP
+
84#define DLAF_MAKE_GPUBLAS_OP(Name, f) \
+
85 DLAF_DECLARE_GPUBLAS_OP(Name); \
+
86 DLAF_DEFINE_GPUBLAS_OP(Name, float, s##f); \
+
87 DLAF_DEFINE_GPUBLAS_OP(Name, double, d##f); \
+
88 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<float>, c##f); \
+
89 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<double>, z##f)
+
90
+
91#define DLAF_MAKE_GPUBLAS_SYHE_OP(Name, f) \
+
92 DLAF_DECLARE_GPUBLAS_OP(Name); \
+
93 DLAF_DEFINE_GPUBLAS_OP(Name, float, ssy##f); \
+
94 DLAF_DEFINE_GPUBLAS_OP(Name, double, dsy##f); \
+
95 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<float>, che##f); \
+
96 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<double>, zhe##f)
+
97
+
98#elif defined(DLAF_WITH_CUDA)
+
99#define DLAF_MAKE_GPUBLAS_OP(Name, f) \
+
100 DLAF_DECLARE_GPUBLAS_OP(Name); \
+
101 DLAF_DEFINE_GPUBLAS_OP(Name, float, S##f); \
+
102 DLAF_DEFINE_GPUBLAS_OP(Name, double, D##f); \
+
103 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<float>, C##f); \
+
104 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<double>, Z##f)
+
105
+
106#define DLAF_MAKE_GPUBLAS_SYHE_OP(Name, f) \
+
107 DLAF_DECLARE_GPUBLAS_OP(Name); \
+
108 DLAF_DEFINE_GPUBLAS_OP(Name, float, Ssy##f); \
+
109 DLAF_DEFINE_GPUBLAS_OP(Name, double, Dsy##f); \
+
110 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<float>, Che##f); \
+
111 DLAF_DEFINE_GPUBLAS_OP(Name, std::complex<double>, Zhe##f)
+
112#endif
+
113
+
114namespace dlaf::gpublas::internal {
+
115
+
116// Level 1
+
117DLAF_MAKE_GPUBLAS_OP(Axpy, axpy);
+
118
+
119// Level 2
+
120DLAF_MAKE_GPUBLAS_OP(Gemv, gemv);
+
121
+
122DLAF_MAKE_GPUBLAS_OP(Trmv, trmv);
+
123
+
124// Level 3
+
125DLAF_MAKE_GPUBLAS_OP(Gemm, gemm);
+
126
+
127DLAF_MAKE_GPUBLAS_SYHE_OP(Hemm, mm);
+
128
+
129DLAF_MAKE_GPUBLAS_SYHE_OP(Her2k, r2k);
+
130
+
131DLAF_MAKE_GPUBLAS_SYHE_OP(Herk, rk);
+
132
+
133#if defined(DLAF_WITH_CUDA)
+
134DLAF_MAKE_GPUBLAS_OP(Trmm, trmm);
+
135#elif defined(DLAF_WITH_HIP)
+
136
+
137#if ROCBLAS_VERSION_MAJOR >= 3 && defined(ROCBLAS_V3)
+
138DLAF_MAKE_GPUBLAS_OP(Trmm, trmm);
+
139#else
+
140DLAF_MAKE_GPUBLAS_OP(Trmm, trmm_outofplace);
+
141#endif
+
142
+
143#endif
+
144
+
145DLAF_MAKE_GPUBLAS_OP(Trsm, trsm);
+
146}
+
147#endif
+
dlaf::memory::MemoryView
Definition memory_view.h:32
+
memory_view.h
+
+ +
+Generated by doxygen 1.9.8 +
+ + diff --git a/master/larft_8h_source.html b/master/larft_8h_source.html new file mode 100644 index 0000000000..bf7e585559 --- /dev/null +++ b/master/larft_8h_source.html @@ -0,0 +1,142 @@ + + + + + + + +DLAF: /home/runner/work/DLA-Future/DLA-Future/include/dlaf/lapack/gpu/larft.h Source File + + + + + + + + + +
+
+ + + + + + +
+
DLAF +
+
+
+ + + + + + + + + +
+
+ + +
+
+
+
+
+
Loading...
+
Searching...
+
No Matches
+
+
+
+
+ + +
+
+
larft.h
+
+
+
1//
+
2// Distributed Linear Algebra with Future (DLAF)
+
3//
+
4// Copyright (c) 2018-2024, ETH Zurich
+
5// All rights reserved.
+
6//
+
7// Please, refer to the LICENSE file in the root directory.
+
8// SPDX-License-Identifier: BSD-3-Clause
+
9//
+
10
+
11#pragma once
+
12
+
13#ifdef DLAF_WITH_GPU
+
14
+
15#include <blas.hh>
+
16#include <whip.hpp>
+
17
+
18#include <dlaf/gpu/blas/api.h>
+
19#include <dlaf/types.h>
+
20
+
21namespace dlaf::gpulapack {
+
22
+
23template <class T>
+
24void larft_gemv0(cublasHandle_t handle, const SizeType m, SizeType k, const T* v, const SizeType ldv,
+
25 const T* tau, T* t, const SizeType ldt);
+
26
+
27template <class T>
+
28void larft_gemv1_notau(cublasHandle_t handle, const SizeType m, const SizeType k, const T* v,
+
29 const SizeType ldv, T* t, const SizeType ldt);
+
30
+
31template <class T>
+
32void larft_gemv1_fixtau(const SizeType k, const T* tau, const SizeType inctau, T* t, const SizeType ldt,
+
33 whip::stream_t stream);
+
34
+
35#define DLAF_CUBLAS_LARFT_GEMV_ETI(kword, Type) \
+
36 kword template void larft_gemv0(cublasHandle_t handle, const SizeType n, SizeType k, const Type* v, \
+
37 const SizeType ldv, const Type* tau, Type* t, const SizeType ldt); \
+
38 kword template void larft_gemv1_notau(cublasHandle_t handle, const SizeType m, const SizeType k, \
+
39 const Type* v, const SizeType ldv, Type* t, \
+
40 const SizeType ldt); \
+
41 kword template void larft_gemv1_fixtau(const SizeType k, const Type* tau, const SizeType inctau, \
+
42 Type* t, const SizeType ldt, whip::stream_t stream)
+
43
+
44DLAF_CUBLAS_LARFT_GEMV_ETI(extern, float);
+
45DLAF_CUBLAS_LARFT_GEMV_ETI(extern, double);
+
46DLAF_CUBLAS_LARFT_GEMV_ETI(extern, std::complex<float>);
+
47DLAF_CUBLAS_LARFT_GEMV_ETI(extern, std::complex<double>);
+
48
+
49}
+
50
+
51#endif
+ +
+ +
+Generated by doxygen 1.9.8 +
+ + diff --git a/master/qr_8h_source.html b/master/qr_8h_source.html index 1575836033..354051df01 100644 --- a/master/qr_8h_source.html +++ b/master/qr_8h_source.html @@ -126,7 +126,7 @@
communicator_pipeline.h
index.h
tile.h
-
dlaf::factorization::internal::QR_Tfactor::call
static void call(matrix::Panel< Coord::Col, T, device > &panel_view, matrix::ReadOnlyTileSender< T, Device::CPU > taus, matrix::ReadWriteTileSender< T, device > t)
Definition t_factor_impl.h:260
+
dlaf::factorization::internal::QR_Tfactor::call
static void call(matrix::Panel< Coord::Col, T, device > &panel_view, matrix::ReadOnlyTileSender< T, Device::CPU > taus, matrix::ReadWriteTileSender< T, device > t)
Definition t_factor_impl.h:248
diff --git a/master/search/all_6.js b/master/search/all_6.js index 62947929de..c6a8ba989e 100644 --- a/master/search/all_6.js +++ b/master/search/all_6.js @@ -11,10 +11,10 @@ var searchData= ['generalmatrix_8',['generalmatrix',['../multiplication_2general_8h.html#a834bc28949a9b0aa7000f8f52822e7a4',1,'dlaf::multiplication::internal::generalMatrix(const blas::Op opA, const blas::Op opB, const T alpha, MatrixRef< const T, D > &mat_a, MatrixRef< const T, D > &mat_b, const T beta, MatrixRef< T, D > &mat_c)'],['../multiplication_2general_8h.html#a475d73bf4d853e54c47e93b495904581',1,'dlaf::multiplication::internal::generalMatrix(comm::CommunicatorPipeline< comm::CommunicatorType::Row > &row_task_chain, comm::CommunicatorPipeline< comm::CommunicatorType::Col > &col_task_chain, const T alpha, MatrixRef< const T, D > &mat_a, MatrixRef< const T, D > &mat_b, const T beta, MatrixRef< T, D > &mat_c)']]], ['generic_5finteger_5fop_9',['generic_integer_op',['../util__math_8h.html#a4a1c1045d16c92d7fef9c553ae8e6c90',1,'dlaf::util::internal']]], ['gentostd_10',['GenToStd',['../structdlaf_1_1eigensolver_1_1internal_1_1_gen_to_std.html',1,'dlaf::eigensolver::internal']]], - ['get_11',['get',['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01const_01_t_00_01_source_target_00_01_source_target_01_4.html#ac9deda6f1d85a6c884deb18255530ae1',1,'dlaf::matrix::MatrixMirror< const T, SourceTarget, SourceTarget >::get()'],['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01_t_00_01_source_target_00_01_source_target_01_4.html#a655dc5435eeba0a3dbad440e7079c193',1,'dlaf::matrix::MatrixMirror< T, SourceTarget, SourceTarget >::get()'],['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01const_01_t_00_01_target_00_01_source_01_4.html#a9bab6d25f671bad8086967f7d395cf39',1,'dlaf::matrix::MatrixMirror< const T, Target, Source >::get()'],['../classdlaf_1_1matrix_1_1_matrix_mirror.html#a0adf7757a3236ff337c5c9ee9ef209eb',1,'dlaf::matrix::MatrixMirror::get()'],['../classdlaf_1_1common_1_1internal_1_1basic__coords.html#a23701a66458fad55204f90b84674a8e2',1,'dlaf::common::internal::basic_coords::get()']]], + ['get_11',['get',['../classdlaf_1_1common_1_1internal_1_1basic__coords.html#a23701a66458fad55204f90b84674a8e2',1,'dlaf::common::internal::basic_coords::get()'],['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01const_01_t_00_01_source_target_00_01_source_target_01_4.html#ac9deda6f1d85a6c884deb18255530ae1',1,'dlaf::matrix::MatrixMirror< const T, SourceTarget, SourceTarget >::get()'],['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01_t_00_01_source_target_00_01_source_target_01_4.html#a655dc5435eeba0a3dbad440e7079c193',1,'dlaf::matrix::MatrixMirror< T, SourceTarget, SourceTarget >::get()'],['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01const_01_t_00_01_target_00_01_source_01_4.html#a9bab6d25f671bad8086967f7d395cf39',1,'dlaf::matrix::MatrixMirror< const T, Target, Source >::get()'],['../classdlaf_1_1matrix_1_1_matrix_mirror.html#a0adf7757a3236ff337c5c9ee9ef209eb',1,'dlaf::matrix::MatrixMirror::get()']]], ['getheight_12',['getHeight',['../structdlaf_1_1matrix_1_1_panel_3_01axis_00_01const_01_t_00_01_d_00_01_store_transposed_1_1_no_01_4.html#ad59bcb70b44cf2a4a61119d8ce4eb5db',1,'dlaf::matrix::Panel< axis, const T, D, StoreTransposed::No >']]], ['getoptionsdescription_13',['getOptionsDescription',['../dlaf_2init_8h.html#a6f4711b407727e4b9317d252b451bd89',1,'dlaf']]], - ['getsource_14',['getsource',['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01const_01_t_00_01_target_00_01_source_01_4.html#a631c46e4f2b01ce329ad887869ba1cff',1,'dlaf::matrix::MatrixMirror< const T, Target, Source >::getSource()'],['../classdlaf_1_1matrix_1_1_matrix_mirror.html#a505b265f6138344924570816271e84a7',1,'dlaf::matrix::MatrixMirror::getSource()'],['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01_t_00_01_source_target_00_01_source_target_01_4.html#a98c6c5777c0fe2c35cd652e4f6b47ab4',1,'dlaf::matrix::MatrixMirror< T, SourceTarget, SourceTarget >::getSource()'],['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01const_01_t_00_01_source_target_00_01_source_target_01_4.html#a91c5417e16e33a3ffa6f7bf9b34a3e4f',1,'dlaf::matrix::MatrixMirror< const T, SourceTarget, SourceTarget >::getSource()']]], + ['getsource_14',['getsource',['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01_t_00_01_source_target_00_01_source_target_01_4.html#a98c6c5777c0fe2c35cd652e4f6b47ab4',1,'dlaf::matrix::MatrixMirror< T, SourceTarget, SourceTarget >::getSource()'],['../classdlaf_1_1matrix_1_1_matrix_mirror.html#a505b265f6138344924570816271e84a7',1,'dlaf::matrix::MatrixMirror::getSource()'],['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01const_01_t_00_01_target_00_01_source_01_4.html#a631c46e4f2b01ce329ad887869ba1cff',1,'dlaf::matrix::MatrixMirror< const T, Target, Source >::getSource()'],['../classdlaf_1_1matrix_1_1_matrix_mirror_3_01const_01_t_00_01_source_target_00_01_source_target_01_4.html#a91c5417e16e33a3ffa6f7bf9b34a3e4f',1,'dlaf::matrix::MatrixMirror< const T, SourceTarget, SourceTarget >::getSource()']]], ['getter_5frandom_15',['getter_random',['../classdlaf_1_1matrix_1_1util_1_1internal_1_1getter__random.html',1,'dlaf::matrix::util::internal']]], ['getter_5frandom_3c_20std_3a_3acomplex_3c_20t_20_3e_20_3e_16',['getter_random< std::complex< T > >',['../classdlaf_1_1matrix_1_1util_1_1internal_1_1getter__random_3_01std_1_1complex_3_01_t_01_4_01_4.html',1,'dlaf::matrix::util::internal']]], ['getwidth_17',['getWidth',['../structdlaf_1_1matrix_1_1_panel_3_01axis_00_01const_01_t_00_01_d_00_01_store_transposed_1_1_no_01_4.html#aa34a2e93ba64d958df32428f3928ccba',1,'dlaf::matrix::Panel< axis, const T, D, StoreTransposed::No >']]], @@ -33,7 +33,8 @@ var searchData= ['global_5ftile_5findex_5ffrom_5fsub_5fdistribution_30',['global_tile_index_from_sub_distribution',['../distribution_8h.html#ad40ebeee4d978542c2a1460c05c1c3fa',1,'dlaf::matrix::internal::distribution']]], ['global_5ftile_5flinear_5findex_31',['global_tile_linear_index',['../distribution_8h.html#a4dbe5cf03961dec5b6b381a4f8415314',1,'dlaf::matrix::internal::distribution']]], ['global_5ftile_5foffset_32',['global_tile_offset',['../distribution__extensions_8h.html#ae8c5ed233fad3336d331da897a56a8a9',1,'dlaf::matrix::internal::distribution']]], - ['grid_2eh_33',['grid.h',['../grid_8h.html',1,'']]], - ['grid_5fordering_34',['grid_ordering',['../grid_8h.html#a438a7328e8679dfd1bc4758e96688d6b',1,'grid.h']]], - ['grid_5fsize_35',['grid_size',['../classdlaf_1_1matrix_1_1internal_1_1_matrix_base.html#aab63e932ea19be17ee87fca99879266c',1,'dlaf::matrix::internal::MatrixBase']]] + ['gpublas_2eh_33',['gpublas.h',['../gpublas_8h.html',1,'']]], + ['grid_2eh_34',['grid.h',['../grid_8h.html',1,'']]], + ['grid_5fordering_35',['grid_ordering',['../grid_8h.html#a438a7328e8679dfd1bc4758e96688d6b',1,'grid.h']]], + ['grid_5fsize_36',['grid_size',['../classdlaf_1_1matrix_1_1internal_1_1_matrix_base.html#aab63e932ea19be17ee87fca99879266c',1,'dlaf::matrix::internal::MatrixBase']]] ]; diff --git a/master/search/files_6.js b/master/search/files_6.js index 16cc0e9ca3..a4aa1a1abb 100644 --- a/master/search/files_6.js +++ b/master/search/files_6.js @@ -3,5 +3,6 @@ var searchData= ['gen_5feigensolver_2eh_0',['gen_eigensolver.h',['../eigensolver_2gen__eigensolver_8h.html',1,'(Global Namespace)'],['../c_2eigensolver_2gen__eigensolver_8h.html',1,'(Global Namespace)']]], ['gen_5fto_5fstd_2eh_1',['gen_to_std.h',['../gen__to__std_8h.html',1,'']]], ['general_2eh_2',['general.h',['../multiplication_2general_8h.html',1,'(Global Namespace)'],['../permutations_2general_8h.html',1,'(Global Namespace)']]], - ['grid_2eh_3',['grid.h',['../grid_8h.html',1,'']]] + ['gpublas_2eh_3',['gpublas.h',['../gpublas_8h.html',1,'']]], + ['grid_2eh_4',['grid.h',['../grid_8h.html',1,'']]] ]; diff --git a/master/t__factor__impl_8h_source.html b/master/t__factor__impl_8h_source.html index 5e5b277867..3c17d10a44 100644 --- a/master/t__factor__impl_8h_source.html +++ b/master/t__factor__impl_8h_source.html @@ -107,345 +107,333 @@
24#include <dlaf/communication/communicator_pipeline.h>
25#include <dlaf/communication/kernels/all_reduce.h>
26#include <dlaf/factorization/qr/api.h>
-
27#include <dlaf/lapack/tile.h>
-
28#include <dlaf/matrix/matrix.h>
-
29#include <dlaf/matrix/tile.h>
-
30#include <dlaf/matrix/views.h>
-
31#include <dlaf/types.h>
-
32#include <dlaf/util_matrix.h>
-
33
-
34#ifdef DLAF_WITH_GPU
-
35#include <whip.hpp>
-
36
-
37#include <dlaf/blas/tile.h>
-
38#endif
-
39
-
40namespace dlaf::factorization::internal {
-
41
-
42namespace tfactor_l {
-
43template <Backend backend, Device device, class T>
-
44struct Helpers {};
-
45
-
46template <class T>
-
-
47struct Helpers<Backend::MC, Device::CPU, T> {
-
48 static auto prepareT(matrix::ReadOnlyTileSender<T, Device::CPU> taus,
-
49 matrix::ReadWriteTileSender<T, Device::CPU> tile_t) {
-
50 namespace ex = pika::execution::experimental;
-
51 namespace di = dlaf::internal;
-
52 return ex::when_all(std::move(taus), std::move(tile_t)) |
-
53 di::transform(di::Policy<Backend::MC>(pika::execution::thread_priority::high),
-
54 [](const matrix::Tile<const T, Device::CPU>& taus,
-
55 matrix::Tile<T, Device::CPU>&& tile_t) {
-
56 tile::internal::set0<T>(tile_t);
-
57
-
58 const SizeType k = tile_t.size().cols();
-
59 lapack::lacpy(blas::Uplo::General, 1, k, taus.ptr(), 1, tile_t.ptr(),
-
60 tile_t.ld() + 1);
-
61
-
62 return std::move(tile_t);
-
63 });
-
64 }
-
65
-
66 static auto stepGEMV(matrix::ReadOnlyTileSender<T, Device::CPU> tile_vi,
-
67 matrix::ReadOnlyTileSender<T, Device::CPU> taus,
-
68 matrix::ReadWriteTileSender<T, Device::CPU> tile_t) {
-
69 auto gemv_func = [](const matrix::Tile<const T, Device::CPU>& tile_v,
-
70 const matrix::Tile<const T, Device::CPU>& taus,
-
71 matrix::Tile<T, Device::CPU> tile_t) noexcept {
-
72 const SizeType k = tile_t.size().cols();
-
73
-
74 DLAF_ASSERT(tile_v.size().cols() == k, tile_v.size().cols(), k);
-
75 DLAF_ASSERT(taus.size().rows() == k, taus.size().rows(), k);
-
76
-
77 common::internal::SingleThreadedBlasScope single;
-
78 for (SizeType j = 0; j < k; ++j) {
-
79 // T(0:j, j) = -tau . V(j:, 0:j)* . V(j:, j)
-
80 // [j x 1] = [(n-j) x j]* . [(n-j) x 1]
-
81 const TileElementIndex t_start{0, j};
-
82 const TileElementIndex va_start{0, 0};
-
83 const TileElementIndex vb_start{0, j};
-
84 const TileElementSize va_size{tile_v.size().rows(), j};
-
85 const T tau = tile_t({j, j});
-
86
-
87 blas::gemv(blas::Layout::ColMajor, blas::Op::ConjTrans, va_size.rows(), va_size.cols(), -tau,
-
88 tile_v.ptr(va_start), tile_v.ld(), tile_v.ptr(vb_start), 1, 1, tile_t.ptr(t_start),
-
89 1);
-
90 }
-
91 return tile_t;
-
92 };
-
93
-
94 namespace ex = pika::execution::experimental;
-
95 namespace di = dlaf::internal;
-
96
-
97 return ex::when_all(tile_vi, std::move(taus), std::move(tile_t)) |
-
98 di::transform(di::Policy<Backend::MC>(pika::execution::thread_priority::high),
-
99 std::move(gemv_func));
-
100 }
-
101
-
102 static void trmvLoop(const matrix::Tile<T, Device::CPU>& tile_t) {
-
103 common::internal::SingleThreadedBlasScope single;
-
104
-
105 const SizeType k = tile_t.size().cols();
-
106
-
107 // Update each column (in order) t = T . t
-
108 // remember that T is upper triangular, so it is possible to use TRMV
-
109 for (SizeType j = 0; j < k; ++j) {
-
110 const TileElementIndex t_start{0, j};
-
111 const TileElementSize t_size{j, 1};
-
112
-
113 blas::trmv(blas::Layout::ColMajor, blas::Uplo::Upper, blas::Op::NoTrans, blas::Diag::NonUnit,
-
114 t_size.rows(), tile_t.ptr(), tile_t.ld(), tile_t.ptr(t_start), 1);
-
115 }
-
116 }
-
117
-
118 static auto stepTRMV(matrix::ReadWriteTileSender<T, Device::CPU> tile_t) noexcept {
-
119 namespace di = dlaf::internal;
-
120
-
121 return std::move(tile_t) |
-
122 di::transform(di::Policy<Backend::MC>(pika::execution::thread_priority::high), trmvLoop);
-
123 }
-
124
-
125 static auto stepCopyDiagAndTRMV(matrix::ReadOnlyTileSender<T, Device::CPU> taus,
-
126 matrix::ReadWriteTileSender<T, Device::CPU> tile_t) noexcept {
-
127 auto tausdiag_trmvloop = [](const matrix::Tile<const T, Device::CPU>& taus,
-
128 matrix::Tile<T, Device::CPU> tile_t) {
-
129 common::internal::SingleThreadedBlasScope single;
-
130
-
131 const SizeType k = tile_t.size().cols();
-
132 lapack::lacpy(blas::Uplo::General, 1, k, taus.ptr(), 1, tile_t.ptr(), tile_t.ld() + 1);
-
133
-
134 trmvLoop(tile_t);
-
135 };
-
136
-
137 namespace di = dlaf::internal;
-
138 namespace ex = pika::execution::experimental;
-
139
-
140 return ex::when_all(std::move(taus), std::move(tile_t)) |
-
141 di::transform(di::Policy<Backend::MC>(pika::execution::thread_priority::high),
-
142 std::move(tausdiag_trmvloop));
-
143 }
-
144};
+
27#include <dlaf/lapack/gpu/larft.h>
+
28#include <dlaf/lapack/tile.h>
+
29#include <dlaf/matrix/matrix.h>
+
30#include <dlaf/matrix/tile.h>
+
31#include <dlaf/matrix/views.h>
+
32#include <dlaf/types.h>
+
33#include <dlaf/util_matrix.h>
+
34
+
35#ifdef DLAF_WITH_GPU
+
36#include <whip.hpp>
+
37
+
38#include <dlaf/blas/tile.h>
+
39#endif
+
40
+
41namespace dlaf::factorization::internal {
+
42
+
43namespace tfactor_l {
+
44template <Backend backend, Device device, class T>
+
45struct Helpers {};
+
46
+
47template <class T>
+
+
48struct Helpers<Backend::MC, Device::CPU, T> {
+
49 static auto prepareT(matrix::ReadOnlyTileSender<T, Device::CPU> taus,
+
50 matrix::ReadWriteTileSender<T, Device::CPU> tile_t) {
+
51 namespace ex = pika::execution::experimental;
+
52 namespace di = dlaf::internal;
+
53 return ex::when_all(std::move(taus), std::move(tile_t)) |
+
54 di::transform(di::Policy<Backend::MC>(pika::execution::thread_priority::high),
+
55 [](const matrix::Tile<const T, Device::CPU>& taus,
+
56 matrix::Tile<T, Device::CPU>&& tile_t) {
+
57 tile::internal::set0<T>(tile_t);
+
58
+
59 const SizeType k = tile_t.size().cols();
+
60 lapack::lacpy(blas::Uplo::General, 1, k, taus.ptr(), 1, tile_t.ptr(),
+
61 tile_t.ld() + 1);
+
62
+
63 return std::move(tile_t);
+
64 });
+
65 }
+
66
+
67 static auto stepGEMV(matrix::ReadOnlyTileSender<T, Device::CPU> tile_vi,
+
68 matrix::ReadOnlyTileSender<T, Device::CPU> taus,
+
69 matrix::ReadWriteTileSender<T, Device::CPU> tile_t) {
+
70 auto gemv_func = [](const matrix::Tile<const T, Device::CPU>& tile_v,
+
71 const matrix::Tile<const T, Device::CPU>& taus,
+
72 matrix::Tile<T, Device::CPU> tile_t) noexcept {
+
73 const SizeType k = tile_t.size().cols();
+
74
+
75 DLAF_ASSERT(tile_v.size().cols() == k, tile_v.size().cols(), k);
+
76 DLAF_ASSERT(taus.size().rows() == k, taus.size().rows(), k);
+
77
+
78 common::internal::SingleThreadedBlasScope single;
+
79 for (SizeType j = 0; j < k; ++j) {
+
80 // T(0:j, j) = -tau . V(j:, 0:j)* . V(j:, j)
+
81 // [j x 1] = [(n-j) x j]* . [(n-j) x 1]
+
82 const TileElementIndex t_start{0, j};
+
83 const TileElementIndex va_start{0, 0};
+
84 const TileElementIndex vb_start{0, j};
+
85 const TileElementSize va_size{tile_v.size().rows(), j};
+
86 const T tau = tile_t({j, j});
+
87
+
88 blas::gemv(blas::Layout::ColMajor, blas::Op::ConjTrans, va_size.rows(), va_size.cols(), -tau,
+
89 tile_v.ptr(va_start), tile_v.ld(), tile_v.ptr(vb_start), 1, 1, tile_t.ptr(t_start),
+
90 1);
+
91 }
+
92 return tile_t;
+
93 };
+
94
+
95 namespace ex = pika::execution::experimental;
+
96 namespace di = dlaf::internal;
+
97
+
98 return ex::when_all(tile_vi, std::move(taus), std::move(tile_t)) |
+
99 di::transform(di::Policy<Backend::MC>(pika::execution::thread_priority::high),
+
100 std::move(gemv_func));
+
101 }
+
102
+
103 static void trmvLoop(const matrix::Tile<T, Device::CPU>& tile_t) {
+
104 common::internal::SingleThreadedBlasScope single;
+
105
+
106 const SizeType k = tile_t.size().cols();
+
107
+
108 // Update each column (in order) t = T . t
+
109 // remember that T is upper triangular, so it is possible to use TRMV
+
110 for (SizeType j = 0; j < k; ++j) {
+
111 const TileElementIndex t_start{0, j};
+
112 const TileElementSize t_size{j, 1};
+
113
+
114 blas::trmv(blas::Layout::ColMajor, blas::Uplo::Upper, blas::Op::NoTrans, blas::Diag::NonUnit,
+
115 t_size.rows(), tile_t.ptr(), tile_t.ld(), tile_t.ptr(t_start), 1);
+
116 }
+
117 }
+
118
+
119 static auto stepTRMV(matrix::ReadWriteTileSender<T, Device::CPU> tile_t) noexcept {
+
120 namespace di = dlaf::internal;
+
121
+
122 return std::move(tile_t) |
+
123 di::transform(di::Policy<Backend::MC>(pika::execution::thread_priority::high), trmvLoop);
+
124 }
+
125
+
126 static auto stepCopyDiagAndTRMV(matrix::ReadOnlyTileSender<T, Device::CPU> taus,
+
127 matrix::ReadWriteTileSender<T, Device::CPU> tile_t) noexcept {
+
128 auto tausdiag_trmvloop = [](const matrix::Tile<const T, Device::CPU>& taus,
+
129 matrix::Tile<T, Device::CPU> tile_t) {
+
130 common::internal::SingleThreadedBlasScope single;
+
131
+
132 const SizeType k = tile_t.size().cols();
+
133 lapack::lacpy(blas::Uplo::General, 1, k, taus.ptr(), 1, tile_t.ptr(), tile_t.ld() + 1);
+
134
+
135 trmvLoop(tile_t);
+
136 };
+
137
+
138 namespace di = dlaf::internal;
+
139 namespace ex = pika::execution::experimental;
+
140
+
141 return ex::when_all(std::move(taus), std::move(tile_t)) |
+
142 di::transform(di::Policy<Backend::MC>(pika::execution::thread_priority::high),
+
143 std::move(tausdiag_trmvloop));
+
144 }
+
145};
-
145
-
146#ifdef DLAF_WITH_GPU
-
147template <class T>
-
-
148struct Helpers<Backend::GPU, Device::GPU, T> {
-
149 static auto prepareT(matrix::ReadOnlyTileSender<T, Device::CPU> taus,
-
150 matrix::ReadWriteTileSender<T, Device::GPU> tile_t) {
-
151 namespace di = dlaf::internal;
-
152 namespace ex = pika::execution::experimental;
-
153
-
154 return ex::when_all(std::move(taus), std::move(tile_t)) |
-
155 di::transform(di::Policy<Backend::GPU>(pika::execution::thread_priority::high),
-
156 [](const matrix::Tile<const T, Device::CPU>& taus,
-
157 matrix::Tile<T, Device::GPU>& tile_t, whip::stream_t stream) {
-
158 tile::internal::set0<T>(tile_t, stream);
-
159
-
160 const SizeType k = tile_t.size().cols();
-
161 whip::memcpy_2d_async(tile_t.ptr(), to_sizet(tile_t.ld() + 1) * sizeof(T),
-
162 taus.ptr(), sizeof(T), sizeof(T), to_sizet(k),
-
163 whip::memcpy_host_to_device, stream);
-
164
-
165 return std::move(tile_t);
-
166 });
-
167 }
-
168
-
169 static auto stepGEMV(matrix::ReadOnlyTileSender<T, Device::GPU> tile_vi,
-
170 matrix::ReadOnlyTileSender<T, Device::CPU> taus,
-
171 matrix::ReadWriteTileSender<T, Device::GPU> tile_t) noexcept {
-
172 auto gemv_func = [](cublasHandle_t handle, const matrix::Tile<const T, Device::GPU>& tile_v,
-
173 const matrix::Tile<const T, Device::CPU>& taus,
-
174 matrix::Tile<T, Device::GPU>& tile_t) noexcept {
-
175 const SizeType k = tile_t.size().cols();
-
176 DLAF_ASSERT(tile_v.size().cols() == k, tile_v.size().cols(), k);
-
177 DLAF_ASSERT(taus.size().rows() == k, taus.size().rows(), k);
-
178 DLAF_ASSERT(taus.size().cols() == 1, taus.size().cols());
-
179
-
180 for (SizeType j = 0; j < k; ++j) {
-
181 // T(0:j, j) = -tau . V(j:, 0:j)* . V(j:, j)
-
182 // [j x 1] = [(n-j) x j]* . [(n-j) x 1]
-
183 const TileElementIndex va_start{0, 0};
-
184 const TileElementIndex vb_start{0, j};
-
185 const TileElementSize va_size{tile_v.size().rows(), j};
-
186 const TileElementIndex t_start{0, j};
-
187 const auto neg_tau = util::blasToCublasCast(-taus({j, 0}));
-
188 const auto one = util::blasToCublasCast(T{1});
-
189
-
190 gpublas::internal::Gemv<T>::call(handle, CUBLAS_OP_C, to_int(va_size.rows()),
-
191 to_int(va_size.cols()), &neg_tau,
-
192 util::blasToCublasCast(tile_v.ptr(va_start)),
-
193 to_int(tile_v.ld()),
-
194 util::blasToCublasCast(tile_v.ptr(vb_start)), 1, &one,
-
195 util::blasToCublasCast(tile_t.ptr(t_start)), 1);
-
196 }
-
197 return std::move(tile_t);
-
198 };
-
199
-
200 namespace ex = pika::execution::experimental;
-
201 namespace di = dlaf::internal;
-
202
-
203 return ex::when_all(std::move(tile_vi), std::move(taus), std::move(tile_t)) |
-
204 di::transform<di::TransformDispatchType::Blas>(
-
205 di::Policy<Backend::GPU>(pika::execution::thread_priority::high), std::move(gemv_func));
-
206 }
-
207
-
208 static void trmvLoop(cublasHandle_t handle, const matrix::Tile<T, Device::GPU>& tile_t) {
-
209 const SizeType k = tile_t.size().cols();
-
210
-
211 // Update each column (in order) t = T . t
-
212 // remember that T is upper triangular, so it is possible to use TRMV
-
213 for (SizeType j = 0; j < k; ++j) {
-
214 const TileElementIndex t_start{0, j};
-
215 const TileElementSize t_size{j, 1};
-
216
-
217 gpublas::internal::Trmv<T>::call(handle, CUBLAS_FILL_MODE_UPPER, CUBLAS_OP_N, CUBLAS_DIAG_NON_UNIT,
-
218 to_int(t_size.rows()), util::blasToCublasCast(tile_t.ptr()),
-
219 to_int(tile_t.ld()), util::blasToCublasCast(tile_t.ptr(t_start)),
-
220 1);
-
221 }
-
222 }
-
223
-
224 static auto stepTRMV(matrix::ReadWriteTileSender<T, Device::GPU> tile_t) noexcept {
-
225 namespace di = dlaf::internal;
-
226
-
227 return std::move(tile_t) |
-
228 di::transform<di::TransformDispatchType::Blas>(
-
229 di::Policy<Backend::GPU>(pika::execution::thread_priority::high), trmvLoop);
-
230 }
-
231
-
232 static auto stepCopyDiagAndTRMV(matrix::ReadOnlyTileSender<T, Device::CPU> taus,
-
233 matrix::ReadWriteTileSender<T, Device::GPU> tile_t) noexcept {
-
234 // Update each column (in order) t = T . t
-
235 // remember that T is upper triangular, so it is possible to use TRMV
-
236 auto trmv_func = [](cublasHandle_t handle, const matrix::Tile<const T, Device::CPU>& taus,
-
237 matrix::Tile<T, Device::GPU>& tile_t) {
-
238 whip::stream_t stream;
-
239 DLAF_GPUBLAS_CHECK_ERROR(cublasGetStream(handle, &stream));
-
240
-
241 const SizeType k = tile_t.size().cols();
-
242 whip::memcpy_2d_async(tile_t.ptr(), to_sizet(tile_t.ld() + 1) * sizeof(T), taus.ptr(), sizeof(T),
-
243 sizeof(T), to_sizet(k), whip::memcpy_host_to_device, stream);
-
244
-
245 trmvLoop(handle, tile_t);
-
246 };
-
247
-
248 namespace ex = pika::execution::experimental;
-
249 namespace di = dlaf::internal;
-
250
-
251 return ex::when_all(std::move(taus), std::move(tile_t)) |
-
252 di::transform<di::TransformDispatchType::Blas>(
-
253 di::Policy<Backend::GPU>(pika::execution::thread_priority::high), std::move(trmv_func));
-
254 }
-
255};
+
146
+
147#ifdef DLAF_WITH_GPU
+
148template <class T>
+
+
149struct Helpers<Backend::GPU, Device::GPU, T> {
+
150 static auto prepareT(matrix::ReadOnlyTileSender<T, Device::CPU> taus,
+
151 matrix::ReadWriteTileSender<T, Device::GPU> tile_t) {
+
152 namespace di = dlaf::internal;
+
153 namespace ex = pika::execution::experimental;
+
154
+
155 return ex::when_all(std::move(taus), std::move(tile_t)) |
+
156 di::transform(di::Policy<Backend::GPU>(pika::execution::thread_priority::high),
+
157 [](const matrix::Tile<const T, Device::CPU>& taus,
+
158 matrix::Tile<T, Device::GPU>& tile_t, whip::stream_t stream) {
+
159 tile::internal::set0<T>(tile_t, stream);
+
160
+
161 const SizeType k = tile_t.size().cols();
+
162 whip::memcpy_2d_async(tile_t.ptr(), to_sizet(tile_t.ld() + 1) * sizeof(T),
+
163 taus.ptr(), sizeof(T), sizeof(T), to_sizet(k),
+
164 whip::memcpy_host_to_device, stream);
+
165
+
166 return std::move(tile_t);
+
167 });
+
168 }
+
169
+
170 static auto stepGEMV(matrix::ReadOnlyTileSender<T, Device::GPU> tile_vi,
+
171 matrix::ReadOnlyTileSender<T, Device::CPU> taus,
+
172 matrix::ReadWriteTileSender<T, Device::GPU> tile_t) noexcept {
+
173 auto gemv_func = [](cublasHandle_t handle, const matrix::Tile<const T, Device::GPU>& tile_v,
+
174 const matrix::Tile<const T, Device::CPU>& taus,
+
175 matrix::Tile<T, Device::GPU>& tile_t) noexcept {
+
176 const SizeType m = tile_v.size().rows();
+
177 const SizeType k = tile_t.size().cols();
+
178 DLAF_ASSERT(tile_v.size().cols() == k, tile_v.size().cols(), k);
+
179 DLAF_ASSERT(taus.size().rows() == k, taus.size().rows(), k);
+
180 DLAF_ASSERT(taus.size().cols() == 1, taus.size().cols());
+
181
+
182 gpulapack::larft_gemv0(handle, m, k, tile_v.ptr(), tile_v.ld(), taus.ptr(), tile_t.ptr(),
+
183 tile_t.ld());
+
184
+
185 return std::move(tile_t);
+
186 };
+
187
+
188 namespace ex = pika::execution::experimental;
+
189 namespace di = dlaf::internal;
+
190
+
191 return ex::when_all(std::move(tile_vi), std::move(taus), std::move(tile_t)) |
+
192 di::transform<di::TransformDispatchType::Blas>(
+
193 di::Policy<Backend::GPU>(pika::execution::thread_priority::high), std::move(gemv_func));
+
194 }
+
195
+
196 static void trmvLoop(cublasHandle_t handle, const matrix::Tile<T, Device::GPU>& tile_t) {
+
197 const SizeType k = tile_t.size().cols();
+
198
+
199 // Update each column (in order) t = T . t
+
200 // remember that T is upper triangular, so it is possible to use TRMV
+
201 for (SizeType j = 0; j < k; ++j) {
+
202 const TileElementIndex t_start{0, j};
+
203 const TileElementSize t_size{j, 1};
+
204
+
205 gpublas::internal::Trmv<T>::call(handle, CUBLAS_FILL_MODE_UPPER, CUBLAS_OP_N, CUBLAS_DIAG_NON_UNIT,
+
206 to_int(t_size.rows()), util::blasToCublasCast(tile_t.ptr()),
+
207 to_int(tile_t.ld()), util::blasToCublasCast(tile_t.ptr(t_start)),
+
208 1);
+
209 }
+
210 }
+
211
+
212 static auto stepTRMV(matrix::ReadWriteTileSender<T, Device::GPU> tile_t) noexcept {
+
213 namespace di = dlaf::internal;
+
214
+
215 return std::move(tile_t) |
+
216 di::transform<di::TransformDispatchType::Blas>(
+
217 di::Policy<Backend::GPU>(pika::execution::thread_priority::high), trmvLoop);
+
218 }
+
219
+
220 static auto stepCopyDiagAndTRMV(matrix::ReadOnlyTileSender<T, Device::CPU> taus,
+
221 matrix::ReadWriteTileSender<T, Device::GPU> tile_t) noexcept {
+
222 // Update each column (in order) t = T . t
+
223 // remember that T is upper triangular, so it is possible to use TRMV
+
224 auto trmv_func = [](cublasHandle_t handle, const matrix::Tile<const T, Device::CPU>& taus,
+
225 matrix::Tile<T, Device::GPU>& tile_t) {
+
226 whip::stream_t stream;
+
227 DLAF_GPUBLAS_CHECK_ERROR(cublasGetStream(handle, &stream));
+
228
+
229 const SizeType k = tile_t.size().cols();
+
230 whip::memcpy_2d_async(tile_t.ptr(), to_sizet(tile_t.ld() + 1) * sizeof(T), taus.ptr(), sizeof(T),
+
231 sizeof(T), to_sizet(k), whip::memcpy_host_to_device, stream);
+
232
+
233 trmvLoop(handle, tile_t);
+
234 };
+
235
+
236 namespace ex = pika::execution::experimental;
+
237 namespace di = dlaf::internal;
+
238
+
239 return ex::when_all(std::move(taus), std::move(tile_t)) |
+
240 di::transform<di::TransformDispatchType::Blas>(
+
241 di::Policy<Backend::GPU>(pika::execution::thread_priority::high), std::move(trmv_func));
+
242 }
+
243};
-
256#endif
-
257}
+
244#endif
+
245}
+
246
+
247template <Backend backend, Device device, class T>
+
+
248void QR_Tfactor<backend, device, T>::call(matrix::Panel<Coord::Col, T, device>& hh_panel,
+
249 matrix::ReadOnlyTileSender<T, Device::CPU> taus,
+
250 matrix::ReadWriteTileSender<T, device> tile_t) {
+
251 namespace ex = pika::execution::experimental;
+
252
+
253 using Helpers = tfactor_l::Helpers<backend, device, T>;
+
254
+
255 // Fast return in case of no reflectors
+
256 if (hh_panel.getWidth() == 0)
+
257 return;
258
-
259template <Backend backend, Device device, class T>
-
-
260void QR_Tfactor<backend, device, T>::call(matrix::Panel<Coord::Col, T, device>& hh_panel,
-
261 matrix::ReadOnlyTileSender<T, Device::CPU> taus,
-
262 matrix::ReadWriteTileSender<T, device> tile_t) {
-
263 namespace ex = pika::execution::experimental;
-
264
-
265 using Helpers = tfactor_l::Helpers<backend, device, T>;
-
266
-
267 // Fast return in case of no reflectors
-
268 if (hh_panel.getWidth() == 0)
-
269 return;
-
270
-
271 tile_t = Helpers::prepareT(taus, std::move(tile_t));
-
272
-
273 // Note:
-
274 // T factor is an upper triangular square matrix, built column by column
-
275 // with taus values on the diagonal
-
276 //
-
277 // T(j,j) = tau(j)
-
278 //
-
279 // and in the upper triangular part the following formula applies
-
280 //
-
281 // T(0:j, j) = T(0:j, 0:j) . -tau(j) . V(j:, 0:j)* . V(j:, j)
-
282 //
-
283 //
-
284 // The result is achieved in two main steps:
-
285 // 1) t = -tau(j) . V(j:, 0:j)* . V(j:, j)
-
286 // 2) T(0:j, j) = T(0:j, 0:j) . t
-
287
-
288 // 1st step: compute the column partial result `t`
-
289 // First we compute the matrix vector multiplication for each column
-
290 // -tau(j) . V(j:, 0:j)* . V(j:, j)
-
291 for (const auto& i_lc : hh_panel.iteratorLocal()) {
-
292 // Note:
-
293 // Since we are writing always on the same t, the gemv are serialized
-
294 // A possible solution to this would be to have multiple places where to store partial
-
295 // results, and then locally reduce them just before the reduce over ranks
-
296 tile_t = Helpers::stepGEMV(hh_panel.read(i_lc), taus, std::move(tile_t));
-
297 }
-
298
-
299 // 2nd step: compute the T factor, by performing the last step on each column
-
300 // each column depends on the previous part (all reflectors that comes before)
-
301 // so it is performed sequentially
-
302 ex::start_detached(Helpers::stepTRMV(std::move(tile_t)));
-
303}
+
259 tile_t = Helpers::prepareT(taus, std::move(tile_t));
+
260
+
261 // Note:
+
262 // T factor is an upper triangular square matrix, built column by column
+
263 // with taus values on the diagonal
+
264 //
+
265 // T(j,j) = tau(j)
+
266 //
+
267 // and in the upper triangular part the following formula applies
+
268 //
+
269 // T(0:j, j) = T(0:j, 0:j) . -tau(j) . V(j:, 0:j)* . V(j:, j)
+
270 //
+
271 //
+
272 // The result is achieved in two main steps:
+
273 // 1) t = -tau(j) . V(j:, 0:j)* . V(j:, j)
+
274 // 2) T(0:j, j) = T(0:j, 0:j) . t
+
275
+
276 // 1st step: compute the column partial result `t`
+
277 // First we compute the matrix vector multiplication for each column
+
278 // -tau(j) . V(j:, 0:j)* . V(j:, j)
+
279 for (const auto& i_lc : hh_panel.iteratorLocal()) {
+
280 // Note:
+
281 // Since we are writing always on the same t, the gemv are serialized
+
282 // A possible solution to this would be to have multiple places where to store partial
+
283 // results, and then locally reduce them just before the reduce over ranks
+
284 tile_t = Helpers::stepGEMV(hh_panel.read(i_lc), taus, std::move(tile_t));
+
285 }
+
286
+
287 // 2nd step: compute the T factor, by performing the last step on each column
+
288 // each column depends on the previous part (all reflectors that comes before)
+
289 // so it is performed sequentially
+
290 ex::start_detached(Helpers::stepTRMV(std::move(tile_t)));
+
291}
-
304
-
305template <Backend backend, Device device, class T>
-
-
306void QR_Tfactor<backend, device, T>::call(
-
307 matrix::Panel<Coord::Col, T, device>& hh_panel, matrix::ReadOnlyTileSender<T, Device::CPU> taus,
-
308 matrix::ReadWriteTileSender<T, device> tile_t,
-
309 comm::CommunicatorPipeline<comm::CommunicatorType::Col>& mpi_col_task_chain) {
-
310 namespace ex = pika::execution::experimental;
-
311
-
312 using Helpers = tfactor_l::Helpers<backend, device, T>;
-
313
-
314 // Fast return in case of no reflectors
-
315 if (hh_panel.getWidth() == 0)
-
316 return;
-
317
-
318 tile_t = Helpers::prepareT(taus, std::move(tile_t));
-
319
-
320 // Note:
-
321 // T factor is an upper triangular square matrix, built column by column
-
322 // with taus values on the diagonal
-
323 //
-
324 // T(j,j) = tau(j)
-
325 //
-
326 // and in the upper triangular part the following formula applies
-
327 //
-
328 // T(0:j, j) = T(0:j, 0:j) . -tau(j) . V(j:, 0:j)* . V(j:, j)
-
329 //
-
330 //
-
331 // The result is achieved in two main steps:
-
332 // 1) t = -tau(j) . V(j:, 0:j)* . V(j:, j)
-
333 // 2) T(0:j, j) = T(0:j, 0:j) . t
-
334
-
335 // 1st step: compute the column partial result `t`
-
336 // First we compute the matrix vector multiplication for each column
-
337 // -tau(j) . V(j:, 0:j)* . V(j:, j)
-
338 for (const auto& i_lc : hh_panel.iteratorLocal()) {
-
339 // Note:
-
340 // Since we are writing always on the same t, the gemv are serialized
-
341 // A possible solution to this would be to have multiple places where to store partial
-
342 // results, and then locally reduce them just before the reduce over ranks
-
343 tile_t = Helpers::stepGEMV(hh_panel.read(i_lc), taus, std::move(tile_t));
-
344 }
-
345
-
346 // at this point each rank has its partial result for each column
-
347 // so, let's reduce the results (on all ranks, so that everyone can independently compute T factor)
-
348 if (mpi_col_task_chain.size() > 1)
-
349 tile_t = schedule_all_reduce_in_place(mpi_col_task_chain.exclusive(), MPI_SUM, std::move(tile_t));
-
350
-
351 // 2nd step: compute the T factor, by performing the last step on each column
-
352 // each column depends on the previous part (all reflectors that comes before)
-
353 // so it is performed sequentially
-
354 ex::start_detached(Helpers::stepCopyDiagAndTRMV(taus, std::move(tile_t)));
-
355}
+
292
+
293template <Backend backend, Device device, class T>
+
+
294void QR_Tfactor<backend, device, T>::call(
+
295 matrix::Panel<Coord::Col, T, device>& hh_panel, matrix::ReadOnlyTileSender<T, Device::CPU> taus,
+
296 matrix::ReadWriteTileSender<T, device> tile_t,
+
297 comm::CommunicatorPipeline<comm::CommunicatorType::Col>& mpi_col_task_chain) {
+
298 namespace ex = pika::execution::experimental;
+
299
+
300 using Helpers = tfactor_l::Helpers<backend, device, T>;
+
301
+
302 // Fast return in case of no reflectors
+
303 if (hh_panel.getWidth() == 0)
+
304 return;
+
305
+
306 tile_t = Helpers::prepareT(taus, std::move(tile_t));
+
307
+
308 // Note:
+
309 // T factor is an upper triangular square matrix, built column by column
+
310 // with taus values on the diagonal
+
311 //
+
312 // T(j,j) = tau(j)
+
313 //
+
314 // and in the upper triangular part the following formula applies
+
315 //
+
316 // T(0:j, j) = T(0:j, 0:j) . -tau(j) . V(j:, 0:j)* . V(j:, j)
+
317 //
+
318 //
+
319 // The result is achieved in two main steps:
+
320 // 1) t = -tau(j) . V(j:, 0:j)* . V(j:, j)
+
321 // 2) T(0:j, j) = T(0:j, 0:j) . t
+
322
+
323 // 1st step: compute the column partial result `t`
+
324 // First we compute the matrix vector multiplication for each column
+
325 // -tau(j) . V(j:, 0:j)* . V(j:, j)
+
326 for (const auto& i_lc : hh_panel.iteratorLocal()) {
+
327 // Note:
+
328 // Since we are writing always on the same t, the gemv are serialized
+
329 // A possible solution to this would be to have multiple places where to store partial
+
330 // results, and then locally reduce them just before the reduce over ranks
+
331 tile_t = Helpers::stepGEMV(hh_panel.read(i_lc), taus, std::move(tile_t));
+
332 }
+
333
+
334 // at this point each rank has its partial result for each column
+
335 // so, let's reduce the results (on all ranks, so that everyone can independently compute T factor)
+
336 if (mpi_col_task_chain.size() > 1)
+
337 tile_t = schedule_all_reduce_in_place(mpi_col_task_chain.exclusive(), MPI_SUM, std::move(tile_t));
+
338
+
339 // 2nd step: compute the T factor, by performing the last step on each column
+
340 // each column depends on the previous part (all reflectors that comes before)
+
341 // so it is performed sequentially
+
342 ex::start_detached(Helpers::stepCopyDiagAndTRMV(taus, std::move(tile_t)));
+
343}
-
356
-
357}
+
344
+
345}
dlaf::comm::CommunicatorPipeline
Definition communicator_pipeline.h:41
dlaf::comm::CommunicatorPipeline::size
IndexT_MPI size() const noexcept
Return the size of the grid.
Definition communicator_pipeline.h:68
@@ -463,8 +451,8 @@
single_threaded_blas.h
-
dlaf::factorization::internal::QR_Tfactor::call
static void call(matrix::Panel< Coord::Col, T, device > &panel_view, matrix::ReadOnlyTileSender< T, Device::CPU > taus, matrix::ReadWriteTileSender< T, device > t)
Definition t_factor_impl.h:260
-
dlaf::factorization::internal::tfactor_l::Helpers
Definition t_factor_impl.h:44
+
dlaf::factorization::internal::QR_Tfactor::call
static void call(matrix::Panel< Coord::Col, T, device > &panel_view, matrix::ReadOnlyTileSender< T, Device::CPU > taus, matrix::ReadWriteTileSender< T, device > t)
Definition t_factor_impl.h:248
+
dlaf::factorization::internal::tfactor_l::Helpers
Definition t_factor_impl.h:45
dlaf::to_sizet
auto to_sizet(const T signed_value)
Definition types.h:266
dlaf::to_int
auto to_int(const T unsigned_value)
Definition types.h:248