Skip to content

Fix compiling for CUDA with 2024.1 compiler #1630

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Apr 3, 2024
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
Expand Up @@ -31,8 +31,8 @@
#include <sycl/sycl.hpp>
#include <type_traits>

#include "cabs_impl.hpp"
#include "kernels/elementwise_functions/common.hpp"
#include "sycl_complex.hpp"

#include "kernels/dpctl_tensor_types.hpp"
#include "utils/offset_utils.hpp"
Expand Down Expand Up @@ -73,59 +73,18 @@ template <typename argT, typename resT> struct AbsFunctor
}
else {
if constexpr (is_complex<argT>::value) {
return cabs(x);
return detail::cabs(x);
}
else if constexpr (std::is_same_v<argT, sycl::half> ||
std::is_floating_point_v<argT>)
{
return (std::signbit(x) ? -x : x);
}
else {
return std::abs(x);
return sycl::abs(x);
}
}
}

private:
template <typename realT> realT cabs(std::complex<realT> const &z) const
{
// Special values for cabs( x + y * 1j):
// * If x is either +infinity or -infinity and y is any value
// (including NaN), the result is +infinity.
// * If x is any value (including NaN) and y is either +infinity or
// -infinity, the result is +infinity.
// * If x is either +0 or -0, the result is equal to abs(y).
// * If y is either +0 or -0, the result is equal to abs(x).
// * If x is NaN and y is a finite number, the result is NaN.
// * If x is a finite number and y is NaN, the result is NaN.
// * If x is NaN and y is NaN, the result is NaN.

const realT x = std::real(z);
const realT y = std::imag(z);

constexpr realT q_nan = std::numeric_limits<realT>::quiet_NaN();
constexpr realT p_inf = std::numeric_limits<realT>::infinity();

if (std::isinf(x)) {
return p_inf;
}
else if (std::isinf(y)) {
return p_inf;
}
else if (std::isnan(x)) {
return q_nan;
}
else if (std::isnan(y)) {
return q_nan;
}
else {
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
return exprm_ns::abs(exprm_ns::complex<realT>(z));
#else
return std::hypot(std::real(z), std::imag(z));
#endif
}
}
};

template <typename argT,
Expand Down Expand Up @@ -195,7 +154,7 @@ template <typename fnT, typename T> struct AbsContigFactory

