/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hip/checkouts/llvm-project/clang/lib/Headers/__clang_hip_math.h Source File

• llvm-project
• clang
• lib
• Headers

__clang_hip_math.h

Go to the documentation of this file.

1 /*===---- __clang_hip_math.h - Device-side HIP math support ----------------===

2  *

3  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

4  * See https://llvm.org/LICENSE.txt for license information.

5  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

6  *

7  *===-----------------------------------------------------------------------===

8  */

9 #ifndef __CLANG_HIP_MATH_H__

10 #define __CLANG_HIP_MATH_H__

11 

12 #if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)

13 #error "This file is for HIP and OpenMP AMDGCN device compilation only."

14 #endif

15 

16 #if !defined(__HIPCC_RTC__)

17 #include <limits.h>

18 #include <stdint.h>

19 #ifdef __OPENMP_AMDGCN__

20 #include <omp.h>

21 #endif

22 #endif // !defined(__HIPCC_RTC__)

23 

24 #pragma push_macro("__DEVICE__")

25 

26 #ifdef __OPENMP_AMDGCN__

27 #define __DEVICE__ static inline __attribute__((always_inline, nothrow))

28 #else

29 #define __DEVICE__ static __device__ inline __attribute__((always_inline))

30 #endif

31 

32 // Device library provides fast low precision and slow full-recision

33 // implementations for some functions. Which one gets selected depends on

34 // __CLANG_GPU_APPROX_TRANSCENDENTALS__ which gets defined by clang if

35 // -ffast-math or -fgpu-approx-transcendentals are in effect.

36 #pragma push_macro("__FAST_OR_SLOW")

37 #if defined(__CLANG_GPU_APPROX_TRANSCENDENTALS__)

38 #define __FAST_OR_SLOW(fast, slow) fast

39 #else

40 #define __FAST_OR_SLOW(fast, slow) slow

41 #endif

42 

43 // A few functions return bool type starting only in C++11.

44 #pragma push_macro("__RETURN_TYPE")

45 #ifdef __OPENMP_AMDGCN__

46 #define __RETURN_TYPE int

47 #else

48 #if defined(__cplusplus)

49 #define __RETURN_TYPE bool

50 #else

51 #define __RETURN_TYPE int

52 #endif

53 #endif // __OPENMP_AMDGCN__

54 

55 #if defined (__cplusplus) && __cplusplus < 201103L

56 // emulate static_assert on type sizes

57 template<bool>

58 struct __compare_result{};

59 template<>

60 struct __compare_result<true> {

61  static const __device__ bool valid;

62 };

63 

64 __DEVICE__

65 void __suppress_unused_warning(bool b){};

66 template <unsigned int S, unsigned int T>

67 __DEVICE__ void __static_assert_equal_size() {

68  __suppress_unused_warning(__compare_result<S == T>::valid);

69 }

70 

71 #define __static_assert_type_size_equal(A, B) \

72  __static_assert_equal_size<A,B>()

73 

74 #else

75 #define __static_assert_type_size_equal(A,B) \

76  static_assert((A) == (B), "")

77 

78 #endif

79 

87 

91 

93 __DEVICE__

94 uint64_t __make_mantissa_base8(const char *__tagp __attribute__((nonnull))) {

95  uint64_t __r = 0;

96  while (*__tagp != '\0') {

97  char __tmp = *__tagp;

98 

99  if (__tmp >= '0' && __tmp <= '7')

100  __r = (__r * 8u) + __tmp - '0';

101  else

102  return 0;

103 

104  ++__tagp;

105  }

106 

107  return __r;

108 }

110 __DEVICE__

111 uint64_t __make_mantissa_base10(const char *__tagp __attribute__((nonnull))) {

112  uint64_t __r = 0;

113  while (*__tagp != '\0') {

114  char __tmp = *__tagp;

115 

116  if (__tmp >= '0' && __tmp <= '9')

117  __r = (__r * 10u) + __tmp - '0';

118  else

119  return 0;

120 

121  ++__tagp;

122  }

123 

124  return __r;

125 }

127 __DEVICE__

128 uint64_t __make_mantissa_base16(const char *__tagp __attribute__((nonnull))) {

129  uint64_t __r = 0;

130  while (*__tagp != '\0') {

131  char __tmp = *__tagp;

132 

133  if (__tmp >= '0' && __tmp <= '9')

134  __r = (__r * 16u) + __tmp - '0';

135  else if (__tmp >= 'a' && __tmp <= 'f')

136  __r = (__r * 16u) + __tmp - 'a' + 10;

137  else if (__tmp >= 'A' && __tmp <= 'F')

138  __r = (__r * 16u) + __tmp - 'A' + 10;

139  else

140  return 0;

141 

142  ++__tagp;

143  }

144 

145  return __r;

146 }

148 __DEVICE__

149 uint64_t __make_mantissa(const char *__tagp __attribute__((nonnull))) {

150  if (*__tagp == '0') {

151  ++__tagp;

152 

153  if (*__tagp == 'x' || *__tagp == 'X')

154  return __make_mantissa_base16(__tagp);

155  else

156  return __make_mantissa_base8(__tagp);

157  }

158 

159  return __make_mantissa_base10(__tagp);

160 }

162 

166 

167 // BEGIN FLOAT

168 

169 // BEGIN INTRINSICS

170 

172 __DEVICE__

173 float __cosf(float __x) { return __ocml_native_cos_f32(__x); }

175 __DEVICE__

176 float __exp10f(float __x) {

177  const float __log2_10 = 0x1.a934f0p+1f;

178  return __builtin_amdgcn_exp2f(__log2_10 * __x);

179 }

181 __DEVICE__

182 float __expf(float __x) {

183  const float __log2_e = 0x1.715476p+0;

184  return __builtin_amdgcn_exp2f(__log2_e * __x);

185 }

186 

187 #if defined OCML_BASIC_ROUNDED_OPERATIONS

189 __DEVICE__

190 float __fadd_rd(float __x, float __y) { return __ocml_add_rtn_f32(__x, __y); }

192 __DEVICE__

193 float __fadd_rn(float __x, float __y) { return __ocml_add_rte_f32(__x, __y); }

195 __DEVICE__

196 float __fadd_ru(float __x, float __y) { return __ocml_add_rtp_f32(__x, __y); }

198 __DEVICE__

199 float __fadd_rz(float __x, float __y) { return __ocml_add_rtz_f32(__x, __y); }

200 #else

202 __DEVICE__

203 float __fadd_rn(float __x, float __y) { return __x + __y; }

204 #endif

205 

206 #if defined OCML_BASIC_ROUNDED_OPERATIONS

208 __DEVICE__

209 float __fdiv_rd(float __x, float __y) { return __ocml_div_rtn_f32(__x, __y); }

211 __DEVICE__

212 float __fdiv_rn(float __x, float __y) { return __ocml_div_rte_f32(__x, __y); }

214 __DEVICE__

215 float __fdiv_ru(float __x, float __y) { return __ocml_div_rtp_f32(__x, __y); }

217 __DEVICE__

218 float __fdiv_rz(float __x, float __y) { return __ocml_div_rtz_f32(__x, __y); }

219 #else

221 __DEVICE__

222 float __fdiv_rn(float __x, float __y) { return __x / __y; }

223 #endif

225 __DEVICE__

226 float __fdividef(float __x, float __y) { return __x / __y; }

227 

228 #if defined OCML_BASIC_ROUNDED_OPERATIONS

230 __DEVICE__

231 float __fmaf_rd(float __x, float __y, float __z) {

232  return __ocml_fma_rtn_f32(__x, __y, __z);

233 }

235 __DEVICE__

236 float __fmaf_rn(float __x, float __y, float __z) {

237  return __ocml_fma_rte_f32(__x, __y, __z);

238 }

240 __DEVICE__

241 float __fmaf_ru(float __x, float __y, float __z) {

242  return __ocml_fma_rtp_f32(__x, __y, __z);

243 }

245 __DEVICE__

246 float __fmaf_rz(float __x, float __y, float __z) {

247  return __ocml_fma_rtz_f32(__x, __y, __z);

248 }

249 #else

251 __DEVICE__

252 float __fmaf_rn(float __x, float __y, float __z) {

253  return __builtin_fmaf(__x, __y, __z);

254 }

255 #endif

256 

257 #if defined OCML_BASIC_ROUNDED_OPERATIONS

259 __DEVICE__

260 float __fmul_rd(float __x, float __y) { return __ocml_mul_rtn_f32(__x, __y); }

262 __DEVICE__

263 float __fmul_rn(float __x, float __y) { return __ocml_mul_rte_f32(__x, __y); }

265 __DEVICE__

266 float __fmul_ru(float __x, float __y) { return __ocml_mul_rtp_f32(__x, __y); }

268 __DEVICE__

269 float __fmul_rz(float __x, float __y) { return __ocml_mul_rtz_f32(__x, __y); }

270 #else

272 __DEVICE__

273 float __fmul_rn(float __x, float __y) { return __x * __y; }

274 #endif

275 

276 #if defined OCML_BASIC_ROUNDED_OPERATIONS

278 __DEVICE__

279 float __frcp_rd(float __x) { return __ocml_div_rtn_f32(1.0f, __x); }

281 __DEVICE__

282 float __frcp_rn(float __x) { return __ocml_div_rte_f32(1.0f, __x); }

284 __DEVICE__

285 float __frcp_ru(float __x) { return __ocml_div_rtp_f32(1.0f, __x); }

287 __DEVICE__

288 float __frcp_rz(float __x) { return __ocml_div_rtz_f32(1.0f, __x); }

289 #else

291 __DEVICE__

292 float __frcp_rn(float __x) { return 1.0f / __x; }

293 #endif

294 

296 __DEVICE__

297 float __frsqrt_rn(float __x) { return __builtin_amdgcn_rsqf(__x); }

298 

