Skip to content
This repository was archived by the owner on Mar 21, 2024. It is now read-only.

[WIP] Dirty support for NVRTC on complex headers #1035

Closed
emcastillo wants to merge 3 commits into from
Closed

[WIP] Dirty support for NVRTC on complex headers #1035

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
b0344bd
Dirty support for NVRTC on complex headers
Nov 25, 2019
88868ad
Fix complex.h
Nov 25, 2019
c036848
Fixes
Nov 25, 2019
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
156 changes: 88 additions & 68 deletions thrust/complex.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,12 +21,19 @@

#pragma once

#ifndef __CUDACC_RTC__
// Can't include system dependencies with NVRTC
#include <thrust/detail/config.h>
#include <thrust/detail/type_traits.h>

#include <cmath>
#include <complex>
#include <sstream>
#include <thrust/detail/type_traits.h>
#else
#include "detail/type_traits.h"
#endif



#if THRUST_CPP_DIALECT >= 2011
# define THRUST_STD_COMPLEX_REAL(z) \
Expand Down Expand Up @@ -65,76 +72,80 @@ namespace thrust

namespace detail
{

template <typename T, std::size_t Align>
struct complex_storage;

#ifndef __CUDACC_RTC__
struct complex_storage;
#endif
#if __cplusplus >= 201103L \
&& (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC) \
&& (THRUST_GCC_VERSION >= 40800)
// C++11 implementation, excluding GCC 4.7, which doesn't have `alignas`.
template <typename T, std::size_t Align>
struct complex_storage
{
struct alignas(Align) type { T x; T y; };
};
#elif (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC) \
|| ( (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC) \
&& (THRUST_GCC_VERSION < 40600))
// C++03 implementation for MSVC and GCC <= 4.5.
//
// We have to implement `aligned_type` with specializations for MSVC
// and GCC 4.2 and older because they require literals as arguments to
// their alignment attribute.

#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC)
// MSVC implementation.
&& (THRUST_GCC_VERSION >= 40800) \
&& !defined(__CUDACC_RTC__)
// C++11 implementation, excluding GCC 4.7, which doesn't have `alignas`.
template <typename T, std::size_t Align>
struct complex_storage
{
struct alignas(Align) type { T x; T y; };
};
#elif !defined(__CUDACC_RTC__) \
&& ((THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC) \
|| ((THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC) \
&& (THRUST_GCC_VERSION < 40600)))
// C++03 implementation for MSVC and GCC <= 4.5.
//
// We have to implement `aligned_type` with specializations for MSVC
// and GCC 4.2 and older because they require literals as arguments to
// their alignment attribute.

#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC)
// MSVC implementation.
#define THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(X) \
template <typename T> \
struct complex_storage<T, X> \
{ \
__declspec(align(X)) struct type { T x; T y; }; \
}; \
/**/
#else
// GCC <= 4.2 implementation.
/**/
#else
// GCC <= 4.2 implementation.
#define THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(X) \
template <typename T> \
struct complex_storage<T, X> \
{ \
struct type { T x; T y; } __attribute__((aligned(X))); \
}; \
/**/
#endif

// The primary template is a fallback, which doesn't specify any alignment.
// It's only used when T is very large and we're using an older compilers
// which we have to fully specialize each alignment case.
template <typename T, std::size_t Align>
struct complex_storage
{
T x; T y;
};
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(1);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(2);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(4);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(8);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(16);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(32);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(64);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(128);

#undef THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION
/**/
#endif

// The primary template is a fallback, which doesn't specify any alignment.
// It's only used when T is very large and we're using an older compilers
// which we have to fully specialize each alignment case.
template <typename T, std::size_t Align>
struct complex_storage
{
T x; T y;
};

THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(1);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(2);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(4);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(8);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(16);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(32);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(64);
THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION(128);

#undef THRUST_DEFINE_COMPLEX_STORAGE_SPECIALIZATION
#else
// C++03 implementation for GCC > 4.5, Clang, PGI, ICPC, and xlC.
template <typename T, std::size_t Align>
struct complex_storage
{
struct type { T x; T y; } __attribute__((aligned(Align)));
};

template <typename T, ::size_t Align>
struct complex_storage;
// C++03 implementation for GCC > 4.5, Clang, PGI, ICPC, and xlC.
template <typename T, ::size_t Align>
struct complex_storage
{
struct type { T x; T y; } __attribute__((aligned(Align)));
};
#endif

} // end namespace detail

