Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
gmres_device.F90
Go to the documentation of this file.
1! Copyright (c) 2022-2024, The Neko Authors
2! All rights reserved.
3!
4! Redistribution and use in source and binary forms, with or without
5! modification, are permitted provided that the following conditions
6! are met:
7!
8! * Redistributions of source code must retain the above copyright
9! notice, this list of conditions and the following disclaimer.
10!
11! * Redistributions in binary form must reproduce the above
12! copyright notice, this list of conditions and the following
13! disclaimer in the documentation and/or other materials provided
14! with the distribution.
15!
16! * Neither the name of the authors nor the names of its
17! contributors may be used to endorse or promote products derived
18! from this software without specific prior written permission.
19!
20! THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
21! "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
22! LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
23! FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
24! COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
25! INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
26! BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
27! LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
28! CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
29! LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
30! ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
31! POSSIBILITY OF SUCH DAMAGE.
32!
36 use krylov, only : ksp_t, ksp_monitor_t
37 use gmres, only : gmres_space_size
38 use precon, only : pc_t
39 use ax_product, only : ax_t
40 use num_types, only : rp, c_rp
41 use field, only : field_t
42 use coefs, only : coef_t
43 use gather_scatter, only : gs_t, gs_op_add
48 use math, only : rone, rzero, abscmp
56 use utils, only : neko_error
57 use json_module, only : json_file
59 use logger, only : neko_log, log_size
61 use mpi_f08, only : mpi_in_place, mpi_sum, mpi_allreduce
62 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_loc, &
63 c_associated, c_int, c_size_t, c_sizeof
64 implicit none
65 private
66
68 type, public, extends(ksp_t) :: gmres_device_t
69 integer :: m_restart = gmres_space_size
70 real(kind=rp), allocatable :: w(:)
71 real(kind=rp), allocatable :: c(:)
72 real(kind=rp), allocatable :: r(:)
73 real(kind=rp), allocatable :: z(:,:)
74 real(kind=rp), allocatable :: h(:,:)
75 real(kind=rp), allocatable :: v(:,:)
76 real(kind=rp), allocatable :: s(:)
77 real(kind=rp), allocatable :: gam(:)
78 type(c_ptr) :: w_d = c_null_ptr
79 type(c_ptr) :: c_d = c_null_ptr
80 type(c_ptr) :: r_d = c_null_ptr
81 type(c_ptr) :: s_d = c_null_ptr
82 type(c_ptr) :: gam_d = c_null_ptr
83 type(c_ptr), allocatable :: z_d(:), h_d(:), v_d(:)
84 type(c_ptr) :: z_d_d = c_null_ptr
85 type(c_ptr) :: h_d_d = c_null_ptr
86 type(c_ptr) :: v_d_d = c_null_ptr
87 type(c_ptr) :: gs_event = c_null_ptr
88 contains
89 procedure, pass(this) :: init => gmres_device_init
90 procedure, pass(this) :: init_from_json => gmres_device_init_from_json
91 procedure, pass(this) :: free => gmres_device_free
92 procedure, pass(this) :: solve => gmres_device_solve
93 procedure, pass(this) :: solve_coupled => gmres_device_solve_coupled
94 end type gmres_device_t
95
96#ifdef HAVE_HIP
97 interface
98 real(c_rp) function hip_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n) &
99 bind(c, name = 'hip_gmres_part2')
100 use, intrinsic :: iso_c_binding
101 import c_rp
102 implicit none
103 type(c_ptr), value :: h_d, w_d, v_d_d, mult_d
104 integer(c_int) :: j, n
105 end function hip_gmres_part2
106 end interface
107#elif HAVE_CUDA
108
109 interface
110 real(c_rp) function cuda_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n) &
111 bind(c, name = 'cuda_gmres_part2')
112 use, intrinsic :: iso_c_binding
113 import c_rp
114 implicit none
115 type(c_ptr), value :: h_d, w_d, v_d_d, mult_d
116 integer(c_int) :: j, n
117 end function cuda_gmres_part2
118 end interface
119#endif
120
121contains
122
123 function device_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n) result(alpha)
124 type(c_ptr), value :: h_d, w_d, v_d_d, mult_d
125 integer(c_int) :: j, n
126 real(c_rp) :: alpha
127 integer :: ierr
128#ifdef HAVE_HIP
129 alpha = hip_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n)
130#elif HAVE_CUDA
131 alpha = cuda_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n)
132#else
133 call neko_error('No device backend configured')
134#endif
135
136#ifndef HAVE_DEVICE_MPI
137 if (pe_size .gt. 1) then
138 call mpi_allreduce(mpi_in_place, alpha, 1, &
139 mpi_real_precision, mpi_sum, neko_comm, ierr)
140 end if
141#endif
142
143 end function device_gmres_part2
144
146 subroutine gmres_device_init(this, n, max_iter, M, rel_tol, abs_tol, monitor)
147 class(gmres_device_t), target, intent(inout) :: this
148 integer, intent(in) :: n
149 integer, intent(in) :: max_iter
150 class(pc_t), optional, intent(in), target :: M
151 real(kind=rp), optional, intent(in) :: rel_tol
152 real(kind=rp), optional, intent(in) :: abs_tol
153 logical, optional, intent(in) :: monitor
154 type(device_ident_t), target :: M_ident
155 type(c_ptr) :: ptr
156 integer(c_size_t) :: z_size
157 integer :: i
158
159 call this%free()
160
161 if (present(m)) then
162 this%M => m
163 else
164 this%M => m_ident
165 end if
166
167 allocate(this%w(n))
168 allocate(this%r(n))
169 call device_map(this%w, this%w_d, n)
170 call device_map(this%r, this%r_d, n)
171
172 allocate(this%c(this%m_restart))
173 allocate(this%s(this%m_restart))
174 allocate(this%gam(this%m_restart + 1))
175 call device_map(this%c, this%c_d, this%m_restart)
176 call device_map(this%s, this%s_d, this%m_restart)
177 call device_map(this%gam, this%gam_d, this%m_restart+1)
178
179 allocate(this%z(n, this%m_restart))
180 allocate(this%v(n, this%m_restart))
181 allocate(this%h(this%m_restart, this%m_restart))
182 allocate(this%z_d(this%m_restart))
183 allocate(this%v_d(this%m_restart))
184 allocate(this%h_d(this%m_restart))
185 do i = 1, this%m_restart
186 this%z_d(i) = c_null_ptr
187 call device_map(this%z(:,i), this%z_d(i), n)
188
189 this%v_d(i) = c_null_ptr
190 call device_map(this%v(:,i), this%v_d(i), n)
191
192 this%h_d(i) = c_null_ptr
193 call device_map(this%h(:,i), this%h_d(i), this%m_restart)
194 end do
195
196 z_size = c_sizeof(c_null_ptr) * (this%m_restart)
197 call device_alloc(this%z_d_d, z_size)
198 call device_alloc(this%v_d_d, z_size)
199 call device_alloc(this%h_d_d, z_size)
200 ptr = c_loc(this%z_d)
201 call device_memcpy(ptr, this%z_d_d, z_size, &
202 host_to_device, sync = .false.)
203 ptr = c_loc(this%v_d)
204 call device_memcpy(ptr, this%v_d_d, z_size, &
205 host_to_device, sync = .false.)
206 ptr = c_loc(this%h_d)
207 call device_memcpy(ptr, this%h_d_d, z_size, &
208 host_to_device, sync = .false.)
209
210
211 if (present(rel_tol) .and. present(abs_tol) .and. present(monitor)) then
212 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
213 else if (present(rel_tol) .and. present(abs_tol)) then
214 call this%ksp_init(max_iter, rel_tol, abs_tol)
215 else if (present(monitor) .and. present(abs_tol)) then
216 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
217 else if (present(rel_tol) .and. present(monitor)) then
218 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
219 else if (present(rel_tol)) then
220 call this%ksp_init(max_iter, rel_tol = rel_tol)
221 else if (present(abs_tol)) then
222 call this%ksp_init(max_iter, abs_tol = abs_tol)
223 else if (present(monitor)) then
224 call this%ksp_init(max_iter, monitor = monitor)
225 else
226 call this%ksp_init(max_iter)
227 end if
228
229 call device_event_create(this%gs_event, 2)
230
231 end subroutine gmres_device_init
232
236 subroutine gmres_device_init_from_json(this, n, json, M)
237 class(gmres_device_t), target, intent(inout) :: this
238 integer, intent(in) :: n
239 type(json_file), intent(inout) :: json
240 class(pc_t), optional, intent(in), target :: M
241 character(len=LOG_SIZE) :: log_buf
242
243 call json_get_or_lookup_or_default(json, 'space_size', &
244 this%m_restart, gmres_space_size)
245 if (this%m_restart .lt. 1) then
246 call neko_error('GMRES space_size must be at least 1')
247 end if
248 write(log_buf, '(A,I0)') 'GMRES space: ', this%m_restart
249 call neko_log%message(log_buf)
250
251 call this%ksp_init_from_json(n, json, m = m)
252
253 end subroutine gmres_device_init_from_json
254
256 subroutine gmres_device_free(this)
257 class(gmres_device_t), intent(inout) :: this
258 integer :: i
259
260 call this%ksp_free()
261
262 if (allocated(this%w)) then
263 if (c_associated(this%w_d)) then
264 call device_unmap(this%w, this%w_d)
265 end if
266 deallocate(this%w)
267 end if
268
269 if (allocated(this%c)) then
270 if (c_associated(this%c_d)) then
271 call device_unmap(this%c, this%c_d)
272 end if
273 deallocate(this%c)
274 end if
275
276 if (allocated(this%r)) then
277 if (c_associated(this%r_d)) then
278 call device_unmap(this%r, this%r_d)
279 end if
280 deallocate(this%r)
281 end if
282
283 if (allocated(this%z)) then
284 if (allocated(this%z_d)) then
285 do i = 1, this%m_restart
286 if (c_associated(this%z_d(i))) then
287 call device_unmap(this%z(:,i), this%z_d(i))
288 end if
289 end do
290 end if
291 deallocate(this%z)
292 end if
293
294 if (allocated(this%h)) then
295 if (allocated(this%h_d)) then
296 do i = 1, this%m_restart
297 if (c_associated(this%h_d(i))) then
298 call device_unmap(this%h(:,i), this%h_d(i))
299 end if
300 end do
301 end if
302 deallocate(this%h)
303 end if
304
305 if (allocated(this%v)) then
306 if (allocated(this%v_d)) then
307 do i = 1, this%m_restart
308 if (c_associated(this%v_d(i))) then
309 call device_unmap(this%v(:,i), this%v_d(i))
310 end if
311 end do
312 end if
313 deallocate(this%v)
314 end if
315
316 if (allocated(this%s)) then
317 if (c_associated(this%s_d)) then
318 call device_unmap(this%s, this%s_d)
319 end if
320 deallocate(this%s)
321 end if
322 if (allocated(this%gam)) then
323 if (c_associated(this%gam_d)) then
324 call device_unmap(this%gam, this%gam_d)
325 end if
326 deallocate(this%gam)
327 end if
328
329 if (c_associated(this%z_d_d)) then
330 call device_free(this%z_d_d)
331 end if
332 if (c_associated(this%v_d_d)) then
333 call device_free(this%v_d_d)
334 end if
335 if (c_associated(this%h_d_d)) then
336 call device_free(this%h_d_d)
337 end if
338
339 nullify(this%M)
340
341 if (c_associated(this%gs_event)) then
342 call device_event_destroy(this%gs_event)
343 end if
344
345 end subroutine gmres_device_free
346
348 function gmres_device_solve(this, Ax, x, f, n, coef, bc_projector, gs_h, &
349 niter) result(ksp_results)
350 class(gmres_device_t), intent(inout) :: this
351 class(ax_t), intent(in) :: ax
352 type(field_t), intent(inout) :: x
353 integer, intent(in) :: n
354 real(kind=rp), dimension(n), intent(in) :: f
355 type(coef_t), intent(inout) :: coef
356 class(scalar_bc_projector_t), intent(inout) :: bc_projector
357 type(gs_t), intent(inout) :: gs_h
358 type(ksp_monitor_t) :: ksp_results
359 integer, optional, intent(in) :: niter
360 integer :: iter, max_iter
361 integer :: i, j, k
362 real(kind=rp) :: rnorm, alpha, temp, lr, alpha2, norm_fac
363 logical :: conv
364 type(c_ptr) :: f_d
365
366 f_d = device_get_ptr(f)
367
368 conv = .false.
369 iter = 0
370 rnorm = 0.0_rp
371
372 if (present(niter)) then
373 max_iter = niter
374 else
375 max_iter = this%max_iter
376 end if
377
378 associate(w => this%w, c => this%c, r => this%r, z => this%z, h => this%h, &
379 v => this%v, s => this%s, gam => this%gam, v_d => this%v_d, &
380 w_d => this%w_d, r_d => this%r_d, h_d => this%h_d, &
381 v_d_d => this%v_d_d, x_d => x%x_d, z_d_d => this%z_d_d, &
382 c_d => this%c_d)
383
384 norm_fac = 1.0_rp / sqrt(coef%volume)
385 call rzero(gam, this%m_restart + 1)
386 call rone(s, this%m_restart)
387 call rone(c, this%m_restart)
388 call rzero(h, this%m_restart * this%m_restart)
389 call device_rzero(x%x_d, n)
390 call device_rzero(this%gam_d, this%m_restart + 1)
391 call device_rone(this%s_d, this%m_restart)
392 call device_rone(this%c_d, this%m_restart)
393
394 call rzero(this%h, this%m_restart**2)
395 ! do j = 1, this%m_restart
396 ! call device_rzero(h_d(j), this%m_restart)
397 ! end do
398
399 call this%monitor_start('GMRES')
400 do while (.not. conv .and. iter .lt. max_iter)
401
402 if (iter .eq. 0) then
403 call device_copy(r_d, f_d, n)
404 else
405 call device_copy(r_d, f_d, n)
406 call ax%compute(w, x%x, coef, x%msh, x%Xh)
407 call gs_h%op(w, n, gs_op_add, this%gs_event)
408 call device_event_sync(this%gs_event)
409 call bc_projector%apply(w, n)
410 call device_sub2(r_d, w_d, n)
411 end if
412
413 gam(1) = sqrt(device_glsc3(r_d, r_d, coef%mult_d, n))
414 if (iter .eq. 0) then
415 ksp_results%res_start = gam(1) * norm_fac
416 end if
417
418 if (abscmp(gam(1), 0.0_rp)) exit
419
420 rnorm = 0.0_rp
421 temp = 1.0_rp / gam(1)
422 call device_cmult2(v_d(1), r_d, temp, n)
423 do j = 1, this%m_restart
424 iter = iter+1
425
426 call this%M%solve(z(1,j), v(1,j), n)
427
428 call ax%compute(w, z(1,j), coef, x%msh, x%Xh)
429 call gs_h%op(w, n, gs_op_add, this%gs_event)
430 call device_event_sync(this%gs_event)
431 call bc_projector%apply(w, n)
432
433 if (neko_bcknd_opencl .eq. 1 .or. neko_bcknd_metal .eq. 1) then
434 do i = 1, j
435 h(i,j) = device_glsc3(w_d, v_d(i), coef%mult_d, n)
436
437 call device_add2s2(w_d, v_d(i), -h(i,j), n)
438
439 alpha2 = device_glsc3(w_d, w_d, coef%mult_d, n)
440 end do
441 else
442 call device_glsc3_many(h(1,j), w_d, v_d_d, coef%mult_d, j, n)
443
444 call device_memcpy(h(:,j), h_d(j), j, &
445 host_to_device, sync = .false.)
446
447 alpha2 = device_gmres_part2(w_d, v_d_d, h_d(j), &
448 coef%mult_d, j, n)
449
450 end if
451
452 alpha = sqrt(alpha2)
453 do i = 1, j-1
454 temp = h(i,j)
455 h(i,j) = c(i)*temp + s(i) * h(i+1,j)
456 h(i+1,j) = -s(i)*temp + c(i) * h(i+1,j)
457 end do
458
459 rnorm = 0.0_rp
460 if (abscmp(alpha, 0.0_rp)) then
461 conv = .true.
462 exit
463 end if
464
465 lr = sqrt(h(j,j) * h(j,j) + alpha2)
466 temp = 1.0_rp / lr
467 c(j) = h(j,j) * temp
468 s(j) = alpha * temp
469 h(j,j) = lr
470 call device_memcpy(h(:,j), h_d(j), j, &
471 host_to_device, sync = .false.)
472 gam(j+1) = -s(j) * gam(j)
473 gam(j) = c(j) * gam(j)
474
475 rnorm = abs(gam(j+1)) * norm_fac
476 call this%monitor_iter(iter, rnorm)
477 if (rnorm .lt. this%abs_tol) then
478 conv = .true.
479 exit
480 end if
481
482 if (iter + 1 .gt. max_iter) exit
483
484 if (j .lt. this%m_restart) then
485 temp = 1.0_rp / alpha
486 call device_cmult2(v_d(j+1), w_d, temp, n)
487 end if
488
489 end do
490
491 j = min(j, this%m_restart)
492 do k = j, 1, -1
493 temp = gam(k)
494 do i = j, k+1, -1
495 temp = temp - h(k,i) * c(i)
496 end do
497 c(k) = temp / h(k,k)
498 end do
499
500 if (neko_bcknd_opencl .eq. 1 .or. neko_bcknd_metal .eq. 1) then
501 do i = 1, j
502 call device_add2s2(x_d, this%z_d(i), c(i), n)
503 end do
504 else
505 call device_memcpy(c, c_d, j, host_to_device, sync = .false.)
506 call device_add2s2_many(x_d, z_d_d, c_d, j, n)
507 end if
508 end do
509
510 end associate
511 call this%monitor_stop()
512 ksp_results%res_final = rnorm
513 ksp_results%iter = iter
514 ksp_results%converged = this%is_converged(iter, rnorm)
515
516 end function gmres_device_solve
517
519 function gmres_device_solve_coupled(this, Ax, x, y, z, fx, fy, fz, &
520 n, coef, bc_projector, gs_h, niter) result(ksp_results)
521 class(gmres_device_t), intent(inout) :: this
522 class(ax_t), intent(in) :: ax
523 type(field_t), intent(inout) :: x
524 type(field_t), intent(inout) :: y
525 type(field_t), intent(inout) :: z
526 integer, intent(in) :: n
527 real(kind=rp), dimension(n), intent(in) :: fx
528 real(kind=rp), dimension(n), intent(in) :: fy
529 real(kind=rp), dimension(n), intent(in) :: fz
530 type(coef_t), intent(inout) :: coef
531 class(vector_bc_projector_t), intent(inout) :: bc_projector
532 type(gs_t), intent(inout) :: gs_h
533 type(ksp_monitor_t), dimension(3) :: ksp_results
534 integer, optional, intent(in) :: niter
535 type(scalar_bc_projector_t), pointer :: bc_x, bc_y, bc_z
536
537 call vector_bc_projector_components(bc_projector, bc_x, bc_y, bc_z)
538 ksp_results(1) = this%solve(ax, x, fx, n, coef, bc_x, gs_h, niter)
539 ksp_results(2) = this%solve(ax, y, fy, n, coef, bc_y, gs_h, niter)
540 ksp_results(3) = this%solve(ax, z, fz, n, coef, bc_z, gs_h, niter)
541
542 end function gmres_device_solve_coupled
543
544end module gmres_device
__device__ T solve(const T u, const T y, const T guess, const T nu, const T kappa, const T B)
real cuda_gmres_part2(void *w, void *v, void *h, void *mult, int *j, int *n)
Definition gmres_aux.cu:62
Return the device pointer for an associated Fortran array.
Definition device.F90:113
Map a Fortran array to a device (allocate and associate)
Definition device.F90:83
Copy data between host and device (or device and device)
Definition device.F90:72
Unmap a Fortran array from a device (deassociate and free)
Definition device.F90:89
Defines a Matrix-vector product.
Definition ax.f90:34
Coefficients.
Definition coef.f90:34
Definition comm.F90:1
type(mpi_datatype), public mpi_real_precision
MPI type for working precision of REAL types.
Definition comm.F90:54
integer, public pe_size
MPI size of communicator.
Definition comm.F90:62
type(mpi_comm), public neko_comm
MPI communicator.
Definition comm.F90:46
Identity Krylov preconditioner for accelerators.
subroutine, public device_add2s1(a_d, b_d, c1, n, strm)
subroutine, public device_add2s2_many(y_d, x_d_d, a_d, j, n, strm)
subroutine, public device_add2s2(a_d, b_d, c1, n, strm)
Vector addition with scalar multiplication (multiplication on first argument)
subroutine, public device_rzero(a_d, n, strm)
Zero a real vector.
subroutine, public device_rone(a_d, n, strm)
Set all elements to one.
subroutine, public device_glsc3_many(h, w_d, v_d_d, mult_d, j, n, strm)
subroutine, public device_sub2(a_d, b_d, n, strm)
Vector substraction .
subroutine, public device_copy(a_d, b_d, n, strm)
Copy a vector .
real(kind=rp) function, public device_glsc3(a_d, b_d, c_d, n, strm)
Weighted inner product .
subroutine, public device_cmult2(a_d, b_d, c, n, strm)
Multiplication by constant c .
Device abstraction, common interface for various accelerators.
Definition device.F90:34
subroutine, public device_event_sync(event)
Synchronize an event.
Definition device.F90:1705
integer, parameter, public host_to_device
Definition device.F90:48
subroutine, public device_free(x_d)
Deallocate memory on the device.
Definition device.F90:281
subroutine, public device_event_destroy(event)
Destroy a device event.
Definition device.F90:1661
subroutine, public device_alloc(x_d, s)
Allocate memory on the device.
Definition device.F90:250
subroutine, public device_event_create(event, flags)
Create a device event queue.
Definition device.F90:1627
Defines a field.
Definition field.f90:34
Gather-scatter.
Defines various GMRES methods.
real(c_rp) function device_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n)
subroutine gmres_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a standard GMRES solver.
subroutine gmres_device_init_from_json(this, n, json, m)
Initialise a standard GMRES solver from a JSON solver block. Reads space_size (default GMRES_SPACE_SI...
type(ksp_monitor_t) function, dimension(3) gmres_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Standard GMRES coupled solve.
type(ksp_monitor_t) function gmres_device_solve(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
Standard GMRES solve.
subroutine gmres_device_free(this)
Deallocate a standard GMRES solver.
Implements gmres_t.
Definition gmres.f90:34
integer, parameter, public gmres_space_size
Default Krylov space size (restart length) of the GMRES solvers.
Definition gmres.f90:60
Utilities for retrieving parameters from the case files.
Implements the base abstract type for Krylov solvers plus helper types.
Definition krylov.f90:34
Logging routines.
Definition log.f90:34
type(log_t), public neko_log
Global log stream.
Definition log.f90:92
integer, parameter, public log_size
Definition log.f90:46
Definition math.F90:60
subroutine, public rone(a, n)
Set all elements to one.
Definition math.F90:291
subroutine, public rzero(a, n)
Zero a real vector.
Definition math.F90:249
Build configurations.
integer, parameter neko_bcknd_opencl
integer, parameter neko_bcknd_metal
integer, parameter, public c_rp
Definition num_types.f90:15
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:14
Krylov preconditioner.
Definition precon.f90:34
Implements scalar_projector_t.
Utilities.
Definition utils.f90:35
Implements boundary condition projectors for vector fields. Two types concrete types are provided: se...
subroutine, public vector_bc_projector_components(this, x, y, z)
Access the component scalar projectors from a segregated vector projector.
Base type for a matrix-vector product providing .
Definition ax.f90:44
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:136
Defines a canonical Krylov preconditioner for accelerators.
Gather-scatter kernel.
Standard preconditioned generalized minimal residual method.
Type for storing initial and final residuals in a Krylov solver.
Definition krylov.f90:60
Base abstract type for a canonical Krylov method, solving .
Definition krylov.f90:77
Defines a canonical Krylov preconditioner.
Definition precon.f90:44
Projector for scalar boundary conditions.
Abstract type for resolving vector boundary conditions.