299 #if defined OCML_BASIC_ROUNDED_OPERATIONS

301 __DEVICE__

302 float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); }

304 __DEVICE__

305 float __fsqrt_rn(float __x) { return __ocml_sqrt_rte_f32(__x); }

307 __DEVICE__

308 float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); }

310 __DEVICE__

311 float __fsqrt_rz(float __x) { return __ocml_sqrt_rtz_f32(__x); }

312 #else

314 __DEVICE__

315 float __fsqrt_rn(float __x) { return __ocml_native_sqrt_f32(__x); }

316 #endif

317 

318 #if defined OCML_BASIC_ROUNDED_OPERATIONS

320 __DEVICE__

321 float __fsub_rd(float __x, float __y) { return __ocml_sub_rtn_f32(__x, __y); }

323 __DEVICE__

324 float __fsub_rn(float __x, float __y) { return __ocml_sub_rte_f32(__x, __y); }

326 __DEVICE__

327 float __fsub_ru(float __x, float __y) { return __ocml_sub_rtp_f32(__x, __y); }

329 __DEVICE__

330 float __fsub_rz(float __x, float __y) { return __ocml_sub_rtz_f32(__x, __y); }

331 #else

333 __DEVICE__

334 float __fsub_rn(float __x, float __y) { return __x - __y; }

335 #endif

336 

338 __DEVICE__

339 float __log10f(float __x) { return __builtin_log10f(__x); }

341 __DEVICE__

342 float __log2f(float __x) { return __builtin_amdgcn_logf(__x); }

344 __DEVICE__

345 float __logf(float __x) { return __builtin_logf(__x); }

347 __DEVICE__

348 float __powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }

350 __DEVICE__

351 float __saturatef(float __x) { return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x); }

353 __DEVICE__

354 void __sincosf(float __x, float *__sinptr, float *__cosptr) {

355  *__sinptr = __ocml_native_sin_f32(__x);

356  *__cosptr = __ocml_native_cos_f32(__x);

357 }

359 __DEVICE__

360 float __sinf(float __x) { return __ocml_native_sin_f32(__x); }

362 __DEVICE__

363 float __tanf(float __x) {

364  return __sinf(__x) * __builtin_amdgcn_rcpf(__cosf(__x));

365 }

366 // END INTRINSICS

367 

369 

370 #if defined(__cplusplus)

372 __DEVICE__

373 int abs(int __x) {

374  return __builtin_abs(__x);

375 }

377 __DEVICE__

378 long labs(long __x) {

379  return __builtin_labs(__x);

380 }

382 __DEVICE__

383 long long llabs(long long __x) {

384  return __builtin_llabs(__x);

385 }

386 #endif

387 

391 

393 __DEVICE__

394 float acosf(float __x) { return __ocml_acos_f32(__x); }

396 __DEVICE__

397 float acoshf(float __x) { return __ocml_acosh_f32(__x); }

399 __DEVICE__

400 float asinf(float __x) { return __ocml_asin_f32(__x); }

402 __DEVICE__

403 float asinhf(float __x) { return __ocml_asinh_f32(__x); }

405 __DEVICE__

406 float atan2f(float __x, float __y) { return __ocml_atan2_f32(__x, __y); }

408 __DEVICE__

409 float atanf(float __x) { return __ocml_atan_f32(__x); }

411 __DEVICE__

412 float atanhf(float __x) { return __ocml_atanh_f32(__x); }

414 __DEVICE__

415 float cbrtf(float __x) { return __ocml_cbrt_f32(__x); }

417 __DEVICE__

418 float ceilf(float __x) { return __builtin_ceilf(__x); }

420 __DEVICE__

421 float copysignf(float __x, float __y) { return __builtin_copysignf(__x, __y); }

423 __DEVICE__

424 float cosf(float __x) { return __FAST_OR_SLOW(__cosf, __ocml_cos_f32)(__x); }

426 __DEVICE__

427 float coshf(float __x) { return __ocml_cosh_f32(__x); }

429 __DEVICE__

430 float cospif(float __x) { return __ocml_cospi_f32(__x); }

432 __DEVICE__

433 float cyl_bessel_i0f(float __x) { return __ocml_i0_f32(__x); }

435 __DEVICE__

436 float cyl_bessel_i1f(float __x) { return __ocml_i1_f32(__x); }

438 __DEVICE__

439 float erfcf(float __x) { return __ocml_erfc_f32(__x); }

441 __DEVICE__

442 float erfcinvf(float __x) { return __ocml_erfcinv_f32(__x); }

444 __DEVICE__

445 float erfcxf(float __x) { return __ocml_erfcx_f32(__x); }

447 __DEVICE__

448 float erff(float __x) { return __ocml_erf_f32(__x); }

450 __DEVICE__

451 float erfinvf(float __x) { return __ocml_erfinv_f32(__x); }

453 __DEVICE__

454 float exp10f(float __x) { return __ocml_exp10_f32(__x); }

456 __DEVICE__

457 float exp2f(float __x) { return __builtin_exp2f(__x); }

459 __DEVICE__

460 float expf(float __x) { return __builtin_expf(__x); }

462 __DEVICE__

463 float expm1f(float __x) { return __ocml_expm1_f32(__x); }

465 __DEVICE__

466 float fabsf(float __x) { return __builtin_fabsf(__x); }

468 __DEVICE__

469 float fdimf(float __x, float __y) { return __ocml_fdim_f32(__x, __y); }

471 __DEVICE__

472 float fdividef(float __x, float __y) { return __x / __y; }

474 __DEVICE__

475 float floorf(float __x) { return __builtin_floorf(__x); }

477 __DEVICE__

478 float fmaf(float __x, float __y, float __z) {

479  return __builtin_fmaf(__x, __y, __z);

480 }

482 __DEVICE__

483 float fmaxf(float __x, float __y) { return __builtin_fmaxf(__x, __y); }

485 __DEVICE__

486 float fminf(float __x, float __y) { return __builtin_fminf(__x, __y); }

488 __DEVICE__

489 float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); }

491 __DEVICE__

492 float frexpf(float __x, int *__nptr) {

493  return __builtin_frexpf(__x, __nptr);

494 }

496 __DEVICE__

497 float hypotf(float __x, float __y) { return __ocml_hypot_f32(__x, __y); }

499 __DEVICE__

500 int ilogbf(float __x) { return __ocml_ilogb_f32(__x); }

502 __DEVICE__

503 __RETURN_TYPE __finitef(float __x) { return __builtin_isfinite(__x); }

505 __DEVICE__

506 __RETURN_TYPE __isinff(float __x) { return __builtin_isinf(__x); }

508 __DEVICE__

509 __RETURN_TYPE __isnanf(float __x) { return __builtin_isnan(__x); }

511 __DEVICE__

512 float j0f(float __x) { return __ocml_j0_f32(__x); }

514 __DEVICE__

515 float j1f(float __x) { return __ocml_j1_f32(__x); }

517 __DEVICE__

518 float jnf(int __n, float __x) { // TODO: we could use Ahmes multiplication

519  // and the Miller & Brown algorithm

520  // for linear recurrences to get O(log n) steps, but it's unclear if

521  // it'd be beneficial in this case.

522  if (__n == 0)

523  return j0f(__x);

524  if (__n == 1)

525  return j1f(__x);

526 

527  float __x0 = j0f(__x);

528  float __x1 = j1f(__x);

529  for (int __i = 1; __i < __n; ++__i) {

530  float __x2 = (2 * __i) / __x * __x1 - __x0;

531  __x0 = __x1;

532  __x1 = __x2;

533  }

534 

535  return __x1;

536 }

538 __DEVICE__

539 float ldexpf(float __x, int __e) { return __builtin_amdgcn_ldexpf(__x, __e); }

541 __DEVICE__

542 float lgammaf(float __x) { return __ocml_lgamma_f32(__x); }

544 __DEVICE__

545 long long int llrintf(float __x) { return __builtin_rintf(__x); }

547 __DEVICE__

548 long long int llroundf(float __x) { return __builtin_roundf(__x); }

550 __DEVICE__

551 float log10f(float __x) { return __builtin_log10f(__x); }

553 __DEVICE__

554 float log1pf(float __x) { return __ocml_log1p_f32(__x); }

556 __DEVICE__

557 float log2f(float __x) { return __FAST_OR_SLOW(__log2f, __ocml_log2_f32)(__x); }

559 __DEVICE__

560 float logbf(float __x) { return __ocml_logb_f32(__x); }

562 __DEVICE__

563 float logf(float __x) { return __FAST_OR_SLOW(__logf, __ocml_log_f32)(__x); }

565 __DEVICE__

566 long int lrintf(float __x) { return __builtin_rintf(__x); }

568 __DEVICE__

569 long int lroundf(float __x) { return __builtin_roundf(__x); }

571 __DEVICE__

572 float modff(float __x, float *__iptr) {

573  float __tmp;

574 #ifdef __OPENMP_AMDGCN__

575 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)

576 #endif

577  float __r =

578  __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);

579  *__iptr = __tmp;

580  return __r;

581 }

583 __DEVICE__

584 float nanf(const char *__tagp __attribute__((nonnull))) {

585  union {

586  float val;

587  struct ieee_float {

588  unsigned int mantissa : 22;

589  unsigned int quiet : 1;

590  unsigned int exponent : 8;

591  unsigned int sign : 1;

592  } bits;

593  } __tmp;

594  __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));

595 

596  __tmp.bits.sign = 0u;

597  __tmp.bits.exponent = ~0u;

598  __tmp.bits.quiet = 1u;

599  __tmp.bits.mantissa = __make_mantissa(__tagp);

600 

601  return __tmp.val;

602 }

604 __DEVICE__

605 float nearbyintf(float __x) { return __builtin_nearbyintf(__x); }

607 __DEVICE__

608 float nextafterf(float __x, float __y) {

609  return __ocml_nextafter_f32(__x, __y);

610 }