/*! \p complex is the Thrust equivalent to <tt>std::complex</tt>. It is
Expand Down Expand Up @@ -209,7 +220,7 @@ struct complex
template <typename U>
__host__ __device__
complex(const complex<U>& z);

#ifndef __CUDACC_RTC__
/*! This converting copy constructor copies from a <tt>std::complex</tt> with
* a type that is convertible to this \p complex's \c value_type.
*
Expand All @@ -228,7 +239,7 @@ struct complex
template <typename U>
__host__ THRUST_STD_COMPLEX_DEVICE
complex(const std::complex<U>& z);

#endif


/* --- Assignment Operators --- */
Expand Down Expand Up @@ -268,7 +279,7 @@ struct complex
template <typename U>
__host__ __device__
complex& operator=(const complex<U>& z);

#ifndef __CUDACC_RTC__
/*! Assign `z.real()` and `z.imag()` to the real and imaginary parts of this
* \p complex respectively.
*
Expand All @@ -287,7 +298,7 @@ struct complex
template <typename U>
__host__ THRUST_STD_COMPLEX_DEVICE
complex& operator=(const std::complex<U>& z);

#endif

/* --- Compound Assignment Operators --- */

Expand Down Expand Up @@ -444,12 +455,12 @@ struct complex


/* --- Casting functions --- */

#ifndef __CUDACC_RTC__
/*! Casts this \p complex to a <tt>std::complex</tt> of the same type.
*/
__host__
operator std::complex<T>() const { return std::complex<T>(real(), imag()); }

#endif
private:
typename detail::complex_storage<T, sizeof(T) * 2>::type data;
};
Expand Down Expand Up @@ -489,6 +500,16 @@ template <typename T>
__host__ __device__
complex<T> conj(const complex<T>& z);

template <typename T>
__host__ __device__ inline T real(const complex<T>& z) {
return z.real();
}

template <typename T>
__host__ __device__ inline T imag(const complex<T>& z) {
return z.imag();
}

/*! Returns a \p complex with the specified magnitude and phase.
*
* \param m The magnitude of the returned \p complex.
Expand Down Expand Up @@ -899,7 +920,7 @@ __host__ __device__
complex<T> atanh(const complex<T>& z);



#ifndef __CUDACC_RTC__
/* --- Stream Operators --- */

/*! Writes to an output stream a \p complex number in the form (real, imaginary).
Expand Down Expand Up @@ -927,8 +948,7 @@ template <typename T, typename CharT, typename Traits>
__host__
std::basic_istream<CharT, Traits>&
operator>>(std::basic_istream<CharT, Traits>& is, complex<T>& z);


#endif

/* --- Equality Operators --- */

Expand All @@ -940,7 +960,7 @@ operator>>(std::basic_istream<CharT, Traits>& is, complex<T>& z);
template <typename T0, typename T1>
__host__ __device__
bool operator==(const complex<T0>& x, const complex<T1>& y);

