Neko 1.99.6
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
fusedcg_cpld_device.F90
Go to the documentation of this file.
1! Copyright (c) 2021-2026, 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 precon, only : pc_t
37 use ax_product, only : ax_t
38 use num_types, only : rp, c_rp
39 use field, only : field_t
40 use coefs, only : coef_t
41 use gather_scatter, only : gs_t, gs_op_add
42 use bc_list, only : bc_list_t
43 use math, only : glsc3, rzero, copy, abscmp
48 use utils, only : neko_error
50 use mpi_f08, only : mpi_in_place, mpi_allreduce, &
51 mpi_sum
52 use operators, only : rotate_cyc
53 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, &
54 c_associated, c_size_t, c_sizeof, c_int, c_loc
55 implicit none
56 private
57
58 integer, parameter :: device_fusedcg_cpld_p_space = 10
59
61 type, public, extends(ksp_t) :: fusedcg_cpld_device_t
62 real(kind=rp), allocatable :: w1(:)
63 real(kind=rp), allocatable :: w2(:)
64 real(kind=rp), allocatable :: w3(:)
65 real(kind=rp), allocatable :: r1(:)
66 real(kind=rp), allocatable :: r2(:)
67 real(kind=rp), allocatable :: r3(:)
68 real(kind=rp), allocatable :: z1(:)
69 real(kind=rp), allocatable :: z2(:)
70 real(kind=rp), allocatable :: z3(:)
71 real(kind=rp), allocatable :: tmp(:)
72 real(kind=rp), allocatable :: p1(:,:)
73 real(kind=rp), allocatable :: p2(:,:)
74 real(kind=rp), allocatable :: p3(:,:)
75 real(kind=rp), allocatable :: alpha(:)
76 type(c_ptr) :: w1_d = c_null_ptr
77 type(c_ptr) :: w2_d = c_null_ptr
78 type(c_ptr) :: w3_d = c_null_ptr
79 type(c_ptr) :: r1_d = c_null_ptr
80 type(c_ptr) :: r2_d = c_null_ptr
81 type(c_ptr) :: r3_d = c_null_ptr
82 type(c_ptr) :: z1_d = c_null_ptr
83 type(c_ptr) :: z2_d = c_null_ptr
84 type(c_ptr) :: z3_d = c_null_ptr
85 type(c_ptr) :: alpha_d = c_null_ptr
86 type(c_ptr) :: p1_d_d = c_null_ptr
87 type(c_ptr) :: p2_d_d = c_null_ptr
88 type(c_ptr) :: p3_d_d = c_null_ptr
89 type(c_ptr) :: tmp_d = c_null_ptr
90 type(c_ptr), allocatable :: p1_d(:)
91 type(c_ptr), allocatable :: p2_d(:)
92 type(c_ptr), allocatable :: p3_d(:)
93 type(c_ptr) :: gs_event1 = c_null_ptr
94 type(c_ptr) :: gs_event2 = c_null_ptr
95 type(c_ptr) :: gs_event3 = c_null_ptr
96 contains
97 procedure, pass(this) :: init => fusedcg_cpld_device_init
98 procedure, pass(this) :: free => fusedcg_cpld_device_free
99 procedure, pass(this) :: solve => fusedcg_cpld_device_solve
100 procedure, pass(this) :: solve_coupled => fusedcg_cpld_device_solve_coupled
101 end type fusedcg_cpld_device_t
102
103#ifdef HAVE_CUDA
104 interface
105 subroutine cuda_fusedcg_cpld_part1(a1_d, a2_d, a3_d, &
106 b1_d, b2_d, b3_d, tmp_d, n) bind(c, name = 'cuda_fusedcg_cpld_part1')
107 use, intrinsic :: iso_c_binding
108 import c_rp
109 implicit none
110 type(c_ptr), value :: a1_d, a2_d, a3_d, b1_d, b2_d, b3_d, tmp_d
111 integer(c_int) :: n
112 end subroutine cuda_fusedcg_cpld_part1
113 end interface
114
115 interface
116 subroutine cuda_fusedcg_cpld_update_p(p1_d, p2_d, p3_d, z1_d, z2_d, z3_d, &
117 po1_d, po2_d, po3_d, beta, n) &
118 bind(c, name = 'cuda_fusedcg_cpld_update_p')
119 use, intrinsic :: iso_c_binding
120 import c_rp
121 implicit none
122 type(c_ptr), value :: p1_d, p2_d, p3_d, z1_d, z2_d, z3_d
123 type(c_ptr), value :: po1_d, po2_d, po3_d
124 real(c_rp) :: beta
125 integer(c_int) :: n
126 end subroutine cuda_fusedcg_cpld_update_p
127 end interface
128
129 interface
130 subroutine cuda_fusedcg_cpld_update_x(x1_d, x2_d, x3_d, p1_d, p2_d, p3_d, &
131 alpha, p_cur, n) bind(c, name = 'cuda_fusedcg_cpld_update_x')
132 use, intrinsic :: iso_c_binding
133 implicit none
134 type(c_ptr), value :: x1_d, x2_d, x3_d, p1_d, p2_d, p3_d, alpha
135 integer(c_int) :: p_cur, n
136 end subroutine cuda_fusedcg_cpld_update_x
137 end interface
138
139 interface
140 real(c_rp) function cuda_fusedcg_cpld_part2(a1_d, a2_d, a3_d, b_d, &
141 c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n) &
142 bind(c, name = 'cuda_fusedcg_cpld_part2')
143 use, intrinsic :: iso_c_binding
144 import c_rp
145 implicit none
146 type(c_ptr), value :: a1_d, a2_d, a3_d, b_d
147 type(c_ptr), value :: c1_d, c2_d, c3_d, alpha_d
148 real(c_rp) :: alpha
149 integer(c_int) :: n, p_cur
150 end function cuda_fusedcg_cpld_part2
151 end interface
152#elif HAVE_HIP
153 interface
154 subroutine hip_fusedcg_cpld_part1(a1_d, a2_d, a3_d, &
155 b1_d, b2_d, b3_d, tmp_d, n) &
156 bind(c, name = 'hip_fusedcg_cpld_part1')
157 use, intrinsic :: iso_c_binding
158 import c_rp
159 implicit none
160 type(c_ptr), value :: a1_d, a2_d, a3_d, b1_d, b2_d, b3_d, tmp_d
161 integer(c_int) :: n
162 end subroutine hip_fusedcg_cpld_part1
163 end interface
164
165 interface
166 subroutine hip_fusedcg_cpld_update_p(p1_d, p2_d, p3_d, z1_d, z2_d, z3_d, &
167 po1_d, po2_d, po3_d, beta, n) &
168 bind(c, name = 'hip_fusedcg_cpld_update_p')
169 use, intrinsic :: iso_c_binding
170 import c_rp
171 implicit none
172 type(c_ptr), value :: p1_d, p2_d, p3_d, z1_d, z2_d, z3_d
173 type(c_ptr), value :: po1_d, po2_d, po3_d
174 real(c_rp) :: beta
175 integer(c_int) :: n
176 end subroutine hip_fusedcg_cpld_update_p
177 end interface
178
179 interface
180 subroutine hip_fusedcg_cpld_update_x(x1_d, x2_d, x3_d, p1_d, p2_d, p3_d, &
181 alpha, p_cur, n) bind(c, name = 'hip_fusedcg_cpld_update_x')
182 use, intrinsic :: iso_c_binding
183 implicit none
184 type(c_ptr), value :: x1_d, x2_d, x3_d, p1_d, p2_d, p3_d, alpha
185 integer(c_int) :: p_cur, n
186 end subroutine hip_fusedcg_cpld_update_x
187 end interface
188
189 interface
190 real(c_rp) function hip_fusedcg_cpld_part2(a1_d, a2_d, a3_d, b_d, &
191 c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n) &
192 bind(c, name = 'hip_fusedcg_cpld_part2')
193 use, intrinsic :: iso_c_binding
194 import c_rp
195 implicit none
196 type(c_ptr), value :: a1_d, a2_d, a3_d, b_d
197 type(c_ptr), value :: c1_d, c2_d, c3_d, alpha_d
198 real(c_rp) :: alpha
199 integer(c_int) :: n, p_cur
200 end function hip_fusedcg_cpld_part2
201 end interface
202#endif
203
204contains
205
206 subroutine device_fusedcg_cpld_part1(a1_d, a2_d, a3_d, &
207 b1_d, b2_d, b3_d, tmp_d, n)
208 type(c_ptr), value :: a1_d, a2_d, a3_d, b1_d, b2_d, b3_d
209 type(c_ptr), value :: tmp_d
210 integer(c_int) :: n
211#ifdef HAVE_HIP
212 call hip_fusedcg_cpld_part1(a1_d, a2_d, a3_d, b1_d, b2_d, b3_d, tmp_d, n)
213#elif HAVE_CUDA
214 call cuda_fusedcg_cpld_part1(a1_d, a2_d, a3_d, b1_d, b2_d, b3_d, tmp_d, n)
215#else
216 call neko_error('No device backend configured')
217#endif
218 end subroutine device_fusedcg_cpld_part1
219
220 subroutine device_fusedcg_cpld_update_p(p1_d, p2_d, p3_d, z1_d, z2_d, z3_d, &
221 po1_d, po2_d, po3_d, beta, n)
222 type(c_ptr), value :: p1_d, p2_d, p3_d, z1_d, z2_d, z3_d
223 type(c_ptr), value :: po1_d, po2_d, po3_d
224 real(c_rp) :: beta
225 integer(c_int) :: n
226#ifdef HAVE_HIP
227 call hip_fusedcg_cpld_update_p(p1_d, p2_d, p3_d, z1_d, z2_d, z3_d, &
228 po1_d, po2_d, po3_d, beta, n)
229#elif HAVE_CUDA
230 call cuda_fusedcg_cpld_update_p(p1_d, p2_d, p3_d, z1_d, z2_d, z3_d, &
231 po1_d, po2_d, po3_d, beta, n)
232#else
233 call neko_error('No device backend configured')
234#endif
235 end subroutine device_fusedcg_cpld_update_p
236
237 subroutine device_fusedcg_cpld_update_x(x1_d, x2_d, x3_d, &
238 p1_d, p2_d, p3_d, alpha, p_cur, n)
239 type(c_ptr), value :: x1_d, x2_d, x3_d, p1_d, p2_d, p3_d, alpha
240 integer(c_int) :: p_cur, n
241#ifdef HAVE_HIP
242 call hip_fusedcg_cpld_update_x(x1_d, x2_d, x3_d, &
243 p1_d, p2_d, p3_d, alpha, p_cur, n)
244#elif HAVE_CUDA
245 call cuda_fusedcg_cpld_update_x(x1_d, x2_d, x3_d, &
246 p1_d, p2_d, p3_d, alpha, p_cur, n)
247#else
248 call neko_error('No device backend configured')
249#endif
250 end subroutine device_fusedcg_cpld_update_x
251
252 function device_fusedcg_cpld_part2(a1_d, a2_d, a3_d, b_d, &
253 c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n) result(res)
254 type(c_ptr), value :: a1_d, a2_d, a3_d, b_d
255 type(c_ptr), value :: c1_d, c2_d, c3_d, alpha_d
256 real(c_rp) :: alpha
257 integer :: n, p_cur
258 real(kind=rp) :: res
259 integer :: ierr
260#ifdef HAVE_HIP
261 res = hip_fusedcg_cpld_part2(a1_d, a2_d, a3_d, b_d, &
262 c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n)
263#elif HAVE_CUDA
264 res = cuda_fusedcg_cpld_part2(a1_d, a2_d, a3_d, b_d, &
265 c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n)
266#else
267 call neko_error('No device backend configured')
268#endif
269
270#ifndef HAVE_DEVICE_MPI
271 if (pe_size .gt. 1) then
272 call mpi_allreduce(mpi_in_place, res, 1, &
273 mpi_real_precision, mpi_sum, neko_comm, ierr)
274 end if
275#endif
276
277 end function device_fusedcg_cpld_part2
278
280 subroutine fusedcg_cpld_device_init(this, n, max_iter, M, &
281 rel_tol, abs_tol, monitor)
282 class(fusedcg_cpld_device_t), target, intent(inout) :: this
283 class(pc_t), optional, intent(in), target :: M
284 integer, intent(in) :: n
285 integer, intent(in) :: max_iter
286 real(kind=rp), optional, intent(in) :: rel_tol
287 real(kind=rp), optional, intent(in) :: abs_tol
288 logical, optional, intent(in) :: monitor
289 type(c_ptr) :: ptr
290 integer(c_size_t) :: p_size
291 integer :: i
292
293 call this%free()
294
295 allocate(this%w1(n))
296 allocate(this%w2(n))
297 allocate(this%w3(n))
298 allocate(this%r1(n))
299 allocate(this%r2(n))
300 allocate(this%r3(n))
301 allocate(this%z1(n))
302 allocate(this%z2(n))
303 allocate(this%z3(n))
304 allocate(this%tmp(n))
305 allocate(this%p1(n, device_fusedcg_cpld_p_space))
306 allocate(this%p2(n, device_fusedcg_cpld_p_space))
307 allocate(this%p3(n, device_fusedcg_cpld_p_space))
308 allocate(this%p1_d(device_fusedcg_cpld_p_space))
309 allocate(this%p2_d(device_fusedcg_cpld_p_space))
310 allocate(this%p3_d(device_fusedcg_cpld_p_space))
311 allocate(this%alpha(device_fusedcg_cpld_p_space))
312
313 if (present(m)) then
314 this%M => m
315 end if
316
317 call device_map(this%w1, this%w1_d, n)
318 call device_map(this%w2, this%w2_d, n)
319 call device_map(this%w3, this%w3_d, n)
320 call device_map(this%r1, this%r1_d, n)
321 call device_map(this%r2, this%r2_d, n)
322 call device_map(this%r3, this%r3_d, n)
323 call device_map(this%z1, this%z1_d, n)
324 call device_map(this%z2, this%z2_d, n)
325 call device_map(this%z3, this%z3_d, n)
326 call device_map(this%tmp, this%tmp_d, n)
327 call device_map(this%alpha, this%alpha_d, device_fusedcg_cpld_p_space)
329 this%p1_d(i) = c_null_ptr
330 call device_map(this%p1(:,i), this%p1_d(i), n)
331
332 this%p2_d(i) = c_null_ptr
333 call device_map(this%p2(:,i), this%p2_d(i), n)
334
335 this%p3_d(i) = c_null_ptr
336 call device_map(this%p3(:,i), this%p3_d(i), n)
337 end do
338
339 p_size = c_sizeof(c_null_ptr) * (device_fusedcg_cpld_p_space)
340 call device_alloc(this%p1_d_d, p_size)
341 call device_alloc(this%p2_d_d, p_size)
342 call device_alloc(this%p3_d_d, p_size)
343 ptr = c_loc(this%p1_d)
344 call device_memcpy(ptr, this%p1_d_d, p_size, &
345 host_to_device, sync=.false.)
346 ptr = c_loc(this%p2_d)
347 call device_memcpy(ptr, this%p2_d_d, p_size, &
348 host_to_device, sync=.false.)
349 ptr = c_loc(this%p3_d)
350 call device_memcpy(ptr, this%p3_d_d, p_size, &
351 host_to_device, sync=.false.)
352 if (present(rel_tol) .and. present(abs_tol) .and. present(monitor)) then
353 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
354 else if (present(rel_tol) .and. present(abs_tol)) then
355 call this%ksp_init(max_iter, rel_tol, abs_tol)
356 else if (present(monitor) .and. present(abs_tol)) then
357 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
358 else if (present(rel_tol) .and. present(monitor)) then
359 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
360 else if (present(rel_tol)) then
361 call this%ksp_init(max_iter, rel_tol = rel_tol)
362 else if (present(abs_tol)) then
363 call this%ksp_init(max_iter, abs_tol = abs_tol)
364 else if (present(monitor)) then
365 call this%ksp_init(max_iter, monitor = monitor)
366 else
367 call this%ksp_init(max_iter)
368 end if
369
370 call device_event_create(this%gs_event1, 2)
371 call device_event_create(this%gs_event2, 2)
372 call device_event_create(this%gs_event3, 2)
373
374 end subroutine fusedcg_cpld_device_init
375
378 class(fusedcg_cpld_device_t), intent(inout) :: this
379 integer :: i
380
381 call this%ksp_free()
382
383 if (allocated(this%w1)) then
384 if (c_associated(this%w1_d)) then
385 call device_unmap(this%w1, this%w1_d)
386 end if
387 deallocate(this%w1)
388 end if
389
390 if (allocated(this%w2)) then
391 if (c_associated(this%w2_d)) then
392 call device_unmap(this%w2, this%w2_d)
393 end if
394 deallocate(this%w2)
395 end if
396
397 if (allocated(this%w3)) then
398 if (c_associated(this%w3_d)) then
399 call device_unmap(this%w3, this%w3_d)
400 end if
401 deallocate(this%w3)
402 end if
403
404 if (allocated(this%r1)) then
405 if (c_associated(this%r1_d)) then
406 call device_unmap(this%r1, this%r1_d)
407 end if
408 deallocate(this%r1)
409 end if
410
411 if (allocated(this%r2)) then
412 if (c_associated(this%r2_d)) then
413 call device_unmap(this%r2, this%r2_d)
414 end if
415 deallocate(this%r2)
416 end if
417
418 if (allocated(this%r3)) then
419 if (c_associated(this%r3_d)) then
420 call device_unmap(this%r3, this%r3_d)
421 end if
422 deallocate(this%r3)
423 end if
424
425 if (allocated(this%z1)) then
426 if (c_associated(this%z1_d)) then
427 call device_unmap(this%z1, this%z1_d)
428 end if
429 deallocate(this%z1)
430 end if
431
432 if (allocated(this%z2)) then
433 if (c_associated(this%z2_d)) then
434 call device_unmap(this%z2, this%z2_d)
435 end if
436 deallocate(this%z2)
437 end if
438
439 if (allocated(this%z3)) then
440 if (c_associated(this%z3_d)) then
441 call device_unmap(this%z3, this%z3_d)
442 end if
443 deallocate(this%z3)
444 end if
445
446 if (allocated(this%tmp)) then
447 if (c_associated(this%tmp_d)) then
448 call device_unmap(this%tmp, this%tmp_d)
449 end if
450 deallocate(this%tmp)
451 end if
452
453 if (allocated(this%alpha)) then
454 if (c_associated(this%alpha_d)) then
455 call device_unmap(this%alpha, this%alpha_d)
456 end if
457 deallocate(this%alpha)
458 end if
459
460 if (allocated(this%p1)) then
461 if (allocated(this%p1_d)) then
463 if (c_associated(this%p1_d(i))) then
464 call device_unmap(this%p1(:,i), this%p1_d(i))
465 end if
466 end do
467 end if
468 deallocate(this%p1)
469 end if
470
471 if (allocated(this%p2)) then
472 if (allocated(this%p2_d)) then
474 if (c_associated(this%p2_d(i))) then
475 call device_unmap(this%p2(:,i), this%p2_d(i))
476 end if
477 end do
478 end if
479 deallocate(this%p2)
480 end if
481
482 if (allocated(this%p3)) then
483 if (allocated(this%p3_d)) then
485 if (c_associated(this%p3_d(i))) then
486 call device_unmap(this%p3(:,i), this%p3_d(i))
487 end if
488 end do
489 end if
490 deallocate(this%p3)
491 end if
492
493 if (allocated(this%p1_d)) then
494 deallocate(this%p1_d)
495 end if
496
497 if (allocated(this%p2_d)) then
498 deallocate(this%p2_d)
499 end if
500
501 if (allocated(this%p3_d)) then
502 deallocate(this%p3_d)
503 end if
504
505 if (c_associated(this%p1_d_d)) then
506 call device_free(this%p1_d_d)
507 end if
508
509 if (c_associated(this%p2_d_d)) then
510 call device_free(this%p2_d_d)
511 end if
512
513 if (c_associated(this%p3_d_d)) then
514 call device_free(this%p3_d_d)
515 end if
516
517 nullify(this%M)
518
519 if (c_associated(this%gs_event1)) then
520 call device_event_destroy(this%gs_event1)
521 end if
522
523 if (c_associated(this%gs_event2)) then
524 call device_event_destroy(this%gs_event2)
525 end if
526
527 if (c_associated(this%gs_event3)) then
528 call device_event_destroy(this%gs_event3)
529 end if
530
531 end subroutine fusedcg_cpld_device_free
532
534 function fusedcg_cpld_device_solve_coupled(this, Ax, x, y, z, fx, fy, fz, &
535 n, coef, blstx, blsty, blstz, gs_h, niter) result(ksp_results)
536 class(fusedcg_cpld_device_t), intent(inout) :: this
537 class(ax_t), intent(in) :: ax
538 type(field_t), intent(inout) :: x
539 type(field_t), intent(inout) :: y
540 type(field_t), intent(inout) :: z
541 integer, intent(in) :: n
542 real(kind=rp), dimension(n), intent(in) :: fx
543 real(kind=rp), dimension(n), intent(in) :: fy
544 real(kind=rp), dimension(n), intent(in) :: fz
545 type(coef_t), intent(inout) :: coef
546 type(bc_list_t), intent(inout) :: blstx
547 type(bc_list_t), intent(inout) :: blsty
548 type(bc_list_t), intent(inout) :: blstz
549 type(gs_t), intent(inout) :: gs_h
550 type(ksp_monitor_t), dimension(3) :: ksp_results
551 integer, optional, intent(in) :: niter
552 integer :: iter, max_iter, ierr, i, p_cur, p_prev
553 real(kind=rp) :: rnorm, rtr, norm_fac, rtz1, rtz2
554 real(kind=rp) :: pap, beta
555 type(c_ptr) :: fx_d
556 type(c_ptr) :: fy_d
557 type(c_ptr) :: fz_d
558
559 fx_d = device_get_ptr(fx)
560 fy_d = device_get_ptr(fy)
561 fz_d = device_get_ptr(fz)
562
563 if (present(niter)) then
564 max_iter = niter
565 else
566 max_iter = ksp_max_iter
567 end if
568 norm_fac = 1.0_rp / sqrt(coef%volume)
569
570 associate(w1 => this%w1, w2 => this%w2, w3 => this%w3, r1 => this%r1, &
571 r2 => this%r2, r3 => this%r3, p1 => this%p1, p2 => this%p2, &
572 p3 => this%p3, z1 => this%z1, z2 => this%z2, z3 => this%z3, &
573 tmp_d => this%tmp_d, alpha => this%alpha, alpha_d => this%alpha_d, &
574 w1_d => this%w1_d, w2_d => this%w2_d, w3_d => this%w3_d, &
575 r1_d => this%r1_d, r2_d => this%r2_d, r3_d => this%r3_d, &
576 z1_d => this%z1_d, z2_d => this%z2_d, z3_d => this%z3_d, &
577 p1_d => this%p1_d, p2_d => this%p2_d, p3_d => this%p3_d, &
578 p1_d_d => this%p1_d_d, p2_d_d => this%p2_d_d, p3_d_d => this%p3_d_d)
579
580 rtz1 = 1.0_rp
582 p_cur = 1
583
584
585 call device_rzero(x%x_d, n)
586 call device_rzero(y%x_d, n)
587 call device_rzero(z%x_d, n)
588 call device_rzero(p1_d(1), n)
589 call device_rzero(p2_d(1), n)
590 call device_rzero(p3_d(1), n)
591 call device_copy(r1_d, fx_d, n)
592 call device_copy(r2_d, fy_d, n)
593 call device_copy(r3_d, fz_d, n)
594
595 call device_fusedcg_cpld_part1(r1_d, r2_d, r3_d, r1_d, &
596 r2_d, r3_d, tmp_d, n)
597
598 rtr = device_glsc2(tmp_d, coef%mult_d, n)
599
600 rnorm = sqrt(rtr)*norm_fac
601 ksp_results%res_start = rnorm
602 ksp_results%res_final = rnorm
603 ksp_results(1)%iter = 0
604 ksp_results(2:3)%iter = -1
605 if (abscmp(rnorm, 0.0_rp)) then
606 ksp_results%converged = .true.
607 return
608 end if
609 call this%monitor_start('fcpldCG')
610 do iter = 1, max_iter
611 call this%M%solve(z1, r1, n)
612 call this%M%solve(z2, r2, n)
613 call this%M%solve(z3, r3, n)
614 rtz2 = rtz1
615 call device_fusedcg_cpld_part1(z1_d, z2_d, z3_d, &
616 r1_d, r2_d, r3_d, tmp_d, n)
617 rtz1 = device_glsc2(tmp_d, coef%mult_d, n)
618
619 beta = rtz1 / rtz2
620 if (iter .eq. 1) beta = 0.0_rp
621
622 call device_fusedcg_cpld_update_p(p1_d(p_cur), p2_d(p_cur), p3_d(p_cur), &
623 z1_d, z2_d, z3_d, p1_d(p_prev), p2_d(p_prev), p3_d(p_prev), beta, n)
624
625 call ax%compute_vector(w1, w2, w3, &
626 p1(1, p_cur), p2(1, p_cur), p3(1, p_cur), coef, x%msh, x%Xh)
627
628 call rotate_cyc(w1_d, w2_d, w3_d, 1, coef)
629 call gs_h%op(w1, w2, w3, n, gs_op_add, this%gs_event1)
630 call device_event_sync(this%gs_event1)
631 call blstx%apply(w1, n)
632 call blsty%apply(w2, n)
633 call blstz%apply(w3, n)
634 call rotate_cyc(w1_d, w2_d, w3_d, 0, coef)
635
636 call device_fusedcg_cpld_part1(w1_d, w2_d, w3_d, p1_d(p_cur), &
637 p2_d(p_cur), p3_d(p_cur), tmp_d, n)
638
639 pap = device_glsc2(tmp_d, coef%mult_d, n)
640
641 alpha(p_cur) = rtz1 / pap
642 rtr = device_fusedcg_cpld_part2(r1_d, r2_d, r3_d, coef%mult_d, &
643 w1_d, w2_d, w3_d, alpha_d, alpha(p_cur), p_cur, n)
644 rnorm = sqrt(rtr)*norm_fac
645 call this%monitor_iter(iter, rnorm)
646 if ((p_cur .eq. device_fusedcg_cpld_p_space) .or. &
647 (rnorm .lt. this%abs_tol) .or. iter .eq. max_iter) then
648 call device_fusedcg_cpld_update_x(x%x_d, y%x_d, z%x_d, &
649 p1_d_d, p2_d_d, p3_d_d, alpha_d, p_cur, n)
650 p_prev = p_cur
651 p_cur = 1
652 if (rnorm .lt. this%abs_tol) exit
653 else
654 p_prev = p_cur
655 p_cur = p_cur + 1
656 end if
657 end do
658 call this%monitor_stop()
659 ksp_results%res_final = rnorm
660 ksp_results%iter = iter
661 ksp_results%converged = this%is_converged(iter, rnorm)
662
663 end associate
664
666
668 function fusedcg_cpld_device_solve(this, Ax, x, f, n, coef, blst, &
669 gs_h, niter) result(ksp_results)
670 class(fusedcg_cpld_device_t), intent(inout) :: this
671 class(ax_t), intent(in) :: ax
672 type(field_t), intent(inout) :: x
673 integer, intent(in) :: n
674 real(kind=rp), dimension(n), intent(in) :: f
675 type(coef_t), intent(inout) :: coef
676 type(bc_list_t), intent(inout) :: blst
677 type(gs_t), intent(inout) :: gs_h
678 type(ksp_monitor_t) :: ksp_results
679 integer, optional, intent(in) :: niter
680
681 ! Throw and error
682 call neko_error('The cpldcg solver is only defined for coupled solves')
683
684 ksp_results%res_final = 0.0
685 ksp_results%iter = 0
686 ksp_results%converged = .false.
687
688 end function fusedcg_cpld_device_solve
689
690end module fusedcg_cpld_device
__device__ T solve(const T u, const T y, const T guess, const T nu, const T kappa, const T B)
void hip_fusedcg_cpld_update_x(void *x1, void *x2, void *x3, void *p1, void *p2, void *p3, void *alpha, int *p_cur, int *n)
void hip_fusedcg_cpld_update_p(void *p1, void *p2, void *p3, void *z1, void *z2, void *z3, void *po1, void *po2, void *po3, real *beta, int *n)
real hip_fusedcg_cpld_part2(void *a1, void *a2, void *a3, void *b, void *c1, void *c2, void *c3, void *alpha_d, real *alpha, int *p_cur, int *n)
void hip_fusedcg_cpld_part1(void *a1, void *a2, void *a3, void *b1, void *b2, void *b3, void *tmp, int *n)
Return the device pointer for an associated Fortran array.
Definition device.F90:108
Map a Fortran array to a device (allocate and associate)
Definition device.F90:78
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:84
Apply cyclic boundary condition to a vector field.
Defines a Matrix-vector product.
Definition ax.f90:34
Defines a list of bc_t.
Definition bc_list.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
subroutine, public device_rzero(a_d, n, strm)
Zero a real vector.
subroutine, public device_copy(a_d, b_d, n, strm)
Copy a vector .
real(kind=rp) function, public device_glsc2(a_d, b_d, n, strm)
Weighted inner product .
Device abstraction, common interface for various accelerators.
Definition device.F90:34
subroutine, public device_event_sync(event)
Synchronize an event.
Definition device.F90:1637
integer, parameter, public host_to_device
Definition device.F90:48
subroutine, public device_free(x_d)
Deallocate memory on the device.
Definition device.F90:238
subroutine, public device_event_destroy(event)
Destroy a device event.
Definition device.F90:1593
subroutine, public device_alloc(x_d, s)
Allocate memory on the device.
Definition device.F90:207
subroutine, public device_event_create(event, flags)
Create a device event queue.
Definition device.F90:1559
Defines a field.
Definition field.f90:34
Defines a fused Conjugate Gradient method for accelerators.
subroutine device_fusedcg_cpld_update_x(x1_d, x2_d, x3_d, p1_d, p2_d, p3_d, alpha, p_cur, n)
type(ksp_monitor_t) function fusedcg_cpld_device_solve(this, ax, x, f, n, coef, blst, gs_h, niter)
Pipelined PCG solve.
subroutine fusedcg_cpld_device_free(this)
Deallocate a pipelined PCG solver.
subroutine device_fusedcg_cpld_update_p(p1_d, p2_d, p3_d, z1_d, z2_d, z3_d, po1_d, po2_d, po3_d, beta, n)
subroutine fusedcg_cpld_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a fused PCG solver.
real(kind=rp) function device_fusedcg_cpld_part2(a1_d, a2_d, a3_d, b_d, c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n)
type(ksp_monitor_t) function, dimension(3) fusedcg_cpld_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, blstx, blsty, blstz, gs_h, niter)
Pipelined PCG solve coupled solve.
integer, parameter device_fusedcg_cpld_p_space
subroutine device_fusedcg_cpld_part1(a1_d, a2_d, a3_d, b1_d, b2_d, b3_d, tmp_d, n)
Gather-scatter.
Implements the base abstract type for Krylov solvers plus helper types.
Definition krylov.f90:34
integer, parameter, public ksp_max_iter
Maximum number of iters.
Definition krylov.f90:51
Definition math.f90:60
real(kind=rp) function, public glsc3(a, b, c, n)
Weighted inner product .
Definition math.f90:1287
subroutine, public copy(a, b, n)
Copy a vector .
Definition math.f90:291
subroutine, public rzero(a, n)
Zero a real vector.
Definition math.f90:235
integer, parameter, public c_rp
Definition num_types.f90:13
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:12
Operators.
Definition operators.f90:34
Krylov preconditioner.
Definition precon.f90:34
Utilities.
Definition utils.f90:35
Base type for a matrix-vector product providing .
Definition ax.f90:43
A list of allocatable `bc_t`. Follows the standard interface of lists.
Definition bc_list.f90:49
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:63
Fused preconditioned conjugate gradient method.
Gather-scatter kernel.
Type for storing initial and final residuals in a Krylov solver.
Definition krylov.f90:56
Base abstract type for a canonical Krylov method, solving .
Definition krylov.f90:73
Defines a canonical Krylov preconditioner.
Definition precon.f90:40