612 __DEVICE__

613 float norm3df(float __x, float __y, float __z) {

614  return __ocml_len3_f32(__x, __y, __z);

615 }

617 __DEVICE__

618 float norm4df(float __x, float __y, float __z, float __w) {

619  return __ocml_len4_f32(__x, __y, __z, __w);

620 }

622 __DEVICE__

623 float normcdff(float __x) { return __ocml_ncdf_f32(__x); }

625 __DEVICE__

626 float normcdfinvf(float __x) { return __ocml_ncdfinv_f32(__x); }

628 __DEVICE__

629 float normf(int __dim,

630  const float *__a) { // TODO: placeholder until OCML adds support.

631  float __r = 0;

632  while (__dim--) {

633  __r += __a[0] * __a[0];

634  ++__a;

635  }

636 

637  return __builtin_sqrtf(__r);

638 }

640 __DEVICE__

641 float powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }

643 __DEVICE__

644 float powif(float __x, int __y) { return __ocml_pown_f32(__x, __y); }

646 __DEVICE__

647 float rcbrtf(float __x) { return __ocml_rcbrt_f32(__x); }

649 __DEVICE__

650 float remainderf(float __x, float __y) {

651  return __ocml_remainder_f32(__x, __y);

652 }

654 __DEVICE__

655 float remquof(float __x, float __y, int *__quo) {

656  int __tmp;

657 #ifdef __OPENMP_AMDGCN__

658 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)

659 #endif

660  float __r = __ocml_remquo_f32(

661  __x, __y, (__attribute__((address_space(5))) int *)&__tmp);

662  *__quo = __tmp;

663 

664  return __r;

665 }

667 __DEVICE__

668 float rhypotf(float __x, float __y) { return __ocml_rhypot_f32(__x, __y); }

670 __DEVICE__

671 float rintf(float __x) { return __builtin_rintf(__x); }

673 __DEVICE__

674 float rnorm3df(float __x, float __y, float __z) {

675  return __ocml_rlen3_f32(__x, __y, __z);

676 }

678 __DEVICE__

679 float rnorm4df(float __x, float __y, float __z, float __w) {

680  return __ocml_rlen4_f32(__x, __y, __z, __w);

681 }

683 __DEVICE__

684 float rnormf(int __dim,

685  const float *__a) { // TODO: placeholder until OCML adds support.

686  float __r = 0;

687  while (__dim--) {

688  __r += __a[0] * __a[0];

689  ++__a;

690  }

691 

692  return __ocml_rsqrt_f32(__r);

693 }

695 __DEVICE__

696 float roundf(float __x) { return __builtin_roundf(__x); }

698 __DEVICE__

699 float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); }

701 __DEVICE__

702 float scalblnf(float __x, long int __n) {

703  return (__n < INT_MAX) ? __builtin_amdgcn_ldexpf(__x, __n)

704  : __ocml_scalb_f32(__x, __n);

705 }

707 __DEVICE__

708 float scalbnf(float __x, int __n) { return __builtin_amdgcn_ldexpf(__x, __n); }

710 __DEVICE__

711 __RETURN_TYPE __signbitf(float __x) { return __builtin_signbitf(__x); }

713 __DEVICE__

714 void sincosf(float __x, float *__sinptr, float *__cosptr) {

715  float __tmp;

716 #ifdef __OPENMP_AMDGCN__

717 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)

718 #endif

719 #ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__

720  __sincosf(__x, __sinptr, __cosptr);

721 #else

722  *__sinptr =

723  __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);

724  *__cosptr = __tmp;

725 #endif

726 }

728 __DEVICE__

729 void sincospif(float __x, float *__sinptr, float *__cosptr) {

730  float __tmp;

731 #ifdef __OPENMP_AMDGCN__

732 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)

733 #endif

734  *__sinptr = __ocml_sincospi_f32(

735  __x, (__attribute__((address_space(5))) float *)&__tmp);

736  *__cosptr = __tmp;

737 }

739 __DEVICE__

740 float sinf(float __x) { return __FAST_OR_SLOW(__sinf, __ocml_sin_f32)(__x); }

742 __DEVICE__

743 float sinhf(float __x) { return __ocml_sinh_f32(__x); }

745 __DEVICE__

746 float sinpif(float __x) { return __ocml_sinpi_f32(__x); }

748 __DEVICE__

749 float sqrtf(float __x) { return __builtin_sqrtf(__x); }

751 __DEVICE__

752 float tanf(float __x) { return __ocml_tan_f32(__x); }

754 __DEVICE__

755 float tanhf(float __x) { return __ocml_tanh_f32(__x); }

757 __DEVICE__

758 float tgammaf(float __x) { return __ocml_tgamma_f32(__x); }

760 __DEVICE__

761 float truncf(float __x) { return __builtin_truncf(__x); }

763 __DEVICE__

764 float y0f(float __x) { return __ocml_y0_f32(__x); }

766 __DEVICE__

767 float y1f(float __x) { return __ocml_y1_f32(__x); }

769 __DEVICE__

770 float ynf(int __n, float __x) { // TODO: we could use Ahmes multiplication

771  // and the Miller & Brown algorithm

772  // for linear recurrences to get O(log n) steps, but it's unclear if

773  // it'd be beneficial in this case. Placeholder until OCML adds

774  // support.

775  if (__n == 0)

776  return y0f(__x);

777  if (__n == 1)

778  return y1f(__x);

779 

780  float __x0 = y0f(__x);

781  float __x1 = y1f(__x);

782  for (int __i = 1; __i < __n; ++__i) {

783  float __x2 = (2 * __i) / __x * __x1 - __x0;

784  __x0 = __x1;

785  __x1 = __x2;

786  }

787 

788  return __x1;

789 }

790 // END FLOAT

791 

793 

797 

798 // BEGIN DOUBLE

799 

801 __DEVICE__

802 double acos(double __x) { return __ocml_acos_f64(__x); }

804 __DEVICE__

805 double acosh(double __x) { return __ocml_acosh_f64(__x); }

807 __DEVICE__

808 double asin(double __x) { return __ocml_asin_f64(__x); }

810 __DEVICE__

811 double asinh(double __x) { return __ocml_asinh_f64(__x); }

813 __DEVICE__

814 double atan(double __x) { return __ocml_atan_f64(__x); }

816 __DEVICE__

817 double atan2(double __x, double __y) { return __ocml_atan2_f64(__x, __y); }

819 __DEVICE__

820 double atanh(double __x) { return __ocml_atanh_f64(__x); }

822 __DEVICE__

823 double cbrt(double __x) { return __ocml_cbrt_f64(__x); }

825 __DEVICE__

826 double ceil(double __x) { return __builtin_ceil(__x); }

828 __DEVICE__

829 double copysign(double __x, double __y) {

830  return __builtin_copysign(__x, __y);

831 }

833 __DEVICE__

834 double cos(double __x) { return __ocml_cos_f64(__x); }

836 __DEVICE__

837 double cosh(double __x) { return __ocml_cosh_f64(__x); }

839 __DEVICE__

840 double cospi(double __x) { return __ocml_cospi_f64(__x); }

842 __DEVICE__

843 double cyl_bessel_i0(double __x) { return __ocml_i0_f64(__x); }

845 __DEVICE__

846 double cyl_bessel_i1(double __x) { return __ocml_i1_f64(__x); }

848 __DEVICE__

849 double erf(double __x) { return __ocml_erf_f64(__x); }

851 __DEVICE__

852 double erfc(double __x) { return __ocml_erfc_f64(__x); }

854 __DEVICE__

855 double erfcinv(double __x) { return __ocml_erfcinv_f64(__x); }

857 __DEVICE__

858 double erfcx(double __x) { return __ocml_erfcx_f64(__x); }

860 __DEVICE__

861 double erfinv(double __x) { return __ocml_erfinv_f64(__x); }

863 __DEVICE__

864 double exp(double __x) { return __ocml_exp_f64(__x); }

866 __DEVICE__

867 double exp10(double __x) { return __ocml_exp10_f64(__x); }

869 __DEVICE__

870 double exp2(double __x) { return __ocml_exp2_f64(__x); }

872 __DEVICE__

873 double expm1(double __x) { return __ocml_expm1_f64(__x); }

875 __DEVICE__

876 double fabs(double __x) { return __builtin_fabs(__x); }

878 __DEVICE__

879 double fdim(double __x, double __y) { return __ocml_fdim_f64(__x, __y); }

881 __DEVICE__

882 double floor(double __x) { return __builtin_floor(__x); }

884 __DEVICE__

885 double fma(double __x, double __y, double __z) {

886  return __builtin_fma(__x, __y, __z);

887 }

889 __DEVICE__

890 double fmax(double __x, double __y) { return __builtin_fmax(__x, __y); }

892 __DEVICE__

893 double fmin(double __x, double __y) { return __builtin_fmin(__x, __y); }

895 __DEVICE__

896 double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); }

898 __DEVICE__

899 double frexp(double __x, int *__nptr) {

900  return __builtin_frexp(__x, __nptr);

901 }

903 __DEVICE__

904 double hypot(double __x, double __y) { return __ocml_hypot_f64(__x, __y); }

906 __DEVICE__

907 int ilogb(double __x) { return __ocml_ilogb_f64(__x); }

909 __DEVICE__

910 __RETURN_TYPE __finite(double __x) { return __builtin_isfinite(__x); }

912 __DEVICE__

913 __RETURN_TYPE __isinf(double __x) { return __builtin_isinf(__x); }

915 __DEVICE__

916 __RETURN_TYPE __isnan(double __x) { return __builtin_isnan(__x); }

918 __DEVICE__

919 double j0(double __x) { return __ocml_j0_f64(__x); }

921 __DEVICE__

922 double j1(double __x) { return __ocml_j1_f64(__x); }

924 __DEVICE__

