37 use json_module,
only : json_file
46 use,
intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
52 real(kind=
rp),
allocatable :: d(:,:,:,:)
53 type(c_ptr) :: d_d = c_null_ptr
54 type(
gs_t),
pointer :: gs_h
57 type(c_ptr) :: gs_event = c_null_ptr
60 procedure, pass(this) :: init_from_components => &
69 G11_d, G22_d, G33_d, G12_d, G13_d, G23_d, nelv, lx) &
70 bind(c, name=
'hip_jacobi_update')
71 use,
intrinsic :: iso_c_binding
72 type(c_ptr),
value :: d_d, dxt_d, dyt_d, dzt_d
73 type(c_ptr),
value :: G11_d, G22_d, G33_d, G12_d, G13_d, G23_d
74 integer(c_int) :: nelv, lx
80 G11_d, G22_d, G33_d, G12_d, G13_d, G23_d, nelv, lx) &
81 bind(c, name=
'cuda_jacobi_update')
82 use,
intrinsic :: iso_c_binding
83 type(c_ptr),
value :: d_d, dxt_d, dyt_d, dzt_d
84 type(c_ptr),
value :: G11_d, G22_d, G33_d, G12_d, G13_d, G23_d
85 integer(c_int) :: nelv, lx
91 G11_d, G22_d, G33_d, G12_d, G13_d, G23_d, nelv, lx) &
92 bind(c, name=
'opencl_jacobi_update')
93 use,
intrinsic :: iso_c_binding
94 type(c_ptr),
value :: d_d, dxt_d, dyt_d, dzt_d
95 type(c_ptr),
value :: G11_d, G22_d, G33_d, G12_d, G13_d, G23_d
96 integer(c_int) :: nelv, lx
102 G11_d, G22_d, G33_d, G12_d, G13_d, G23_d, nelv, lx) &
103 bind(c, name=
'metal_jacobi_update')
104 use,
intrinsic :: iso_c_binding
105 type(c_ptr),
value :: d_d, dxt_d, dyt_d, dzt_d
106 type(c_ptr),
value :: G11_d, G22_d, G33_d, G12_d, G13_d, G23_d
107 integer(c_int) :: nelv, lx
117 type(
coef_t),
intent(in),
target :: coef
118 type(
bc_list_t),
intent(inout),
target :: bclst
119 type(json_file),
intent(inout) :: json
121 call this%init_from_components(coef, coef%dof, coef%gs_h)
127 type(
coef_t),
intent(in),
target :: coef
128 type(
dofmap_t),
intent(in),
target :: dof
129 type(
gs_t),
intent(inout),
target :: gs_h
137 allocate(this%d(dof%Xh%lx,dof%Xh%ly,dof%Xh%lz, dof%msh%nelv))
139 call device_map(this%d, this%d_d,
size(this%d))
150 if (
allocated(this%d))
then
151 if (c_associated(this%d_d))
then
157 if (c_associated(this%gs_event))
then
169 integer,
intent(in) :: n
171 real(kind=
rp),
dimension(n),
intent(inout) :: z
172 real(kind=
rp),
dimension(n),
intent(inout) :: r
173 type(c_ptr) :: z_d, r_d
184 integer :: lz, ly, lx
185 associate(dof => this%dof, coef => this%coef, xh => this%dof%Xh, &
186 gs_h => this%gs_h, nelv => this%dof%msh%nelv)
194 coef%G11_d, coef%G22_d, coef%G33_d, &
195 coef%G12_d, coef%G13_d, coef%G23_d, &
199 coef%G11_d, coef%G22_d, coef%G33_d, &
200 coef%G12_d, coef%G13_d, coef%G23_d, &
204 coef%G11_d, coef%G22_d, coef%G33_d, &
205 coef%G12_d, coef%G13_d, coef%G23_d, &
209 coef%G11_d, coef%G22_d, coef%G33_d, &
210 coef%G12_d, coef%G13_d, coef%G23_d, &
214 call device_col2(this%d_d, coef%h1_d, coef%dof%size())
217 call device_addcol3(this%d_d, coef%h2_d, coef%B_d, coef%dof%size())
220 call gs_h%op(this%d, dof%size(), gs_op_add, this%gs_event)
__device__ T solve(const T u, const T y, const T guess, const T nu, const T kappa, const T B)
Return the device pointer for an associated Fortran array.
Map a Fortran array to a device (allocate and associate)
Unmap a Fortran array from a device (deassociate and free)
Jacobi preconditioner accelerator backend.
subroutine device_jacobi_free(this)
subroutine device_jacobi_init_from_components(this, coef, dof, gs_h)
subroutine device_jacobi_update(this)
subroutine device_jacobi_init(this, coef, bclst, json)
Initialize from the case file. The boundary conditions and the JSON dictionary are not needed by this...
subroutine device_jacobi_solve(this, z, r, n)
The jacobi preconditioner where .
subroutine, public device_addcol3(a_d, b_d, c_d, n, strm)
Returns .
subroutine, public device_invcol1(a_d, n, strm)
Invert a vector .
subroutine, public device_col2(a_d, b_d, n, strm)
Vector multiplication .
subroutine, public device_col3(a_d, b_d, c_d, n, strm)
Vector multiplication with 3 vectors .
Device abstraction, common interface for various accelerators.
subroutine, public device_event_sync(event)
Synchronize an event.
subroutine, public device_event_destroy(event)
Destroy a device event.
subroutine, public device_event_create(event, flags)
Create a device event queue.
Defines a mapping of the degrees of freedom.
integer, parameter, public rp
Global precision used in computations.
A list of allocatable `bc_t`. Follows the standard interface of lists.
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Defines a jacobi preconditioner.
Defines a canonical Krylov preconditioner.