Skip to content

Commit 010ee1f

Browse files
Add amrex::Math::rsqrt (#4777)
This PR adds a function to compute the inverse square root of a single- or double-precision value. On GPUs this should be faster than doing the inverse and square root separately.
1 parent a7f0aed commit 010ee1f

File tree

1 file changed

+28
-0
lines changed

1 file changed

+28
-0
lines changed

Src/Base/AMReX_Math.H

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -325,6 +325,34 @@ T comp_ellint_2 (T k)
325325
return Kcomp*sum_val;
326326
}
327327

328+
//! Return inverse square root of x
329+
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
330+
double rsqrt (double x)
331+
{
332+
double r;
333+
#if defined(AMREX_USE_SYCL)
334+
AMREX_IF_ON_DEVICE(( r = sycl::rsqrt(x); ))
335+
#else
336+
AMREX_IF_ON_DEVICE(( r = ::rsqrt(x); ))
337+
#endif
338+
AMREX_IF_ON_HOST(( r = 1. / std::sqrt(x); ))
339+
return r;
340+
}
341+
342+
//! Return inverse square root of x
343+
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
344+
float rsqrt (float x)
345+
{
346+
float r;
347+
#if defined(AMREX_USE_SYCL)
348+
AMREX_IF_ON_DEVICE(( r = sycl::rsqrt(x); ))
349+
#else
350+
AMREX_IF_ON_DEVICE(( r = ::rsqrtf(x); ))
351+
#endif
352+
AMREX_IF_ON_HOST(( r = 1.F / std::sqrt(x); ))
353+
return r;
354+
}
355+
328356
/***************************************************************************************************
329357
* Copyright (c) 2017 - 2024 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
330358
* SPDX-License-Identifier: BSD-3-Clause

0 commit comments

Comments
 (0)