925 double jn(int __n, double __x) { // TODO: we could use Ahmes multiplication

926  // and the Miller & Brown algorithm

927  // for linear recurrences to get O(log n) steps, but it's unclear if

928  // it'd be beneficial in this case. Placeholder until OCML adds

929  // support.

930  if (__n == 0)

931  return j0(__x);

932  if (__n == 1)

933  return j1(__x);

934 

935  double __x0 = j0(__x);

936  double __x1 = j1(__x);

937  for (int __i = 1; __i < __n; ++__i) {

938  double __x2 = (2 * __i) / __x * __x1 - __x0;

939  __x0 = __x1;

940  __x1 = __x2;

941  }

942  return __x1;

943 }

945 __DEVICE__

946 double ldexp(double __x, int __e) { return __builtin_amdgcn_ldexp(__x, __e); }

948 __DEVICE__

949 double lgamma(double __x) { return __ocml_lgamma_f64(__x); }

951 __DEVICE__

952 long long int llrint(double __x) { return __builtin_rint(__x); }

954 __DEVICE__

955 long long int llround(double __x) { return __builtin_round(__x); }

957 __DEVICE__

958 double log(double __x) { return __ocml_log_f64(__x); }

960 __DEVICE__

961 double log10(double __x) { return __ocml_log10_f64(__x); }

963 __DEVICE__

964 double log1p(double __x) { return __ocml_log1p_f64(__x); }

966 __DEVICE__

967 double log2(double __x) { return __ocml_log2_f64(__x); }

969 __DEVICE__

970 double logb(double __x) { return __ocml_logb_f64(__x); }

972 __DEVICE__

973 long int lrint(double __x) { return __builtin_rint(__x); }

975 __DEVICE__

976 long int lround(double __x) { return __builtin_round(__x); }

978 __DEVICE__

979 double modf(double __x, double *__iptr) {

980  double __tmp;

981 #ifdef __OPENMP_AMDGCN__

982 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)

983 #endif

984  double __r =

985  __ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);

986  *__iptr = __tmp;

987 

988  return __r;

989 }

991 __DEVICE__

992 double nan(const char *__tagp) {

993 #if !_WIN32

994  union {

995  double val;

996  struct ieee_double {

997  uint64_t mantissa : 51;

998  uint32_t quiet : 1;

999  uint32_t exponent : 11;

1000  uint32_t sign : 1;

1001  } bits;

1002  } __tmp;

1003  __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));

1004 

1005  __tmp.bits.sign = 0u;

1006  __tmp.bits.exponent = ~0u;

1007  __tmp.bits.quiet = 1u;

1008  __tmp.bits.mantissa = __make_mantissa(__tagp);

1009 

1010  return __tmp.val;

1011 #else

1012  __static_assert_type_size_equal(sizeof(uint64_t), sizeof(double));

1013  uint64_t __val = __make_mantissa(__tagp);

1014  __val |= 0xFFF << 51;

1015  return *reinterpret_cast<double *>(&__val);

1016 #endif

1017 }

1019 __DEVICE__

1020 double nearbyint(double __x) { return __builtin_nearbyint(__x); }

1022 __DEVICE__

1023 double nextafter(double __x, double __y) {

1024  return __ocml_nextafter_f64(__x, __y);

1025 }

1027 __DEVICE__

1028 double norm(int __dim,

1029  const double *__a) { // TODO: placeholder until OCML adds support.

1030  double __r = 0;

1031  while (__dim--) {

1032  __r += __a[0] * __a[0];

1033  ++__a;

1034  }

1035 

1036  return __builtin_sqrt(__r);

1037 }

1039 __DEVICE__

1040 double norm3d(double __x, double __y, double __z) {

1041  return __ocml_len3_f64(__x, __y, __z);

1042 }

1044 __DEVICE__

1045 double norm4d(double __x, double __y, double __z, double __w) {

1046  return __ocml_len4_f64(__x, __y, __z, __w);

1047 }

1049 __DEVICE__

1050 double normcdf(double __x) { return __ocml_ncdf_f64(__x); }

1052 __DEVICE__

1053 double normcdfinv(double __x) { return __ocml_ncdfinv_f64(__x); }

1055 __DEVICE__

1056 double pow(double __x, double __y) { return __ocml_pow_f64(__x, __y); }

1058 __DEVICE__

1059 double powi(double __x, int __y) { return __ocml_pown_f64(__x, __y); }

1061 __DEVICE__

1062 double rcbrt(double __x) { return __ocml_rcbrt_f64(__x); }

1064 __DEVICE__

1065 double remainder(double __x, double __y) {

1066  return __ocml_remainder_f64(__x, __y);

1067 }

1069 __DEVICE__

1070 double remquo(double __x, double __y, int *__quo) {

1071  int __tmp;

1072 #ifdef __OPENMP_AMDGCN__

1073 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)

1074 #endif

1075  double __r = __ocml_remquo_f64(

1076  __x, __y, (__attribute__((address_space(5))) int *)&__tmp);

1077  *__quo = __tmp;

1078 

1079  return __r;

1080 }

1082 __DEVICE__

1083 double rhypot(double __x, double __y) { return __ocml_rhypot_f64(__x, __y); }

1085 __DEVICE__

1086 double rint(double __x) { return __builtin_rint(__x); }

1088 __DEVICE__

1089 double rnorm(int __dim,

1090  const double *__a) { // TODO: placeholder until OCML adds support.

1091  double __r = 0;

1092  while (__dim--) {

1093  __r += __a[0] * __a[0];

1094  ++__a;

1095  }

1096 

1097  return __ocml_rsqrt_f64(__r);

1098 }

1100 __DEVICE__

1101 double rnorm3d(double __x, double __y, double __z) {

1102  return __ocml_rlen3_f64(__x, __y, __z);

1103 }

1105 __DEVICE__

1106 double rnorm4d(double __x, double __y, double __z, double __w) {

1107  return __ocml_rlen4_f64(__x, __y, __z, __w);

1108 }

1110 __DEVICE__

1111 double round(double __x) { return __builtin_round(__x); }

1113 __DEVICE__

1114 double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); }

1116 __DEVICE__

1117 double scalbln(double __x, long int __n) {

1118  return (__n < INT_MAX) ? __builtin_amdgcn_ldexp(__x, __n)

1119  : __ocml_scalb_f64(__x, __n);

1120 }

1122 __DEVICE__

1123 double scalbn(double __x, int __n) { return __builtin_amdgcn_ldexp(__x, __n); }

1125 __DEVICE__

1126 __RETURN_TYPE __signbit(double __x) { return __builtin_signbit(__x); }

1128 __DEVICE__

1129 double sin(double __x) { return __ocml_sin_f64(__x); }

1131 __DEVICE__

1132 void sincos(double __x, double *__sinptr, double *__cosptr) {

1133  double __tmp;

1134 #ifdef __OPENMP_AMDGCN__

1135 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)

1136 #endif

1137  *__sinptr = __ocml_sincos_f64(

1138  __x, (__attribute__((address_space(5))) double *)&__tmp);

1139  *__cosptr = __tmp;

1140 }

1142 __DEVICE__

1143 void sincospi(double __x, double *__sinptr, double *__cosptr) {

1144  double __tmp;

1145 #ifdef __OPENMP_AMDGCN__

1146 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)

1147 #endif

1148  *__sinptr = __ocml_sincospi_f64(

1149  __x, (__attribute__((address_space(5))) double *)&__tmp);

1150  *__cosptr = __tmp;

1151 }

1153 __DEVICE__

1154 double sinh(double __x) { return __ocml_sinh_f64(__x); }

1156 __DEVICE__

1157 double sinpi(double __x) { return __ocml_sinpi_f64(__x); }

1159 __DEVICE__

1160 double sqrt(double __x) { return __builtin_sqrt(__x); }

1162 __DEVICE__

1163 double tan(double __x) { return __ocml_tan_f64(__x); }

1165 __DEVICE__

1166 double tanh(double __x) { return __ocml_tanh_f64(__x); }

1168 __DEVICE__

1169 double tgamma(double __x) { return __ocml_tgamma_f64(__x); }

1171 __DEVICE__

1172 double trunc(double __x) { return __builtin_trunc(__x); }

1174 __DEVICE__

1175 double y0(double __x) { return __ocml_y0_f64(__x); }

1177 __DEVICE__

1178 double y1(double __x) { return __ocml_y1_f64(__x); }

1180 __DEVICE__

1181 double yn(int __n, double __x) { // TODO: we could use Ahmes multiplication

1182  // and the Miller & Brown algorithm

1183  // for linear recurrences to get O(log n) steps, but it's unclear if

1184  // it'd be beneficial in this case. Placeholder until OCML adds

1185  // support.

1186  if (__n == 0)

1187  return y0(__x);

1188  if (__n == 1)

1189  return y1(__x);

1190 

1191  double __x0 = y0(__x);

1192  double __x1 = y1(__x);

1193  for (int __i = 1; __i < __n; ++__i) {

1194  double __x2 = (2 * __i) / __x * __x1 - __x0;

1195  __x0 = __x1;

1196  __x1 = __x2;

1197  }

1198 

1199  return __x1;

1200 }

1202 

1206 

1207 // BEGIN INTRINSICS

1208 

1209 #if defined OCML_BASIC_ROUNDED_OPERATIONS

1211 __DEVICE__

1212 double __dadd_rd(double __x, double __y) {

1213  return __ocml_add_rtn_f64(__x, __y);

1214 }

1216 __DEVICE__

1217 double __dadd_rn(double __x, double __y) {

1218  return __ocml_add_rte_f64(__x, __y);

1219 }

1221 __DEVICE__

1222 double __dadd_ru(double __x, double __y) {

1223  return __ocml_add_rtp_f64(__x, __y);

1224 }

1226 __DEVICE__

1227 double __dadd_rz(double __x, double __y) {

1228  return __ocml_add_rtz_f64(__x, __y);

1229 }

