HIP: Heterogenous-computing Interface for Portability
Loading...
Searching...
No Matches
amd_hip_complex.h
1/*
2Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
3
4Permission is hereby granted, free of charge, to any person obtaining a copy
5of this software and associated documentation files (the "Software"), to deal
6in the Software without restriction, including without limitation the rights
7to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8copies of the Software, and to permit persons to whom the Software is
9furnished to do so, subject to the following conditions:
10
11The above copyright notice and this permission notice shall be included in
12all copies or substantial portions of the Software.
13
14THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20THE SOFTWARE.
21*/
22
23/* The header defines complex numbers and related functions*/
24
25#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
26#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
27
28#if !defined(__HIPCC_RTC__)
29#include "hip/amd_detail/amd_hip_vector_types.h"
30#endif
31
32#if defined(__HIPCC_RTC__)
33#define __HOST_DEVICE__ __device__
34#else
35#define __HOST_DEVICE__ __host__ __device__
36// TODO: Clang has a bug which allows device functions to call std functions
37// when std functions are introduced into default namespace by using statement.
38// math.h may be included after this bug is fixed.
39#if __cplusplus
40#include <cmath>
41#else
42#include "math.h"
43#endif
44#endif // !defined(__HIPCC_RTC__)
45
46typedef float2 hipFloatComplex;
47
48__HOST_DEVICE__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
49
50__HOST_DEVICE__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
51
52__HOST_DEVICE__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
53 hipFloatComplex z;
54 z.x = a;
55 z.y = b;
56 return z;
57}
58
59__HOST_DEVICE__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
60 hipFloatComplex ret;
61 ret.x = z.x;
62 ret.y = -z.y;
63 return ret;
64}
65
66__HOST_DEVICE__ static inline float hipCsqabsf(hipFloatComplex z) {
67 return z.x * z.x + z.y * z.y;
68}
69
70__HOST_DEVICE__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
71 return make_hipFloatComplex(p.x + q.x, p.y + q.y);
72}
73
74__HOST_DEVICE__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
75 return make_hipFloatComplex(p.x - q.x, p.y - q.y);
76}
77
78__HOST_DEVICE__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
79 return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
80}
81
82__HOST_DEVICE__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
83 float sqabs = hipCsqabsf(q);
84 hipFloatComplex ret;
85 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
86 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
87 return ret;
88}
89
90__HOST_DEVICE__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
91
92
93typedef double2 hipDoubleComplex;
94
95__HOST_DEVICE__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
96
97__HOST_DEVICE__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
98
99__HOST_DEVICE__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
100 hipDoubleComplex z;
101 z.x = a;
102 z.y = b;
103 return z;
104}
105
106__HOST_DEVICE__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
107 hipDoubleComplex ret;
108 ret.x = z.x;
109 ret.y = -z.y;
110 return ret;
111}
112
113__HOST_DEVICE__ static inline double hipCsqabs(hipDoubleComplex z) {
114 return z.x * z.x + z.y * z.y;
115}
116
117__HOST_DEVICE__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
118 return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
119}
120
121__HOST_DEVICE__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
122 return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
123}
124
125__HOST_DEVICE__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
126 return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
127}
128
129__HOST_DEVICE__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
130 double sqabs = hipCsqabs(q);
131 hipDoubleComplex ret;
132 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
133 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
134 return ret;
135}
136
137__HOST_DEVICE__ static inline double hipCabs(hipDoubleComplex z) { return sqrt(hipCsqabs(z)); }
138
139typedef hipFloatComplex hipComplex;
140
141__HOST_DEVICE__ static inline hipComplex make_hipComplex(float x, float y) {
142 return make_hipFloatComplex(x, y);
143}
144
145__HOST_DEVICE__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z) {
146 return make_hipFloatComplex((float)z.x, (float)z.y);
147}
148
149__HOST_DEVICE__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z) {
150 return make_hipDoubleComplex((double)z.x, (double)z.y);
151}
152
153__HOST_DEVICE__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r) {
154 float real = (p.x * q.x) + r.x;
155 float imag = (q.x * p.y) + r.y;
156
157 real = -(p.y * q.y) + real;
158 imag = (p.x * q.y) + imag;
159
160 return make_hipComplex(real, imag);
161}
162
163__HOST_DEVICE__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q,
164 hipDoubleComplex r) {
165 double real = (p.x * q.x) + r.x;
166 double imag = (q.x * p.y) + r.y;
167
168 real = -(p.y * q.y) + real;
169 imag = (p.x * q.y) + imag;
170
171 return make_hipDoubleComplex(real, imag);
172}
173
174#endif //HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
Definition amd_hip_vector_types.h:2035
Definition amd_hip_vector_types.h:2072