#ifndef __CUDACC_RTC__
/*! Returns true if two \p complex numbers are equal and false otherwise.
*
* \param x The first \p complex.
Expand All @@ -949,7 +969,6 @@ bool operator==(const complex<T0>& x, const complex<T1>& y);
template <typename T0, typename T1>
__host__ THRUST_STD_COMPLEX_DEVICE
bool operator==(const complex<T0>& x, const std::complex<T1>& y);

/*! Returns true if two \p complex numbers are equal and false otherwise.
*
* \param x The first \p complex.
Expand All @@ -959,6 +978,7 @@ template <typename T0, typename T1>
__host__ THRUST_STD_COMPLEX_DEVICE
bool operator==(const std::complex<T0>& x, const complex<T1>& y);

#endif
/*! Returns true if the imaginary part of the \p complex number is zero and
* the real part is equal to the scalar. Returns false otherwise.
*
Expand Down Expand Up @@ -987,7 +1007,7 @@ bool operator==(const complex<T0>& x, const T1& y);
template <typename T0, typename T1>
__host__ __device__
bool operator!=(const complex<T0>& x, const complex<T1>& y);

#ifndef __CUDACC_RTC__
/*! Returns true if two \p complex numbers are different and false otherwise.
*
* \param x The first \p complex.
Expand All @@ -1005,7 +1025,7 @@ bool operator!=(const complex<T0>& x, const std::complex<T1>& y);
template <typename T0, typename T1>
__host__ THRUST_STD_COMPLEX_DEVICE
bool operator!=(const std::complex<T0>& x, const complex<T1>& y);

#endif
/*! Returns true if the imaginary part of the \p complex number is not zero or
* the real part is different from the scalar. Returns false otherwise.
*
Expand All @@ -1028,7 +1048,7 @@ bool operator!=(const complex<T0>& x, const T1& y);

} // end namespace thrust

#include <thrust/detail/complex/complex.inl>
#include "detail/complex/complex.inl"

#undef THRUST_STD_COMPLEX_REAL
#undef THRUST_STD_COMPLEX_IMAG
Expand Down
25 changes: 13 additions & 12 deletions thrust/detail/alignment.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,10 +20,11 @@

#pragma once

#ifndef __CUDACC_RTC__
#include <thrust/detail/config.h>
#include <thrust/detail/type_traits.h> // For `integral_constant`.

#include <cstddef> // For `std::size_t` and `std::max_align_t`.
#endif

#if __cplusplus >= 201103L
#include <type_traits> // For `std::alignment_of` and `std::aligned_storage`.
Expand Down Expand Up @@ -61,17 +62,17 @@ namespace detail
template <typename T>
struct alignment_of;

template <typename T, std::size_t size_diff>
template <typename T, ::size_t size_diff>
struct alignment_of_helper
{
static const std::size_t value =
integral_constant<std::size_t, size_diff>::value;
static const ::size_t value =
integral_constant<::size_t, size_diff>::value;
};

template <typename T>
struct alignment_of_helper<T, 0>
{
static const std::size_t value = alignment_of<T>::value;
static const ::size_t value = alignment_of<T>::value;
};

template <typename T>
Expand All @@ -85,7 +86,7 @@ namespace detail
};

public:
static const std::size_t value =
static const ::size_t value =
alignment_of_helper<impl, sizeof(impl) - sizeof(T)>::value;
};
#endif
Expand All @@ -94,14 +95,14 @@ namespace detail
/// type whose alignment requirement is a divisor of `Align`.
///
/// The behavior is undefined if `Align` is not a power of 2.
template <std::size_t Align>
template <::size_t Align>
struct aligned_type;

#if __cplusplus >= 201103L \
&& (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC) \
&& (THRUST_GCC_VERSION >= 40800)
// C++11 implementation, excluding GCC 4.7, which doesn't have `alignas`.
template <std::size_t Align>
template <::size_t Align>
struct aligned_type
{
struct alignas(Align) type {};
Expand Down Expand Up @@ -147,7 +148,7 @@ struct aligned_type;
#undef THRUST_DEFINE_ALIGNED_TYPE_SPECIALIZATION
#else
// C++03 implementation for GCC > 4.5, Clang, PGI, ICPC, and xlC.
template <std::size_t Align>
template <::size_t Align>
struct aligned_type
{
struct type {} __attribute__((aligned(Align)));
Expand All @@ -162,10 +163,10 @@ struct aligned_type;
///
/// It is an implementation of C++11's \p std::aligned_storage.
#if __cplusplus >= 201103L
template <std::size_t Len, std::size_t Align>
template <::size_t Len, ::size_t Align>
using aligned_storage = std::aligned_storage<Len, Align>;
#else
template <std::size_t Len, std::size_t Align>
template <::size_t Len, ::size_t Align>
struct aligned_storage
{
union type
Expand Down Expand Up @@ -220,7 +221,7 @@ T aligned_reinterpret_cast(U u)
}

__host__ __device__
inline std::size_t aligned_storage_size(std::size_t n, std::size_t align)
inline ::size_t aligned_storage_size(::size_t n, ::size_t align)
{
return ((n + align - 1) / align) * align;
}
Expand Down
Loading