1230 #else

1232 __DEVICE__

1233 double __dadd_rn(double __x, double __y) { return __x + __y; }

1234 #endif

1235 

1236 #if defined OCML_BASIC_ROUNDED_OPERATIONS

1238 __DEVICE__

1239 double __ddiv_rd(double __x, double __y) {

1240  return __ocml_div_rtn_f64(__x, __y);

1241 }

1243 __DEVICE__

1244 double __ddiv_rn(double __x, double __y) {

1245  return __ocml_div_rte_f64(__x, __y);

1246 }

1248 __DEVICE__

1249 double __ddiv_ru(double __x, double __y) {

1250  return __ocml_div_rtp_f64(__x, __y);

1251 }

1253 __DEVICE__

1254 double __ddiv_rz(double __x, double __y) {

1255  return __ocml_div_rtz_f64(__x, __y);

1256 }

1257 #else

1259 __DEVICE__

1260 double __ddiv_rn(double __x, double __y) { return __x / __y; }

1261 #endif

1262 

1263 #if defined OCML_BASIC_ROUNDED_OPERATIONS

1265 __DEVICE__

1266 double __dmul_rd(double __x, double __y) {

1267  return __ocml_mul_rtn_f64(__x, __y);

1268 }

1270 __DEVICE__

1271 double __dmul_rn(double __x, double __y) {

1272  return __ocml_mul_rte_f64(__x, __y);

1273 }

1275 __DEVICE__

1276 double __dmul_ru(double __x, double __y) {

1277  return __ocml_mul_rtp_f64(__x, __y);

1278 }

1280 __DEVICE__

1281 double __dmul_rz(double __x, double __y) {

1282  return __ocml_mul_rtz_f64(__x, __y);

1283 }

1284 #else

1286 __DEVICE__

1287 double __dmul_rn(double __x, double __y) { return __x * __y; }

1288 #endif

1289 

1290 #if defined OCML_BASIC_ROUNDED_OPERATIONS

1292 __DEVICE__

1293 double __drcp_rd(double __x) { return __ocml_div_rtn_f64(1.0, __x); }

1295 __DEVICE__

1296 double __drcp_rn(double __x) { return __ocml_div_rte_f64(1.0, __x); }

1298 __DEVICE__

1299 double __drcp_ru(double __x) { return __ocml_div_rtp_f64(1.0, __x); }

1301 __DEVICE__

1302 double __drcp_rz(double __x) { return __ocml_div_rtz_f64(1.0, __x); }

1303 #else

1305 __DEVICE__

1306 double __drcp_rn(double __x) { return 1.0 / __x; }

1307 #endif

1308 

1309 #if defined OCML_BASIC_ROUNDED_OPERATIONS

1311 __DEVICE__

1312 double __dsqrt_rd(double __x) { return __ocml_sqrt_rtn_f64(__x); }

1314 __DEVICE__

1315 double __dsqrt_rn(double __x) { return __ocml_sqrt_rte_f64(__x); }

1317 __DEVICE__

1318 double __dsqrt_ru(double __x) { return __ocml_sqrt_rtp_f64(__x); }

1320 __DEVICE__

1321 double __dsqrt_rz(double __x) { return __ocml_sqrt_rtz_f64(__x); }

1322 #else

1324 __DEVICE__

1325 double __dsqrt_rn(double __x) { return __builtin_sqrt(__x); }

1326 #endif

1327 

1328 #if defined OCML_BASIC_ROUNDED_OPERATIONS

1330 __DEVICE__

1331 double __dsub_rd(double __x, double __y) {

1332  return __ocml_sub_rtn_f64(__x, __y);

1333 }

1335 __DEVICE__

1336 double __dsub_rn(double __x, double __y) {

1337  return __ocml_sub_rte_f64(__x, __y);

1338 }

1340 __DEVICE__

1341 double __dsub_ru(double __x, double __y) {

1342  return __ocml_sub_rtp_f64(__x, __y);

1343 }

1345 __DEVICE__

1346 double __dsub_rz(double __x, double __y) {

1347  return __ocml_sub_rtz_f64(__x, __y);

1348 }

1349 #else

1351 __DEVICE__

1352 double __dsub_rn(double __x, double __y) { return __x - __y; }

1353 #endif

1354 

1355 #if defined OCML_BASIC_ROUNDED_OPERATIONS

1357 __DEVICE__

1358 double __fma_rd(double __x, double __y, double __z) {

1359  return __ocml_fma_rtn_f64(__x, __y, __z);

1360 }

1362 __DEVICE__

1363 double __fma_rn(double __x, double __y, double __z) {

1364  return __ocml_fma_rte_f64(__x, __y, __z);

1365 }

1367 __DEVICE__

1368 double __fma_ru(double __x, double __y, double __z) {

1369  return __ocml_fma_rtp_f64(__x, __y, __z);

1370 }

1372 __DEVICE__

1373 double __fma_rz(double __x, double __y, double __z) {

1374  return __ocml_fma_rtz_f64(__x, __y, __z);

1375 }

1376 #else

1378 __DEVICE__

1379 double __fma_rn(double __x, double __y, double __z) {

1380  return __builtin_fma(__x, __y, __z);

1381 }

1382 #endif

1383 // END INTRINSICS

1384 

1385 // END DOUBLE

1386 

1388 

1389 // C only macros

1390 #if !defined(__cplusplus) && __STDC_VERSION__ >= 201112L

1391 #define isfinite(__x) _Generic((__x), float : __finitef, double : __finite)(__x)

1392 #define isinf(__x) _Generic((__x), float : __isinff, double : __isinf)(__x)

1393 #define isnan(__x) _Generic((__x), float : __isnanf, double : __isnan)(__x)

1394 #define signbit(__x) \

1395  _Generic((__x), float : __signbitf, double : __signbit)(__x)

1396 #endif // !defined(__cplusplus) && __STDC_VERSION__ >= 201112L

1397 

1398 #if defined(__cplusplus)

1401 

1403 template <class T> __DEVICE__ T min(T __arg1, T __arg2) {

1404  return (__arg1 < __arg2) ? __arg1 : __arg2;

1405 }

1407 template <class T> __DEVICE__ T max(T __arg1, T __arg2) {

1408  return (__arg1 > __arg2) ? __arg1 : __arg2;

1409 }

1411 

1414 

1416 __DEVICE__ int min(int __arg1, int __arg2) {

1417  return (__arg1 < __arg2) ? __arg1 : __arg2;

1418 }

1420 __DEVICE__ int max(int __arg1, int __arg2) {

1421  return (__arg1 > __arg2) ? __arg1 : __arg2;

1422 }

1424 

1427 __DEVICE__

1428 float max(float __x, float __y) { return __builtin_fmaxf(__x, __y); }

1431 __DEVICE__

1432 double max(double __x, double __y) { return __builtin_fmax(__x, __y); }

1435 __DEVICE__

1436 float min(float __x, float __y) { return __builtin_fminf(__x, __y); }

1439 __DEVICE__

1440 double min(double __x, double __y) { return __builtin_fmin(__x, __y); }

1441 

1442 #if !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)

1445 

1447 __host__ inline static int min(int __arg1, int __arg2) {

1448  return __arg1 < __arg2 ? __arg1 : __arg2;

1449 }

1451 __host__ inline static int max(int __arg1, int __arg2) {

1452  return __arg1 > __arg2 ? __arg1 : __arg2;

1453 }

1455 #endif // !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)

1456 #endif

1457 

1458 // doxygen end Math API

1460 // doxygen end Device APIs

1462 

1463 #pragma pop_macro("__DEVICE__")

1464 #pragma pop_macro("__RETURN_TYPE")

1465 #pragma pop_macro("__FAST_OR_SLOW")

1466 

1467 #endif // __CLANG_HIP_MATH_H__

__DEVICE__

#define __DEVICE__

Definition: __clang_hip_math.h:29

__FAST_OR_SLOW

#define __FAST_OR_SLOW(fast, slow)

Definition: __clang_hip_math.h:40

__RETURN_TYPE

#define __RETURN_TYPE

Definition: __clang_hip_math.h:51

__static_assert_type_size_equal

#define __static_assert_type_size_equal(A, B)

Definition: __clang_hip_math.h:75

log2

__DEVICE__ double log2(double __x)

Returns the base 2 logarithm of x.

Definition: __clang_hip_math.h:967

ceil

__DEVICE__ double ceil(double __x)

Returns ceiling of x.

Definition: __clang_hip_math.h:826

atan2

__DEVICE__ double atan2(double __x, double __y)

Returns the arc tangent of the ratio of x and y.

Definition: __clang_hip_math.h:817

rhypot

__DEVICE__ double rhypot(double __x, double __y)

Returns one over the square root of the sum of squares of x and y.

Definition: __clang_hip_math.h:1083

normcdfinv

__DEVICE__ double normcdfinv(double __x)

Returns the inverse of the standard normal cumulative distribution function.

Definition: __clang_hip_math.h:1053

norm3d

__DEVICE__ double norm3d(double __x, double __y, double __z)

Returns the square root of the sum of squares of x, y and z.

Definition: __clang_hip_math.h:1040

log

__DEVICE__ double log(double __x)

Returns the natural logarithm of x.

Definition: __clang_hip_math.h:958

cbrt

__DEVICE__ double cbrt(double __x)

Returns the cube root of x.

Definition: __clang_hip_math.h:823

ldexp

__DEVICE__ double ldexp(double __x, int __e)

Returns the value of for x and e.

Definition: __clang_hip_math.h:946

fmax

__DEVICE__ double fmax(double __x, double __y)

Determine the maximum numeric value of x and y.

Definition: __clang_hip_math.h:890

cosh

__DEVICE__ double cosh(double __x)

Returns the hyperbolic cosine of x.

Definition: __clang_hip_math.h:837

tgamma

__DEVICE__ double tgamma(double __x)

Returns the gamma function of x.

