Skip to content

Commit 6abcd34

Browse files
oleksandr-pavlykoleksandr-pavlyk
authored
Merge pull request #1630 from IntelPython/fix-cuda-linking-2024-1
Fix compiling for CUDA with 2024.1 compiler
2 parents fda9837 + fc4a612 commit 6abcd34

File tree

25 files changed

+160
-141
lines changed

25 files changed

+160
-141
lines changed

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp

Lines changed: 4 additions & 45 deletions
Original file line numberDiff line numberDiff line change
@@ -31,8 +31,8 @@
3131
#include <sycl/sycl.hpp>
3232
#include <type_traits>
3333

34+
#include "cabs_impl.hpp"
3435
#include "kernels/elementwise_functions/common.hpp"
35-
#include "sycl_complex.hpp"
3636

3737
#include "kernels/dpctl_tensor_types.hpp"
3838
#include "utils/offset_utils.hpp"
@@ -73,59 +73,18 @@ template <typename argT, typename resT> struct AbsFunctor
7373
}
7474
else {
7575
if constexpr (is_complex<argT>::value) {
76-
return cabs(x);
76+
return detail::cabs(x);
7777
}
7878
else if constexpr (std::is_same_v<argT, sycl::half> ||
7979
std::is_floating_point_v<argT>)
8080
{
8181
return (std::signbit(x) ? -x : x);
8282
}
8383
else {
84-
return std::abs(x);
84+
return sycl::abs(x);
8585
}
8686
}
8787
}
88-
89-
private:
90-
template <typename realT> realT cabs(std::complex<realT> const &z) const
91-
{
92-
// Special values for cabs( x + y * 1j):
93-
// * If x is either +infinity or -infinity and y is any value
94-
// (including NaN), the result is +infinity.
95-
// * If x is any value (including NaN) and y is either +infinity or
96-
// -infinity, the result is +infinity.
97-
// * If x is either +0 or -0, the result is equal to abs(y).
98-
// * If y is either +0 or -0, the result is equal to abs(x).
99-
// * If x is NaN and y is a finite number, the result is NaN.
100-
// * If x is a finite number and y is NaN, the result is NaN.
101-
// * If x is NaN and y is NaN, the result is NaN.
102-
103-
const realT x = std::real(z);
104-
const realT y = std::imag(z);
105-
106-
constexpr realT q_nan = std::numeric_limits<realT>::quiet_NaN();
107-
constexpr realT p_inf = std::numeric_limits<realT>::infinity();
108-
109-
if (std::isinf(x)) {
110-
return p_inf;
111-
}
112-
else if (std::isinf(y)) {
113-
return p_inf;
114-
}
115-
else if (std::isnan(x)) {
116-
return q_nan;
117-
}
118-
else if (std::isnan(y)) {
119-
return q_nan;
120-
}
121-
else {
122-
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
123-
return exprm_ns::abs(exprm_ns::complex<realT>(z));
124-
#else
125-
return std::hypot(std::real(z), std::imag(z));
126-
#endif
127-
}
128-
}
12988
};
13089

13190
template <typename argT,
@@ -195,7 +154,7 @@ template <typename fnT, typename T> struct AbsContigFactory
195154