template <typename fnT, typename T> struct AbsTypeMapFactory
{
/*! @brief get typeid for output type of std::abs(T x) */
/*! @brief get typeid for output type of abs(T x) */
std::enable_if_t<std::is_same<fnT, int>::value, int> get()
{
using rT = typename AbsOutputType<T>::value_type;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -102,23 +102,23 @@ template <typename argT, typename resT> struct AcosFunctor
*/
constexpr realT r_eps =
realT(1) / std::numeric_limits<realT>::epsilon();
if (std::abs(x) > r_eps || std::abs(y) > r_eps) {
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using sycl_complexT = exprm_ns::complex<realT>;
sycl_complexT log_in =
exprm_ns::log(exprm_ns::complex<realT>(in));

const realT wx = log_in.real();
const realT wy = log_in.imag();
const realT rx = std::abs(wy);
const realT rx = sycl::fabs(wy);

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

realT ry = wx + std::log(realT(2));
return resT{rx, (std::signbit(y)) ? ry : -ry};
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -109,7 +109,7 @@ template <typename argT, typename resT> struct AcoshFunctor
/*
* For large x or y including acos(+-Inf + I*+-Inf)
*/
if (std::abs(x) > r_eps || std::abs(y) > r_eps) {
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using sycl_complexT = typename exprm_ns::complex<realT>;
const sycl_complexT log_in = exprm_ns::log(sycl_complexT(in));
Expand All @@ -120,7 +120,7 @@ template <typename argT, typename resT> struct AcoshFunctor
const realT wx = std::real(log_in);
const realT wy = std::imag(log_in);
#endif
const realT rx = std::abs(wy);
const realT rx = sycl::fabs(wy);
realT ry = wx + std::log(realT(2));
acos_in = resT{rx, (std::signbit(y)) ? ry : -ry};
}
Expand All @@ -145,15 +145,15 @@ template <typename argT, typename resT> struct AcoshFunctor
/* acosh(NaN + I*+-Inf) = +Inf + I*NaN */
/* acosh(+-Inf + I*NaN) = +Inf + I*NaN */
if (std::isnan(rx)) {
return resT{std::abs(ry), rx};
return resT{sycl::fabs(ry), rx};
}
/* acosh(0 + I*NaN) = NaN + I*NaN */
if (std::isnan(ry)) {
return resT{ry, ry};
}
/* ordinary cases */
const realT res_im = std::copysign(rx, std::imag(in));
return resT{std::abs(ry), res_im};
const realT res_im = sycl::copysign(rx, std::imag(in));
return resT{sycl::fabs(ry), res_im};
}
else {
static_assert(std::is_floating_point_v<argT> ||
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -116,7 +116,7 @@ template <typename argT, typename resT> struct AsinFunctor
*/
constexpr realT r_eps =
realT(1) / std::numeric_limits<realT>::epsilon();
if (std::abs(x) > r_eps || std::abs(y) > r_eps) {
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using sycl_complexT = exprm_ns::complex<realT>;
const sycl_complexT z{x, y};
Expand Down Expand Up @@ -145,8 +145,8 @@ template <typename argT, typename resT> struct AsinFunctor
wy = std::imag(log_mz);
}
#endif
const realT asinh_re = std::copysign(wx, x);
const realT asinh_im = std::copysign(wy, y);
const realT asinh_re = sycl::copysign(wx, x);
const realT asinh_im = sycl::copysign(wy, y);
return resT{asinh_im, asinh_re};
}
/* ordinary cases */
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -105,7 +105,7 @@ template <typename argT, typename resT> struct AsinhFunctor
constexpr realT r_eps =
realT(1) / std::numeric_limits<realT>::epsilon();

if (std::abs(x) > r_eps || std::abs(y) > r_eps) {
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using sycl_complexT = exprm_ns::complex<realT>;
sycl_complexT log_in = (std::signbit(x))
Expand All @@ -118,8 +118,8 @@ template <typename argT, typename resT> struct AsinhFunctor
realT wx = std::real(log_in) + std::log(realT(2));
realT wy = std::imag(log_in);
#endif
const realT res_re = std::copysign(wx, x);
const realT res_im = std::copysign(wy, y);
const realT res_re = sycl::copysign(wx, x);
const realT res_im = sycl::copysign(wy, y);
return resT{res_re, res_im};
}

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -84,8 +84,8 @@ template <typename argT, typename resT> struct AtanFunctor
if (std::isinf(y)) {
const realT pi_half = std::atan(realT(1)) * 2;

const realT atanh_re = std::copysign(realT(0), x);
const realT atanh_im = std::copysign(pi_half, y);
const realT atanh_re = sycl::copysign(realT(0), x);
const realT atanh_im = sycl::copysign(pi_half, y);
return resT{atanh_im, atanh_re};
}
/*
Expand All @@ -96,7 +96,7 @@ template <typename argT, typename resT> struct AtanFunctor
else if (std::isnan(y)) {
/* atanh(+-Inf + I*NaN) = +-0 + I*NaN */
if (std::isinf(x)) {
const realT atanh_re = std::copysign(realT(0), x);
const realT atanh_re = sycl::copysign(realT(0), x);
const realT atanh_im = q_nan;
return resT{atanh_im, atanh_re};
}
Expand All @@ -118,11 +118,11 @@ template <typename argT, typename resT> struct AtanFunctor
*/
constexpr realT r_eps =
realT(1) / std::numeric_limits<realT>::epsilon();
if (std::abs(x) > r_eps || std::abs(y) > r_eps) {
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
const realT pi_half = std::atan(realT(1)) * 2;

const realT atanh_re = realT(0);
const realT atanh_im = std::copysign(pi_half, y);
const realT atanh_im = sycl::copysign(pi_half, y);
return resT{atanh_im, atanh_re};
}
/* ordinary cases */
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -58,7 +58,7 @@ template <typename argT1, typename argT2, typename resT> struct Atan2Functor
{
if (std::isinf(in2) && !std::signbit(in2)) {
if (std::isfinite(in1)) {
return std::copysign(resT(0), in1);
return sycl::copysign(resT(0), in1);
}
}
return std::atan2(in1, in2);
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -78,8 +78,8 @@ template <typename argT, typename resT> struct AtanhFunctor
if (std::isinf(y)) {
const realT pi_half = std::atan(realT(1)) * 2;

const realT res_re = std::copysign(realT(0), x);
const realT res_im = std::copysign(pi_half, y);
const realT res_re = sycl::copysign(realT(0), x);
const realT res_im = sycl::copysign(pi_half, y);
return resT{res_re, res_im};
}
/*
Expand All @@ -90,7 +90,7 @@ template <typename argT, typename resT> struct AtanhFunctor
else if (std::isnan(y)) {
/* atanh(+-Inf + I*NaN) = +-0 + I*NaN */
if (std::isinf(x)) {
const realT res_re = std::copysign(realT(0), x);
const realT res_re = sycl::copysign(realT(0), x);
return resT{res_re, q_nan};
}
/* atanh(+-0 + I*NaN) = +-0 + I*NaN */
Expand All @@ -111,11 +111,12 @@ template <typename argT, typename resT> struct AtanhFunctor
*/
const realT RECIP_EPSILON =
realT(1) / std::numeric_limits<realT>::epsilon();
if (std::abs(x) > RECIP_EPSILON || std::abs(y) > RECIP_EPSILON) {
if (sycl::fabs(x) > RECIP_EPSILON || sycl::fabs(y) > RECIP_EPSILON)
{
const realT pi_half = std::atan(realT(1)) * 2;

const realT res_re = realT(0);
const realT res_im = std::copysign(pi_half, y);
const realT res_im = sycl::copysign(pi_half, y);
return resT{res_re, res_im};
}
/* ordinary cases */
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,84 @@
//===------- cabs_impl.hpp - Implementation of cabs -------*-C++-*/===//
//
// Data Parallel Control (dpctl)
//
// Copyright 2020-2024 Intel Corporation
//
// 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.
//
//===----------------------------------------------------------------------===//
///
/// \file
/// This file defines an implementation of the complex absolute value.
//===----------------------------------------------------------------------===//

#pragma once
#include <cmath>
#include <complex>
#include <limits>

#include "sycl_complex.hpp"

namespace dpctl
{
namespace tensor
{
namespace kernels
{
namespace detail
{

template <typename realT> realT cabs(std::complex<realT> const &z)
{
// Special values for cabs( x + y * 1j):
// * If x is either +infinity or -infinity and y is any value
// (including NaN), the result is +infinity.
// * If x is any value (including NaN) and y is either +infinity or
// -infinity, the result is +infinity.
// * If x is either +0 or -0, the result is equal to abs(y).
// * If y is either +0 or -0, the result is equal to abs(x).
// * If x is NaN and y is a finite number, the result is NaN.
// * If x is a finite number and y is NaN, the result is NaN.
// * If x is NaN and y is NaN, the result is NaN.

const realT x = std::real(z);
const realT y = std::imag(z);

constexpr realT q_nan = std::numeric_limits<realT>::quiet_NaN();
constexpr realT p_inf = std::numeric_limits<realT>::infinity();

if (std::isinf(x)) {
return p_inf;
}
else if (std::isinf(y)) {
return p_inf;
}
else if (std::isnan(x)) {
return q_nan;
}
else if (std::isnan(y)) {
return q_nan;
}
else {
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
return exprm_ns::abs(exprm_ns::complex<realT>(z));
#else
return std::hypot(std::real(z), std::imag(z));
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
return std::hypot(std::real(z), std::imag(z));
return std::hypot(x, y);

#endif
}
}

} // namespace detail
} // namespace kernels
} // namespace tensor
} // namespace dpctl
Original file line number Diff line number Diff line change
Expand Up @@ -71,7 +71,7 @@ template <typename argT, typename resT> struct CeilFunctor
if (in == 0) {
return in;
}
return std::ceil(in);
return sycl::ceil(in);
}
}
};
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -109,7 +109,7 @@ template <typename argT, typename resT> struct CosFunctor
*/
if (x == realT(0) && !yfinite) {
const realT y_m_y = (y - y);
const realT res_im = std::copysign(realT(0), x * y_m_y);
const realT res_im = sycl::copysign(realT(0), x * y_m_y);
return resT{y_m_y, res_im};
}

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

Expand All @@ -144,7 +144,7 @@ template <typename argT, typename resT> struct CosFunctor
*/
if (std::isinf(x)) {
if (!yfinite) {
return resT{x * x, std::copysign(q_nan, x)};
return resT{x * x, sycl::copysign(q_nan, x)};
}
return resT{(x * x) * std::cos(y), x * std::sin(y)};
}
Expand Down
Loading