Definition: __clang_hip_math.h:1169

sinh

__DEVICE__ double sinh(double __x)

Returns the hyperbolic sine of x.

Definition: __clang_hip_math.h:1154

erfinv

__DEVICE__ double erfinv(double __x)

Returns the inverse error function of x.

Definition: __clang_hip_math.h:861

fdim

__DEVICE__ double fdim(double __x, double __y)

Returns the positive difference between x and y.

Definition: __clang_hip_math.h:879

lrint

__DEVICE__ long int lrint(double __x)

Round x to nearest integer value.

Definition: __clang_hip_math.h:973

j0

__DEVICE__ double j0(double __x)

Returns the value of the Bessel function of the first kind of order 0 for x.

Definition: __clang_hip_math.h:919

exp

__DEVICE__ double exp(double __x)

Returns .

Definition: __clang_hip_math.h:864

j1

__DEVICE__ double j1(double __x)

Returns the value of the Bessel function of the first kind of order 1 for x.

Definition: __clang_hip_math.h:922

__isnan

__DEVICE__ __RETURN_TYPE __isnan(double __x)

Determine whether x is a NaN.

Definition: __clang_hip_math.h:916

cos

__DEVICE__ double cos(double __x)

Returns the cosine of x.

Definition: __clang_hip_math.h:834

norm

__DEVICE__ double norm(int __dim, const double *__a)

Returns the square root of the sum of squares of any number of coordinates.

Definition: __clang_hip_math.h:1028

sin

__DEVICE__ double sin(double __x)

Returns the sine of x.

Definition: __clang_hip_math.h:1129

tan

__DEVICE__ double tan(double __x)

Returns the tangent of x.

Definition: __clang_hip_math.h:1163

hypot

__DEVICE__ double hypot(double __x, double __y)

Returns the square root of the sum of squares of x and y.

Definition: __clang_hip_math.h:904

fmin

__DEVICE__ double fmin(double __x, double __y)

Determine the minimum numeric value of x and y.

Definition: __clang_hip_math.h:893

jn

__DEVICE__ double jn(int __n, double __x)

Returns the value of the Bessel function of the first kind of order n for x.

Definition: __clang_hip_math.h:925

lround

__DEVICE__ long int lround(double __x)

Round to nearest integer value.

Definition: __clang_hip_math.h:976

tanh

__DEVICE__ double tanh(double __x)

Returns the hyperbolic tangent of x.

Definition: __clang_hip_math.h:1166

copysign

__DEVICE__ double copysign(double __x, double __y)

Create value with given magnitude, copying sign of second value.

Definition: __clang_hip_math.h:829

frexp

__DEVICE__ double frexp(double __x, int *__nptr)

Extract mantissa and exponent of x.

Definition: __clang_hip_math.h:899

acosh

__DEVICE__ double acosh(double __x)

Returns the nonnegative arc hyperbolic cosine of x.

Definition: __clang_hip_math.h:805

scalbln

__DEVICE__ double scalbln(double __x, long int __n)

Scale x by .

Definition: __clang_hip_math.h:1117

rint

__DEVICE__ double rint(double __x)

Round x to nearest integer value in floating-point.

Definition: __clang_hip_math.h:1086

logb

__DEVICE__ double logb(double __x)

Returns the floating point representation of the exponent of x.

Definition: __clang_hip_math.h:970

atan

__DEVICE__ double atan(double __x)

Returns the arc tangent of x.

Definition: __clang_hip_math.h:814

normcdf

__DEVICE__ double normcdf(double __x)

Returns the standard normal cumulative distribution function.

Definition: __clang_hip_math.h:1050

fma

__DEVICE__ double fma(double __x, double __y, double __z)

Returns as a single operation.

Definition: __clang_hip_math.h:885

erfc

__DEVICE__ double erfc(double __x)

Returns the complementary error function of x.

Definition: __clang_hip_math.h:852

y1

__DEVICE__ double y1(double __x)

Returns the value of the Bessel function of the second kind of order 1 for x.

Definition: __clang_hip_math.h:1178

asin

__DEVICE__ double asin(double __x)

Returns the arc sine of x.

Definition: __clang_hip_math.h:808

erfcinv

__DEVICE__ double erfcinv(double __x)

Returns the inverse complementary function of x.

Definition: __clang_hip_math.h:855

powi

__DEVICE__ double powi(double __x, int __y)

Returns the value of first argument to the power of second argument.

Definition: __clang_hip_math.h:1059

nextafter

__DEVICE__ double nextafter(double __x, double __y)

Returns next representable single-precision floating-point value after x.

Definition: __clang_hip_math.h:1023

cospi

__DEVICE__ double cospi(double __x)

Returns the cosine of .

Definition: __clang_hip_math.h:840

rsqrt

__DEVICE__ double rsqrt(double __x)

Returns the reciprocal of the square root of x.

Definition: __clang_hip_math.h:1114

pow

__DEVICE__ double pow(double __x, double __y)

Returns .

Definition: __clang_hip_math.h:1056

norm4d

__DEVICE__ double norm4d(double __x, double __y, double __z, double __w)

Returns the square root of the sum of squares of x, y, z and w.

Definition: __clang_hip_math.h:1045

remquo

__DEVICE__ double remquo(double __x, double __y, int *__quo)

Returns double-precision floating-point remainder and part of quotient.

Definition: __clang_hip_math.h:1070

nan

__DEVICE__ double nan(const char *__tagp)

Returns "Not a Number" value.

Definition: __clang_hip_math.h:992

rnorm

__DEVICE__ double rnorm(int __dim, const double *__a)

Returns the reciprocal of square root of the sum of squares of any number of coordinates.

Definition: __clang_hip_math.h:1089

sinpi

__DEVICE__ double sinpi(double __x)

Returns the hyperbolic sine of .

Definition: __clang_hip_math.h:1157

fmod

__DEVICE__ double fmod(double __x, double __y)

Returns the floating-point remainder of x / y.

Definition: __clang_hip_math.h:896

y0

__DEVICE__ double y0(double __x)

Returns the value of the Bessel function of the second kind of order 0 for x.

Definition: __clang_hip_math.h:1175

sqrt

__DEVICE__ double sqrt(double __x)

Returns the square root of x.

Definition: __clang_hip_math.h:1160

acos

__DEVICE__ double acos(double __x)

Returns the arc cosine of x.

Definition: __clang_hip_math.h:802

__finite

__DEVICE__ __RETURN_TYPE __finite(double __x)

Determine whether x is finite.

Definition: __clang_hip_math.h:910

lgamma

__DEVICE__ double lgamma(double __x)

Returns the natural logarithm of the absolute value of the gamma function of x.

Definition: __clang_hip_math.h:949

yn

__DEVICE__ double yn(int __n, double __x)

Returns the value of the Bessel function of the second kind of order n for x.

Definition: __clang_hip_math.h:1181

log10

__DEVICE__ double log10(double __x)

Returns the base 10 logarithm of x.

Definition: __clang_hip_math.h:961

cyl_bessel_i0

__DEVICE__ double cyl_bessel_i0(double __x)

Returns the value of the regular modified cylindrical Bessel function of order 0 for x.

Definition: __clang_hip_math.h:843

scalbn

__DEVICE__ double scalbn(double __x, int __n)

Scale x by .

Definition: __clang_hip_math.h:1123

cyl_bessel_i1

__DEVICE__ double cyl_bessel_i1(double __x)

Returns the value of the regular modified cylindrical Bessel function of order 1 for x.

Definition: __clang_hip_math.h:846

rcbrt

__DEVICE__ double rcbrt(double __x)

Returns the reciprocal cube root function.

Definition: __clang_hip_math.h:1062

rnorm3d

__DEVICE__ double rnorm3d(double __x, double __y, double __z)

Returns one over the square root of the sum of squares of x, y and z.

Definition: __clang_hip_math.h:1101

nearbyint

__DEVICE__ double nearbyint(double __x)

Round x to the nearest integer.

Definition: __clang_hip_math.h:1020

llround

__DEVICE__ long long int llround(double __x)

Round to nearest integer value.

Definition: __clang_hip_math.h:955

sincos

__DEVICE__ void sincos(double __x, double *__sinptr, double *__cosptr)

Returns the sine and cosine of x.

Definition: __clang_hip_math.h:1132

round

__DEVICE__ double round(double __x)

Round to nearest integer value in floating-point.

Definition: __clang_hip_math.h:1111

__signbit

__DEVICE__ __RETURN_TYPE __signbit(double __x)

Return the sign bit of x.

Definition: __clang_hip_math.h:1126

llrint

__DEVICE__ long long int llrint(double __x)

Round x to nearest integer value.

Definition: __clang_hip_math.h:952

remainder

__DEVICE__ double remainder(double __x, double __y)

Returns double-precision floating-point remainder.

Definition: __clang_hip_math.h:1065

fabs

__DEVICE__ double fabs(double __x)

Returns the absolute value of x.

Definition: __clang_hip_math.h:876

modf

__DEVICE__ double modf(double __x, double *__iptr)

Break down x into fractional and integral parts.

Definition: __clang_hip_math.h:979

rnorm4d

__DEVICE__ double rnorm4d(double __x, double __y, double __z, double __w)

Returns one over the square root of the sum of squares of x, y, z and w.

Definition: __clang_hip_math.h:1106

ilogb

__DEVICE__ int ilogb(double __x)

Returns the unbiased integer exponent of x.

Definition: __clang_hip_math.h:907

floor

__DEVICE__ double floor(double __x)

Returns the largest integer less than or equal to x.

Definition: __clang_hip_math.h:882

erfcx

__DEVICE__ double erfcx(double __x)

Returns the scaled complementary error function of x.

Definition: __clang_hip_math.h:858

atanh

__DEVICE__ double atanh(double __x)

Returns the arc hyperbolic tangent of x.

Definition: __clang_hip_math.h:820

