Skip to content

Commit

Permalink
Merge pull request IntelPython#1630 from IntelPython/fix-cuda-linking…
Browse files Browse the repository at this point in the history
…-2024-1

Fix compiling for CUDA with 2024.1 compiler
  • Loading branch information
oleksandr-pavlyk authored Apr 3, 2024
2 parents fda9837 + fc4a612 commit 6abcd34
Show file tree
Hide file tree
Showing 25 changed files with 160 additions and 141 deletions.
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));
#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

0 comments on commit 6abcd34

Please sign in to comment.