196155
template <typename fnT, typename T> struct AbsTypeMapFactory
197156
{
198-
/*! @brief get typeid for output type of std::abs(T x) */
157+
/*! @brief get typeid for output type of abs(T x) */
199158
std::enable_if_t<std::is_same<fnT, int>::value, int> get()
200159
{
201160
using rT = typename AbsOutputType<T>::value_type;

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -102,23 +102,23 @@ template <typename argT, typename resT> struct AcosFunctor
102102
*/
103103
constexpr realT r_eps =
104104
realT(1) / std::numeric_limits<realT>::epsilon();
105-
if (std::abs(x) > r_eps || std::abs(y) > r_eps) {
105+
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
106106
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
107107
using sycl_complexT = exprm_ns::complex<realT>;
108108
sycl_complexT log_in =
109109
exprm_ns::log(exprm_ns::complex<realT>(in));
110110

111111
const realT wx = log_in.real();
112112
const realT wy = log_in.imag();
113-
const realT rx = std::abs(wy);
113+
const realT rx = sycl::fabs(wy);
114114

115115
realT ry = wx + std::log(realT(2));
116116
return resT{rx, (std::signbit(y)) ? ry : -ry};
117117
#else
118118
resT log_in = std::log(in);
119119
const realT wx = std::real(log_in);
120120
const realT wy = std::imag(log_in);
121-
const realT rx = std::abs(wy);
121+
const realT rx = sycl::fabs(wy);
122122

123123
realT ry = wx + std::log(realT(2));
124124
return resT{rx, (std::signbit(y)) ? ry : -ry};

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -109,7 +109,7 @@ template <typename argT, typename resT> struct AcoshFunctor
109109
/*
110110
* For large x or y including acos(+-Inf + I*+-Inf)
111111
*/
112-
if (std::abs(x) > r_eps || std::abs(y) > r_eps) {
112+
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
113113
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
114114
using sycl_complexT = typename exprm_ns::complex<realT>;
115115
const sycl_complexT log_in = exprm_ns::log(sycl_complexT(in));
@@ -120,7 +120,7 @@ template <typename argT, typename resT> struct AcoshFunctor
120120
const realT wx = std::real(log_in);
121121
const realT wy = std::imag(log_in);
122122
#endif
123-
const realT rx = std::abs(wy);
123+
const realT rx = sycl::fabs(wy);
124124
realT ry = wx + std::log(realT(2));
125125
acos_in = resT{rx, (std::signbit(y)) ? ry : -ry};
126126
}
@@ -145,15 +145,15 @@ template <typename argT, typename resT> struct AcoshFunctor
145145
/* acosh(NaN + I*+-Inf) = +Inf + I*NaN */
146146
/* acosh(+-Inf + I*NaN) = +Inf + I*NaN */
147147
if (std::isnan(rx)) {
148-
return resT{std::abs(ry), rx};
148+
return resT{sycl::fabs(ry), rx};
149149
}
150150
/* acosh(0 + I*NaN) = NaN + I*NaN */
151151
if (std::isnan(ry)) {
152152
return resT{ry, ry};
153153
}
154154
/* ordinary cases */
155-
const realT res_im = std::copysign(rx, std::imag(in));
156-
return resT{std::abs(ry), res_im};
155+
const realT res_im = sycl::copysign(rx, std::imag(in));
156+
return resT{sycl::fabs(ry), res_im};
157157
}
158158
else {
159159
static_assert(std::is_floating_point_v<argT> ||

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -116,7 +116,7 @@ template <typename argT, typename resT> struct AsinFunctor
116116
*/
117117
constexpr realT r_eps =
118118
realT(1) / std::numeric_limits<realT>::epsilon();
119-
if (std::abs(x) > r_eps || std::abs(y) > r_eps) {
119+
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
120120
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
121121
using sycl_complexT = exprm_ns::complex<realT>;
122122
const sycl_complexT z{x, y};
@@ -145,8 +145,8 @@ template <typename argT, typename resT> struct AsinFunctor
145145
wy = std::imag(log_mz);
146146
}
147147
#endif
148-
const realT asinh_re = std::copysign(wx, x);
149-
const realT asinh_im = std::copysign(wy, y);
148+
const realT asinh_re = sycl::copysign(wx, x);
149+
const realT asinh_im = sycl::copysign(wy, y);
150150
return resT{asinh_im, asinh_re};
151151
}
152152
/* ordinary cases */

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ template <typename argT, typename resT> struct AsinhFunctor
105105
constexpr realT r_eps =
106106
realT(1) / std::numeric_limits<realT>::epsilon();
107107

108-
if (std::abs(x) > r_eps || std::abs(y) > r_eps) {
108+
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
109109
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
110110
using sycl_complexT = exprm_ns::complex<realT>;
111111
sycl_complexT log_in = (std::signbit(x))
@@ -118,8 +118,8 @@ template <typename argT, typename resT> struct AsinhFunctor
118118
realT wx = std::real(log_in) + std::log(realT(2));
119119
realT wy = std::imag(log_in);
120120
#endif
121-
const realT res_re = std::copysign(wx, x);
122-
const realT res_im = std::copysign(wy, y);
121+
const realT res_re = sycl::copysign(wx, x);
122+
const realT res_im = sycl::copysign(wy, y);
123123
return resT{res_re, res_im};
124124
}
125125

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -84,8 +84,8 @@ template <typename argT, typename resT> struct AtanFunctor
8484
if (std::isinf(y)) {
8585
const realT pi_half = std::atan(realT(1)) * 2;
8686

87-
const realT atanh_re = std::copysign(realT(0), x);
88-
const realT atanh_im = std::copysign(pi_half, y);
87+
const realT atanh_re = sycl::copysign(realT(0), x);
88+
const realT atanh_im = sycl::copysign(pi_half, y);
8989
return resT{atanh_im, atanh_re};
9090
}
9191
/*
@@ -96,7 +96,7 @@ template <typename argT, typename resT> struct AtanFunctor
9696
else if (std::isnan(y)) {
9797
/* atanh(+-Inf + I*NaN) = +-0 + I*NaN */
9898
if (std::isinf(x)) {
99-
const realT atanh_re = std::copysign(realT(0), x);
99+
const realT atanh_re = sycl::copysign(realT(0), x);
100100
const realT atanh_im = q_nan;
101101
return resT{atanh_im, atanh_re};
102102
}
@@ -118,11 +118,11 @@ template <typename argT, typename resT> struct AtanFunctor
118118
*/
119119
constexpr realT r_eps =
120120
realT(1) / std::numeric_limits<realT>::epsilon();
121-
if (std::abs(x) > r_eps || std::abs(y) > r_eps) {
121+
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
122122
const realT pi_half = std::atan(realT(1)) * 2;
123123

124124
const realT atanh_re = realT(0);
125-
const realT atanh_im = std::copysign(pi_half, y);
125+
const realT atanh_im = sycl::copysign(pi_half, y);
126126
return resT{atanh_im, atanh_re};
127127
}
128128
/* ordinary cases */

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan2.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -58,7 +58,7 @@ template <typename argT1, typename argT2, typename resT> struct Atan2Functor
5858
{
5959
if (std::isinf(in2) && !std::signbit(in2)) {
6060
if (std::isfinite(in1)) {
61-
return std::copysign(resT(0), in1);
61+
return sycl::copysign(resT(0), in1);
6262
}
6363
}
6464
return std::atan2(in1, in2);

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -78,8 +78,8 @@ template <typename argT, typename resT> struct AtanhFunctor
7878
if (std::isinf(y)) {
7979
const realT pi_half = std::atan(realT(1)) * 2;
8080

81-
const realT res_re = std::copysign(realT(0), x);
82-
const realT res_im = std::copysign(pi_half, y);
81+
const realT res_re = sycl::copysign(realT(0), x);
82+
const realT res_im = sycl::copysign(pi_half, y);
8383
return resT{res_re, res_im};
8484
}
8585
/*
@@ -90,7 +90,7 @@ template <typename argT, typename resT> struct AtanhFunctor
9090
else if (std::isnan(y)) {
9191
/* atanh(+-Inf + I*NaN) = +-0 + I*NaN */
9292
if (std::isinf(x)) {
93-
const realT res_re = std::copysign(realT(0), x);
93+
const realT res_re = sycl::copysign(realT(0), x);
9494
return resT{res_re, q_nan};
9595
}
9696
/* atanh(+-0 + I*NaN) = +-0 + I*NaN */
@@ -111,11 +111,12 @@ template <typename argT, typename resT> struct AtanhFunctor
111111
*/
112112
const realT RECIP_EPSILON =
113113
realT(1) / std::numeric_limits<realT>::epsilon();
114-
if (std::abs(x) > RECIP_EPSILON || std::abs(y) > RECIP_EPSILON) {
114+
if (sycl::fabs(x) > RECIP_EPSILON || sycl::fabs(y) > RECIP_EPSILON)
115+
{
115116
const realT pi_half = std::atan(realT(1)) * 2;
116117

117118
const realT res_re = realT(0);
118-
const realT res_im = std::copysign(pi_half, y);
119+
const realT res_im = sycl::copysign(pi_half, y);
119120
return resT{res_re, res_im};
120121
}
121122
/* ordinary cases */

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/cabs_impl.hpp

Lines changed: 84 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,84 @@
1+
//===------- cabs_impl.hpp - Implementation of cabs -------*-C++-*/===//
2+
//
3+
// Data Parallel Control (dpctl)
4+
//
5+
// Copyright 2020-2024 Intel Corporation
6+
//
7+
// Licensed under the Apache License, Version 2.0 (the "License");
8+
// you may not use this file except in compliance with the License.
9+
// You may obtain a copy of the License at
10+
//
11+
// http://www.apache.org/licenses/LICENSE-2.0
12+
//
13+
// Unless required by applicable law or agreed to in writing, software
14+
// distributed under the License is distributed on an "AS IS" BASIS,
15+
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
16+
// See the License for the specific language governing permissions and
17+
// limitations under the License.
18+
//
19+
//===----------------------------------------------------------------------===//
20+
///
21+
/// \file
22+
/// This file defines an implementation of the complex absolute value.
23+
//===----------------------------------------------------------------------===//
24+
25+
#pragma once
26+
#include <cmath>
27+
#include <complex>
28+
#include <limits>
29+
30+
#include "sycl_complex.hpp"
31+
32+
namespace dpctl
33+
{
34+
namespace tensor
35+
{
36+
namespace kernels
37+
{
38+
namespace detail
39+
{
40+
41+
template <typename realT> realT cabs(std::complex<realT> const &z)
42+
{
43+
// Special values for cabs( x + y * 1j):
44+
// * If x is either +infinity or -infinity and y is any value
45+
// (including NaN), the result is +infinity.
46+
// * If x is any value (including NaN) and y is either +infinity or
47+
// -infinity, the result is +infinity.
48+
// * If x is either +0 or -0, the result is equal to abs(y).
49+
// * If y is either +0 or -0, the result is equal to abs(x).
50+
// * If x is NaN and y is a finite number, the result is NaN.
51+
// * If x is a finite number and y is NaN, the result is NaN.
52+
// * If x is NaN and y is NaN, the result is NaN.
53+
54+
const realT x = std::real(z);
55+
const realT y = std::imag(z);
56+
57+
constexpr realT q_nan = std::numeric_limits<realT>::quiet_NaN();
58+
constexpr realT p_inf = std::numeric_limits<realT>::infinity();
59+
60+
if (std::isinf(x)) {
61+
return p_inf;
62+
}
63+
else if (std::isinf(y)) {
64+
return p_inf;
65+
}
66+
else if (std::isnan(x)) {
67+
return q_nan;
68+
}
69+
else if (std::isnan(y)) {
70+
return q_nan;
71+
}
72+
else {
73+
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
74+
return exprm_ns::abs(exprm_ns::complex<realT>(z));
75+
#else
76+
return std::hypot(std::real(z), std::imag(z));
77+
#endif
78+
}
79+
}
80+
81+
} // namespace detail
82+
} // namespace kernels
83+
} // namespace tensor
84+
} // namespace dpctl

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/ceil.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -71,7 +71,7 @@ template <typename argT, typename resT> struct CeilFunctor
7171
if (in == 0) {
7272
return in;
7373
}
74-
return std::ceil(in);
74+
return sycl::ceil(in);
7575
}
7676
}
7777
};

‎dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -109,7 +109,7 @@ template <typename argT, typename resT> struct CosFunctor
109109
*/
110110
if (x == realT(0) && !yfinite) {
111111
const realT y_m_y = (y - y);
112-
const realT res_im = std::copysign(realT(0), x * y_m_y);
112+
const realT res_im = sycl::copysign(realT(0), x * y_m_y);
113113
return resT{y_m_y, res_im};
114114
}
115115

@@ -120,7 +120,7 @@ template <typename argT, typename resT> struct CosFunctor
120120
* The sign of 0 in the result is unspecified.
121121
*/
122122
if (y == realT(0) && !xfinite) {
123-
const realT res_im = std::copysign(realT(0), x) * y;
123+
const realT res_im = sycl::copysign(realT(0), x) * y;
124124
return resT{x * x, res_im};
125125
}
126126

@@ -144,7 +144,7 @@ template <typename argT, typename resT> struct CosFunctor
144144
*/
145145
if (std::isinf(x)) {
146146
if (!yfinite) {
147-
return resT{x * x, std::copysign(q_nan, x)};
147+
return resT{x * x, sycl::copysign(q_nan, x)};
148148
}
149149
return resT{(x * x) * std::cos(y), x * std::sin(y)};
150150
}

0 commit comments

Comments
 (0)