erf

__DEVICE__ double erf(double __x)

Returns the error function of x.

Definition: __clang_hip_math.h:849

exp10

__DEVICE__ double exp10(double __x)

Returns .

Definition: __clang_hip_math.h:867

sincospi

__DEVICE__ void sincospi(double __x, double *__sinptr, double *__cosptr)

Returns the sine and cosine of .

Definition: __clang_hip_math.h:1143

__isinf

__DEVICE__ __RETURN_TYPE __isinf(double __x)

Determine whether x is infinite.

Definition: __clang_hip_math.h:913

log1p

__DEVICE__ double log1p(double __x)

Returns the natural logarithm of x + 1.

Definition: __clang_hip_math.h:964

expm1

__DEVICE__ double expm1(double __x)

Returns for x.

Definition: __clang_hip_math.h:873

exp2

__DEVICE__ double exp2(double __x)

Returns .

Definition: __clang_hip_math.h:870

trunc

__DEVICE__ double trunc(double __x)

Truncate x to the integral part.

Definition: __clang_hip_math.h:1172

asinh

__DEVICE__ double asinh(double __x)

Returns the arc hyperbolic sine of x.

Definition: __clang_hip_math.h:811

__dsub_rn

__DEVICE__ double __dsub_rn(double __x, double __y)

Subtract two floating-point values in round-to-nearest-even mode.

Definition: __clang_hip_math.h:1352

__dsqrt_rn

__DEVICE__ double __dsqrt_rn(double __x)

Returns in round-to-nearest-even mode.

Definition: __clang_hip_math.h:1325

__dadd_rn

__DEVICE__ double __dadd_rn(double __x, double __y)

Add two floating-point values in round-to-nearest-even mode.

Definition: __clang_hip_math.h:1233

__fma_rn

__DEVICE__ double __fma_rn(double __x, double __y, double __z)

Returns as a single operation in round-to-nearest-even mode.

Definition: __clang_hip_math.h:1379

__ddiv_rn

__DEVICE__ double __ddiv_rn(double __x, double __y)

Divide two floating-point values in round-to-nearest-even mode.

Definition: __clang_hip_math.h:1260

__dmul_rn

__DEVICE__ double __dmul_rn(double __x, double __y)

Multiply two floating-point values in round-to-nearest-even mode.

Definition: __clang_hip_math.h:1287

__drcp_rn

__DEVICE__ double __drcp_rn(double __x)

Returns 1 / x in round-to-nearest-even mode.

Definition: __clang_hip_math.h:1306

__isinff

__DEVICE__ __RETURN_TYPE __isinff(float __x)

Determine whether x is infinite.

Definition: __clang_hip_math.h:506

sinpif

__DEVICE__ float sinpif(float __x)

Returns the hyperbolic sine of .

Definition: __clang_hip_math.h:746

tanf

__DEVICE__ float tanf(float __x)

Returns the tangent of x.

Definition: __clang_hip_math.h:752

log2f

__DEVICE__ float log2f(float __x)

Returns the base 2 logarithm of x.

Definition: __clang_hip_math.h:557

y0f

__DEVICE__ float y0f(float __x)

Returns the value of the Bessel function of the second kind of order 0 for x.

Definition: __clang_hip_math.h:764

tanhf

__DEVICE__ float tanhf(float __x)

Returns the hyperbolic tangent of x.

Definition: __clang_hip_math.h:755

coshf

__DEVICE__ float coshf(float __x)

Returns the hyperbolic cosine of x.

Definition: __clang_hip_math.h:427

log10f

__DEVICE__ float log10f(float __x)

Returns the base 10 logarithm of x.

Definition: __clang_hip_math.h:551

j1f

__DEVICE__ float j1f(float __x)

Returns the value of the Bessel function of the first kind of order 1 for x.

Definition: __clang_hip_math.h:515

__finitef

__DEVICE__ __RETURN_TYPE __finitef(float __x)

Determine whether x is finite.

Definition: __clang_hip_math.h:503

ldexpf

__DEVICE__ float ldexpf(float __x, int __e)

Returns the value of for x and e.

Definition: __clang_hip_math.h:539

llroundf

__DEVICE__ long long int llroundf(float __x)

Round to nearest integer value.

Definition: __clang_hip_math.h:548

truncf

__DEVICE__ float truncf(float __x)

Truncate x to the integral part.

Definition: __clang_hip_math.h:761

remainderf

__DEVICE__ float remainderf(float __x, float __y)

Returns single-precision floating-point remainder.

Definition: __clang_hip_math.h:650

fabsf

__DEVICE__ float fabsf(float __x)

Returns the absolute value of x

Definition: __clang_hip_math.h:466

scalbnf

__DEVICE__ float scalbnf(float __x, int __n)

Scale x by .

Definition: __clang_hip_math.h:708

cyl_bessel_i0f

__DEVICE__ float cyl_bessel_i0f(float __x)

Returns the value of the regular modified cylindrical Bessel function of order 0 for x.

Definition: __clang_hip_math.h:433

nanf

__DEVICE__ float nanf(const char *__tagp __attribute__((nonnull)))

Returns "Not a Number" value.

Definition: __clang_hip_math.h:584

lgammaf

__DEVICE__ float lgammaf(float __x)

Returns the natural logarithm of the absolute value of the gamma function of x.

Definition: __clang_hip_math.h:542

cospif

__DEVICE__ float cospif(float __x)

Returns the cosine of .

Definition: __clang_hip_math.h:430

__signbitf

__DEVICE__ __RETURN_TYPE __signbitf(float __x)

Return the sign bit of x.

Definition: __clang_hip_math.h:711

frexpf

__DEVICE__ float frexpf(float __x, int *__nptr)

Extract mantissa and exponent of x.

Definition: __clang_hip_math.h:492

tgammaf

__DEVICE__ float tgammaf(float __x)

Returns the gamma function of x.

Definition: __clang_hip_math.h:758

erfinvf

__DEVICE__ float erfinvf(float __x)

Returns the inverse error function of x.

Definition: __clang_hip_math.h:451

modff

__DEVICE__ float modff(float __x, float *__iptr)

Break down x into fractional and integral parts.

Definition: __clang_hip_math.h:572

expm1f

__DEVICE__ float expm1f(float __x)

Returns .

Definition: __clang_hip_math.h:463

sinhf

__DEVICE__ float sinhf(float __x)

Returns the hyperbolic sine of x.

Definition: __clang_hip_math.h:743

y1f

__DEVICE__ float y1f(float __x)

Returns the value of the Bessel function of the second kind of order 1 for x.

Definition: __clang_hip_math.h:767

acosf

__DEVICE__ float acosf(float __x)

Returns the arc cosine of x.

Definition: __clang_hip_math.h:394

fmaf

__DEVICE__ float fmaf(float __x, float __y, float __z)

Returns as a single operation.

Definition: __clang_hip_math.h:478

cyl_bessel_i1f

__DEVICE__ float cyl_bessel_i1f(float __x)

Returns the value of the regular modified cylindrical Bessel function of order 1 for x.

Definition: __clang_hip_math.h:436

fmodf

__DEVICE__ float fmodf(float __x, float __y)

Returns the floating-point remainder of x / y.

Definition: __clang_hip_math.h:489

log1pf

__DEVICE__ float log1pf(float __x)

Returns the natural logarithm of x + 1.

Definition: __clang_hip_math.h:554

atan2f

__DEVICE__ float atan2f(float __x, float __y)

Returns the arc tangent of the ratio of x and y.

Definition: __clang_hip_math.h:406

copysignf

__DEVICE__ float copysignf(float __x, float __y)

Create value with given magnitude, copying sign of second value.

Definition: __clang_hip_math.h:421

rnormf

__DEVICE__ float rnormf(int __dim, const float *__a)

Returns the reciprocal of square root of the sum of squares of any number of coordinates.

Definition: __clang_hip_math.h:684

rnorm4df

__DEVICE__ float rnorm4df(float __x, float __y, float __z, float __w)

Returns one over the square root of the sum of squares of x, y, z and w.

Definition: __clang_hip_math.h:679

erff

__DEVICE__ float erff(float __x)

Returns the error function of x.

Definition: __clang_hip_math.h:448

atanf

__DEVICE__ float atanf(float __x)

Returns the arc tangent of x.

Definition: __clang_hip_math.h:409

rnorm3df

__DEVICE__ float rnorm3df(float __x, float __y, float __z)

Returns one over the square root of the sum of squares of x, y and z.

Definition: __clang_hip_math.h:674

erfcxf

__DEVICE__ float erfcxf(float __x)

Returns the scaled complementary error function of x.

Definition: __clang_hip_math.h:445

erfcinvf

__DEVICE__ float erfcinvf(float __x)

Returns the inverse complementary function of x.

Definition: __clang_hip_math.h:442

asinf

__DEVICE__ float asinf(float __x)

Returns the arc sine of x.

Definition: __clang_hip_math.h:400

lroundf

__DEVICE__ long int lroundf(float __x)

Round to nearest integer value.

Definition: __clang_hip_math.h:569

norm4df

__DEVICE__ float norm4df(float __x, float __y, float __z, float __w)

Returns the square root of the sum of squares of x, y, z and w.

Definition: __clang_hip_math.h:618

__isnanf

__DEVICE__ __RETURN_TYPE __isnanf(float __x)

Determine whether x is a NaN.

Definition: __clang_hip_math.h:509

ynf

__DEVICE__ float ynf(int __n, float __x)

Returns the value of the Bessel function of the second kind of order n for x.

Definition: __clang_hip_math.h:770

powf

__DEVICE__ float powf(float __x, float __y)

Returns .

Definition: __clang_hip_math.h:641

sinf

__DEVICE__ float sinf(float __x)

Returns the sine of x.

Definition: __clang_hip_math.h:740

remquof

__DEVICE__ float remquof(float __x, float __y, int *__quo)

