// Copyright (C) 2021 - 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal // in the Software without restriction, including without limitation the rights // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell // copies of the Software, and to permit persons to whom the Software is // furnished to do so, subject to the following conditions: // // The above copyright notice and this permission notice shall be included in // all copies or substantial portions of the Software. // // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN // THE SOFTWARE. #ifndef ROCFFT_COMPLEX_H #define ROCFFT_COMPLEX_H #include #if !defined(__HIPCC_RTC__) #include #endif #include #include #if defined(__HIP_PLATFORM_NVIDIA__) typedef __half rocfft_fp16; #elif defined(__HIP_PLATFORM_AMD__) typedef _Float16 rocfft_fp16; #else // HIP AOT compilation or HIP RTC would set one of the two macros above // so if we're here, we must be doing NVRTC #define ROCFFT_COMPLEX_NVRTC #endif template struct rocfft_complex { Treal x; // Real part Treal y; // Imaginary part // Constructors // Do not initialize the members x or y by default, to ensure that it can // be used in __shared__ and that it is a trivial class compatible with C. rocfft_complex() = default; rocfft_complex(const rocfft_complex&) = default; rocfft_complex(rocfft_complex&&) = default; rocfft_complex& operator=(const rocfft_complex& rhs) & = default; rocfft_complex& operator=(rocfft_complex&& rhs) & = default; ~rocfft_complex() = default; // Constructor from real and imaginary parts __device__ __host__ constexpr rocfft_complex(Treal real, Treal imag) : x{real} , y{imag} { } // Conversion from different precision template __device__ __host__ explicit constexpr rocfft_complex(const rocfft_complex& z) : x(z.x) , y(z.y) { } // Accessors __device__ __host__ constexpr Treal real() const { return x; } __device__ __host__ constexpr Treal imag() const { return y; } // Mutators __device__ __host__ void real(const Treal new_x) { x = new_x; } __device__ __host__ void imag(const Treal new_y) { y = new_y; } // Unary operations __forceinline__ __device__ __host__ rocfft_complex operator-() const { return {-x, -y}; } __forceinline__ __device__ __host__ rocfft_complex operator+() const { return *this; } __device__ __host__ Treal asum(const rocfft_complex& z) { return abs(z.x) + abs(z.y); } // Internal real functions static __forceinline__ __device__ __host__ Treal abs(Treal x) { return x < 0 ? -x : x; } static __forceinline__ __device__ __host__ float sqrt(float x) { return ::sqrtf(x); } static __forceinline__ __device__ __host__ double sqrt(double x) { return ::sqrt(x); } // Addition operators __device__ __host__ auto& operator+=(const rocfft_complex& rhs) { return *this = {x + rhs.x, y + rhs.y}; } __device__ __host__ auto operator+(const rocfft_complex& rhs) const { auto lhs = *this; return lhs += rhs; } // Subtraction operators __device__ __host__ auto& operator-=(const rocfft_complex& rhs) { return *this = {x - rhs.x, y - rhs.y}; } __device__ __host__ auto operator-(const rocfft_complex& rhs) const { auto lhs = *this; return lhs -= rhs; } // Multiplication operators __device__ __host__ auto& operator*=(const rocfft_complex& rhs) { return *this = {x * rhs.x - y * rhs.y, y * rhs.x + x * rhs.y}; } __device__ __host__ auto operator*(const rocfft_complex& rhs) const { auto lhs = *this; return lhs *= rhs; } // Division operators __device__ __host__ auto& operator/=(const rocfft_complex& rhs) { // Form of Robert L. Smith's Algorithm 116 if(abs(rhs.x) > abs(rhs.y)) { Treal ratio = rhs.y / rhs.x; Treal scale = 1 / (rhs.x + rhs.y * ratio); *this = {(x + y * ratio) * scale, (y - x * ratio) * scale}; } else { Treal ratio = rhs.x / rhs.y; Treal scale = 1 / (rhs.x * ratio + rhs.y); *this = {(y + x * ratio) * scale, (y * ratio - x) * scale}; } return *this; } __device__ __host__ auto operator/(const rocfft_complex& rhs) const { auto lhs = *this; return lhs /= rhs; } // Comparison operators __device__ __host__ constexpr bool operator==(const rocfft_complex& rhs) const { return x == rhs.x && y == rhs.y; } __device__ __host__ constexpr bool operator!=(const rocfft_complex& rhs) const { return !(*this == rhs); } // Operators for complex-real computations template __device__ __host__ auto& operator+=(const U& rhs) { return (x += Treal(rhs)), *this; } template __device__ __host__ auto& operator-=(const U& rhs) { return (x -= Treal(rhs)), *this; } __device__ __host__ auto operator+(const Treal& rhs) const { auto lhs = *this; return lhs += rhs; } __device__ __host__ auto operator-(const Treal& rhs) const { auto lhs = *this; return lhs -= rhs; } template __device__ __host__ auto& operator*=(const U& rhs) { return (x *= Treal(rhs)), (y *= Treal(rhs)), *this; } template __device__ __host__ auto operator*(const U& rhs) const { auto lhs = *this; return lhs *= Treal(rhs); } template __device__ __host__ auto& operator/=(const U& rhs) { return (x /= Treal(rhs)), (y /= Treal(rhs)), *this; } template __device__ __host__ auto operator/(const U& rhs) const { auto lhs = *this; return lhs /= Treal(rhs); } template __device__ __host__ constexpr bool operator==(const U& rhs) const { return x == Treal(rhs) && y == 0; } template __device__ __host__ constexpr bool operator!=(const U& rhs) const { return !(*this == rhs); } }; // Stream operators #if !defined(__HIPCC_RTC__) && !defined(ROCFFT_COMPLEX_NVRTC) static std::ostream& operator<<(std::ostream& stream, const rocfft_fp16& f) { return stream << static_cast(f); } template std::ostream& operator<<(std::ostream& out, const rocfft_complex& z) { return out << '(' << static_cast(z.x) << ',' << static_cast(z.y) << ')'; } #endif // Operators for real-complex computations template __device__ __host__ rocfft_complex operator+(const U& lhs, const rocfft_complex& rhs) { return {Treal(lhs) + rhs.x, rhs.y}; } template __device__ __host__ rocfft_complex operator-(const U& lhs, const rocfft_complex& rhs) { return {Treal(lhs) - rhs.x, -rhs.y}; } template __device__ __host__ rocfft_complex operator*(const U& lhs, const rocfft_complex& rhs) { return {Treal(lhs) * rhs.x, Treal(lhs) * rhs.y}; } template __device__ __host__ rocfft_complex operator/(const U& lhs, const rocfft_complex& rhs) { // Form of Robert L. Smith's Algorithm 116 if(rocfft_complex::abs(rhs.x) > rocfft_complex::abs(rhs.y)) { Treal ratio = rhs.y / rhs.x; Treal scale = Treal(lhs) / (rhs.x + rhs.y * ratio); return {scale, -scale * ratio}; } else { Treal ratio = rhs.x / rhs.y; Treal scale = Treal(lhs) / (rhs.x * ratio + rhs.y); return {ratio * scale, -scale}; } } template __device__ __host__ constexpr bool operator==(const U& lhs, const rocfft_complex& rhs) { return Treal(lhs) == rhs.x && 0 == rhs.y; } template __device__ __host__ constexpr bool operator!=(const U& lhs, const rocfft_complex& rhs) { return !(lhs == rhs); } // Extending std namespace to handle rocfft_complex datatype namespace std { template __device__ __host__ constexpr Treal real(const rocfft_complex& z) { return z.x; } // NVRTC does not provide an fp16 type #ifndef ROCFFT_COMPLEX_NVRTC __device__ __host__ constexpr rocfft_fp16 real(const rocfft_fp16 x) { return x; } #endif template __device__ __host__ constexpr Treal imag(const rocfft_complex& z) { return z.y; } template __device__ __host__ constexpr rocfft_complex conj(const rocfft_complex& z) { return {z.x, -z.y}; } template __device__ __host__ inline Treal norm(const rocfft_complex& z) { return (z.x * z.x) + (z.y * z.y); } template __device__ __host__ inline Treal abs(const rocfft_complex& z) { Treal tr = rocfft_complex::abs(z.x), ti = rocfft_complex::abs(z.y); return tr > ti ? (ti /= tr, tr * rocfft_complex::sqrt(ti * ti + 1)) : ti ? (tr /= ti, ti * rocfft_complex::sqrt(tr * tr + 1)) : 0; } } #endif // ROCFFT_COMPLEX_H