ArthurY's picture
update source
c3d0544
# SPDX-FileCopyrightText: Copyright (c) 2023 - 2025 NVIDIA CORPORATION & AFFILIATES.
# SPDX-FileCopyrightText: All rights reserved.
# SPDX-License-Identifier: Apache-2.0
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
# You may obtain a copy of the License at
#
# http://www.apache.org/licenses/LICENSE-2.0
#
# Unless required by applicable law or agreed to in writing, software
# distributed under the License is distributed on an "AS IS" BASIS,
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
# See the License for the specific language governing permissions and
# limitations under the License.
try:
import warp as wp
except ImportError:
print(
"""NVIDIA WARP is required for this datapipe. This package is under the
NVIDIA Source Code License (NVSCL). To install use:
pip install warp-lang
"""
)
raise SystemExit(1)
from .indexing import index_clamped_edges_batched_2d, index_zero_edges_batched_2d
@wp.kernel
def darcy_mgrid_jacobi_iterative_batched_2d(
darcy0: wp.array3d(dtype=float),
darcy1: wp.array3d(dtype=float),
permeability: wp.array3d(dtype=float),
source: float,
lx: int,
ly: int,
dx: float,
mgrid_reduction_factor: int,
): # pragma: no cover
"""Mult-grid jacobi step for Darcy equation.
Parameters
----------
darcy0 : wp.array3d
Darcy solution previous step
darcy1 : wp.array3d
Darcy solution for next step
permeability : wp.array3d
Permeability field for Darcy equation
source : float
Source value for Darcy equation
lx : int
Length of domain in x dim
ly : int
Length of domain in y dim
dx : float
Grid cell size
mgrid_reduction_factor : int
Current multi-grid running at
"""
# get index
b, x, y = wp.tid()
# update index from grid reduction factor
gx = mgrid_reduction_factor * x + (mgrid_reduction_factor - 1)
gy = mgrid_reduction_factor * y + (mgrid_reduction_factor - 1)
gdx = dx * wp.float32(mgrid_reduction_factor)
# compute darcy stensil
d_0_1 = index_zero_edges_batched_2d(
darcy0, b, gx - mgrid_reduction_factor, gy, lx, ly
)
d_2_1 = index_zero_edges_batched_2d(
darcy0, b, gx + mgrid_reduction_factor, gy, lx, ly
)
d_1_0 = index_zero_edges_batched_2d(
darcy0, b, gx, gy - mgrid_reduction_factor, lx, ly
)
d_1_2 = index_zero_edges_batched_2d(
darcy0, b, gx, gy + mgrid_reduction_factor, lx, ly
)
# compute permeability stensil
p_1_1 = index_clamped_edges_batched_2d(permeability, b, gx, gy, lx, ly)
p_0_1 = index_clamped_edges_batched_2d(
permeability, b, gx - mgrid_reduction_factor, gy, lx, ly
)
p_2_1 = index_clamped_edges_batched_2d(
permeability, b, gx + mgrid_reduction_factor, gy, lx, ly
)
p_1_0 = index_clamped_edges_batched_2d(
permeability, b, gx, gy - mgrid_reduction_factor, lx, ly
)
p_1_2 = index_clamped_edges_batched_2d(
permeability, b, gx, gy + mgrid_reduction_factor, lx, ly
)
# compute terms
dx_squared = gdx * gdx
t_1 = p_1_1 * (d_0_1 + d_2_1 + d_1_0 + d_1_2) / dx_squared
t_2 = ((p_2_1 - p_0_1) * (d_2_1 - d_0_1)) / (2.0 * gdx)
t_3 = ((p_1_2 - p_1_0) * (d_1_2 - d_1_0)) / (2.0 * gdx)
# jacobi iterative method
d_star = (t_1 + t_2 + t_3 + source) / (p_1_1 * 4.0 / dx_squared)
# buffers get swapped each iteration
darcy1[b, gx, gy] = d_star
@wp.kernel
def mgrid_inf_residual_batched_2d(
phi0: wp.array3d(dtype=float),
phi1: wp.array3d(dtype=float),
inf_res: wp.array(dtype=float),
mgrid_reduction_factor: int,
): # pragma: no cover
"""Infinity norm for checking multi-grid solutions.
Parameters
----------
phi0 : wp.array3d
Previous solution
phi1 : wp.array3d
Current solution
inf_res : wp.array
Array to hold infinity norm value in
mgrid_reduction_factor : int
Current multi-grid running at
"""
b, x, y = wp.tid()
gx = mgrid_reduction_factor * x + (mgrid_reduction_factor - 1)
gy = mgrid_reduction_factor * y + (mgrid_reduction_factor - 1)
wp.atomic_max(inf_res, 0, wp.abs(phi0[b, gx, gy] - phi1[b, gx, gy]))