Returns single-precision floating-point remainder and part of quotient.

Definition: __clang_hip_math.h:655

hypotf

__DEVICE__ float hypotf(float __x, float __y)

Returns the square root of the sum of squares of x and y.

Definition: __clang_hip_math.h:497

sincosf

__DEVICE__ void sincosf(float __x, float *__sinptr, float *__cosptr)

Returns the sine and cosine of x.

Definition: __clang_hip_math.h:714

exp10f

__DEVICE__ float exp10f(float __x)

Returns .

Definition: __clang_hip_math.h:454

fmaxf

__DEVICE__ float fmaxf(float __x, float __y)

Determine the maximum numeric value of x and y.

Definition: __clang_hip_math.h:483

fminf

__DEVICE__ float fminf(float __x, float __y)

Determine the minimum numeric value of x and y.

Definition: __clang_hip_math.h:486

logf

__DEVICE__ float logf(float __x)

Returns the natural logarithm of x.

Definition: __clang_hip_math.h:563

erfcf

__DEVICE__ float erfcf(float __x)

Returns the complementary error function of x.

Definition: __clang_hip_math.h:439

atanhf

__DEVICE__ float atanhf(float __x)

Returns the arc hyperbolic tangent of x.

Definition: __clang_hip_math.h:412

asinhf

__DEVICE__ float asinhf(float __x)

Returns the arc hyperbolic sine of x.

Definition: __clang_hip_math.h:403

j0f

__DEVICE__ float j0f(float __x)

Returns the value of the Bessel function of the first kind of order 0 for x.

Definition: __clang_hip_math.h:512

rsqrtf

__DEVICE__ float rsqrtf(float __x)

Returns the reciprocal of the square root of x.

Definition: __clang_hip_math.h:699

jnf

__DEVICE__ float jnf(int __n, float __x)

Returns the value of the Bessel function of the first kind of order n for x.

Definition: __clang_hip_math.h:518

logbf

__DEVICE__ float logbf(float __x)

Returns the floating point representation of the exponent of x.

Definition: __clang_hip_math.h:560

rhypotf

__DEVICE__ float rhypotf(float __x, float __y)

Returns one over the square root of the sum of squares of x and y.

Definition: __clang_hip_math.h:668

exp2f

__DEVICE__ float exp2f(float __x)

Returns .

Definition: __clang_hip_math.h:457

powif

__DEVICE__ float powif(float __x, int __y)

Returns the value of first argument to the power of second argument.

Definition: __clang_hip_math.h:644

ceilf

__DEVICE__ float ceilf(float __x)

Returns ceiling of x.

Definition: __clang_hip_math.h:418

normcdfinvf

__DEVICE__ float normcdfinvf(float __x)

Returns the inverse of the standard normal cumulative distribution function.

Definition: __clang_hip_math.h:626

norm3df

__DEVICE__ float norm3df(float __x, float __y, float __z)

Returns the square root of the sum of squares of x, y and z.

Definition: __clang_hip_math.h:613

fdimf

__DEVICE__ float fdimf(float __x, float __y)

Returns the positive difference between x and y.

Definition: __clang_hip_math.h:469

normf

__DEVICE__ float normf(int __dim, const float *__a)

Returns the square root of the sum of squares of any number of coordinates.

Definition: __clang_hip_math.h:629

nearbyintf

__DEVICE__ float nearbyintf(float __x)

Round x to the nearest integer.

Definition: __clang_hip_math.h:605

ilogbf

__DEVICE__ int ilogbf(float __x)

Returns the unbiased integer exponent of x.

Definition: __clang_hip_math.h:500

floorf

__DEVICE__ float floorf(float __x)

Returns the largest integer less than or equal to x.

Definition: __clang_hip_math.h:475

sqrtf

__DEVICE__ float sqrtf(float __x)

Returns the square root of x.

Definition: __clang_hip_math.h:749

roundf

__DEVICE__ float roundf(float __x)

Round to nearest integer value in floating-point.

Definition: __clang_hip_math.h:696

sincospif

__DEVICE__ void sincospif(float __x, float *__sinptr, float *__cosptr)

Returns the sine and cosine of .

Definition: __clang_hip_math.h:729

lrintf

__DEVICE__ long int lrintf(float __x)

Round x to nearest integer value.

Definition: __clang_hip_math.h:566

acoshf

__DEVICE__ float acoshf(float __x)

Returns the nonnegative arc hyperbolic cosine of x.

Definition: __clang_hip_math.h:397

cosf

__DEVICE__ float cosf(float __x)

Returns the cosine of x.

Definition: __clang_hip_math.h:424

expf

__DEVICE__ float expf(float __x)

Returns .

Definition: __clang_hip_math.h:460

nextafterf

__DEVICE__ float nextafterf(float __x, float __y)

Returns next representable single-precision floating-point value after x.

Definition: __clang_hip_math.h:608

llrintf

__DEVICE__ long long int llrintf(float __x)

Round x to nearest integer value.

Definition: __clang_hip_math.h:545

fdividef

__DEVICE__ float fdividef(float __x, float __y)

Divide two floating point values.

Definition: __clang_hip_math.h:472

rcbrtf

__DEVICE__ float rcbrtf(float __x)

Returns the reciprocal cube root function.

Definition: __clang_hip_math.h:647

cbrtf

__DEVICE__ float cbrtf(float __x)

Returns the cube root of x.

Definition: __clang_hip_math.h:415

scalblnf

__DEVICE__ float scalblnf(float __x, long int __n)

Scale x by .

Definition: __clang_hip_math.h:702

rintf

__DEVICE__ float rintf(float __x)

Round x to nearest integer value in floating-point.

Definition: __clang_hip_math.h:671

normcdff

__DEVICE__ float normcdff(float __x)

Returns the standard normal cumulative distribution function.

Definition: __clang_hip_math.h:623

__fdiv_rn

__DEVICE__ float __fdiv_rn(float __x, float __y)

Divide two floating-point values in round-to-nearest-even mode.

Definition: __clang_hip_math.h:222

__sinf

__DEVICE__ float __sinf(float __x)

Returns the fast approximate sine of x.

Definition: __clang_hip_math.h:360

__cosf

__DEVICE__ float __cosf(float __x)

Returns the fast approximate cosine of x.

Definition: __clang_hip_math.h:173

__fdividef

__DEVICE__ float __fdividef(float __x, float __y)

Returns the fast approximate of x / y.

Definition: __clang_hip_math.h:226

__frsqrt_rn

__DEVICE__ float __frsqrt_rn(float __x)

Returns in round-to-nearest-even mode.

Definition: __clang_hip_math.h:297

__log2f

__DEVICE__ float __log2f(float __x)

Returns the fast approximate for base 2 logarithm of x.

Definition: __clang_hip_math.h:342

__exp10f

__DEVICE__ float __exp10f(float __x)

Returns the fast approximate for .

Definition: __clang_hip_math.h:176

__frcp_rn

__DEVICE__ float __frcp_rn(float __x)

Returns 1 / x in round-to-nearest-even mod.

Definition: __clang_hip_math.h:292

__fsub_rn

__DEVICE__ float __fsub_rn(float __x, float __y)

Subtract two floating-point values in round-to-nearest-even mode.

Definition: __clang_hip_math.h:334

__tanf

__DEVICE__ float __tanf(float __x)

Returns the fast approximate tangent of x.

Definition: __clang_hip_math.h:363

__fsqrt_rn

__DEVICE__ float __fsqrt_rn(float __x)

Returns in round-to-nearest-even mode.

Definition: __clang_hip_math.h:315

__fmaf_rn

__DEVICE__ float __fmaf_rn(float __x, float __y, float __z)

Returns as a single operation, in round-to-nearest-even mode.

Definition: __clang_hip_math.h:252

__fadd_rn

__DEVICE__ float __fadd_rn(float __x, float __y)

Add two floating-point values in round-to-nearest-even mode.

Definition: __clang_hip_math.h:203

__expf

__DEVICE__ float __expf(float __x)

Returns the fast approximate for .

Definition: __clang_hip_math.h:182

__logf

__DEVICE__ float __logf(float __x)

Returns the fast approximate for natural logarithm of x.

Definition: __clang_hip_math.h:345

__sincosf

__DEVICE__ void __sincosf(float __x, float *__sinptr, float *__cosptr)

Returns the fast approximate of sine and cosine of x.

Definition: __clang_hip_math.h:354

__log10f

__DEVICE__ float __log10f(float __x)

Returns the fast approximate for base 10 logarithm of x.

Definition: __clang_hip_math.h:339

__fmul_rn

__DEVICE__ float __fmul_rn(float __x, float __y)

Multiply two floating-point values in round-to-nearest-even mode.

Definition: __clang_hip_math.h:273

__saturatef

__DEVICE__ float __saturatef(float __x)

Clamp x to [+0.0, 1.0].

Definition: __clang_hip_math.h:351

__powf

__DEVICE__ float __powf(float __x, float __y)

Returns the fast approximate of .

Definition: __clang_hip_math.h:348

__make_mantissa_base10

__DEVICE__ uint64_t __make_mantissa_base10(const char *__tagp __attribute__((nonnull)))

Make base 10 (decimal) mantissa char array.

Definition: __clang_hip_math.h:111

__make_mantissa_base8

__DEVICE__ uint64_t __make_mantissa_base8(const char *__tagp __attribute__((nonnull)))

Make base 8 (octal) mantissa from char array.

Definition: __clang_hip_math.h:94

__make_mantissa

__DEVICE__ uint64_t __make_mantissa(const char *__tagp __attribute__((nonnull)))

Make mantissa based on number format char array.

Definition: __clang_hip_math.h:149

__make_mantissa_base16

__DEVICE__ uint64_t __make_mantissa_base16(const char *__tagp __attribute__((nonnull)))

Make base 16 (hexadecimal) mantissa char array.

Definition: __clang_hip_math.h:128