23 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H 24 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H 38 #define COMPLEX_NEG_OP_OVERLOAD(type) \ 39 __device__ __host__ static inline type operator-(const type& op) { \ 46 #define COMPLEX_EQ_OP_OVERLOAD(type) \ 47 __device__ __host__ static inline bool operator==(const type& lhs, const type& rhs) { \ 48 return lhs.x == rhs.x && lhs.y == rhs.y; \ 51 #define COMPLEX_NE_OP_OVERLOAD(type) \ 52 __device__ __host__ static inline bool operator!=(const type& lhs, const type& rhs) { \ 53 return !(lhs == rhs); \ 56 #define COMPLEX_ADD_OP_OVERLOAD(type) \ 57 __device__ __host__ static inline type operator+(const type& lhs, const type& rhs) { \ 59 ret.x = lhs.x + rhs.x; \ 60 ret.y = lhs.y + rhs.y; \ 64 #define COMPLEX_SUB_OP_OVERLOAD(type) \ 65 __device__ __host__ static inline type operator-(const type& lhs, const type& rhs) { \ 67 ret.x = lhs.x - rhs.x; \ 68 ret.y = lhs.y - rhs.y; \ 72 #define COMPLEX_MUL_OP_OVERLOAD(type) \ 73 __device__ __host__ static inline type operator*(const type& lhs, const type& rhs) { \ 75 ret.x = lhs.x * rhs.x - lhs.y * rhs.y; \ 76 ret.y = lhs.x * rhs.y + lhs.y * rhs.x; \ 80 #define COMPLEX_DIV_OP_OVERLOAD(type) \ 81 __device__ __host__ static inline type operator/(const type& lhs, const type& rhs) { \ 83 ret.x = (lhs.x * rhs.x + lhs.y * rhs.y); \ 84 ret.y = (rhs.x * lhs.y - lhs.x * rhs.y); \ 85 ret.x = ret.x / (rhs.x * rhs.x + rhs.y * rhs.y); \ 86 ret.y = ret.y / (rhs.x * rhs.x + rhs.y * rhs.y); \ 90 #define COMPLEX_ADD_PREOP_OVERLOAD(type) \ 91 __device__ __host__ static inline type& operator+=(type& lhs, const type& rhs) { \ 97 #define COMPLEX_SUB_PREOP_OVERLOAD(type) \ 98 __device__ __host__ static inline type& operator-=(type& lhs, const type& rhs) { \ 104 #define COMPLEX_MUL_PREOP_OVERLOAD(type) \ 105 __device__ __host__ static inline type& operator*=(type& lhs, const type& rhs) { \ 110 #define COMPLEX_DIV_PREOP_OVERLOAD(type) \ 111 __device__ __host__ static inline type& operator/=(type& lhs, const type& rhs) { \ 116 #define COMPLEX_SCALAR_PRODUCT(type, type1) \ 117 __device__ __host__ static inline type operator*(const type& lhs, type1 rhs) { \ 119 ret.x = lhs.x * rhs; \ 120 ret.y = lhs.y * rhs; \ 126 typedef float2 hipFloatComplex;
128 __device__
__host__ static inline float hipCrealf(hipFloatComplex z) {
return z.x; }
130 __device__
__host__ static inline float hipCimagf(hipFloatComplex z) {
return z.y; }
132 __device__
__host__ static inline hipFloatComplex make_hipFloatComplex(
float a,
float b) {
139 __device__
__host__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
146 __device__
__host__ static inline float hipCsqabsf(hipFloatComplex z) {
147 return z.x * z.x + z.y * z.y;
150 __device__
__host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
151 return make_hipFloatComplex(p.x + q.x, p.y + q.y);
154 __device__
__host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
155 return make_hipFloatComplex(p.x - q.x, p.y - q.y);
158 __device__
__host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
159 return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
162 __device__
__host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
163 float sqabs = hipCsqabsf(q);
165 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
166 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
170 __device__
__host__ static inline float hipCabsf(hipFloatComplex z) {
return sqrtf(hipCsqabsf(z)); }
173 typedef double2 hipDoubleComplex;
175 __device__
__host__ static inline double hipCreal(hipDoubleComplex z) {
return z.x; }
177 __device__
__host__ static inline double hipCimag(hipDoubleComplex z) {
return z.y; }
179 __device__
__host__ static inline hipDoubleComplex make_hipDoubleComplex(
double a,
double b) {
186 __device__
__host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
187 hipDoubleComplex ret;
193 __device__
__host__ static inline double hipCsqabs(hipDoubleComplex z) {
194 return z.x * z.x + z.y * z.y;
197 __device__
__host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
198 return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
201 __device__
__host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
202 return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
205 __device__
__host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
206 return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
209 __device__
__host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
210 double sqabs = hipCsqabs(q);
211 hipDoubleComplex ret;
212 ret.x = (p.x * q.x + p.y * q.y) / sqabs;
213 ret.y = (p.y * q.x - p.x * q.y) / sqabs;
217 __device__
__host__ static inline double hipCabs(hipDoubleComplex z) {
return sqrtf(hipCsqabs(z)); }
222 COMPLEX_NEG_OP_OVERLOAD(hipFloatComplex)
223 COMPLEX_EQ_OP_OVERLOAD(hipFloatComplex)
224 COMPLEX_NE_OP_OVERLOAD(hipFloatComplex)
225 COMPLEX_ADD_OP_OVERLOAD(hipFloatComplex)
226 COMPLEX_SUB_OP_OVERLOAD(hipFloatComplex)
227 COMPLEX_MUL_OP_OVERLOAD(hipFloatComplex)
228 COMPLEX_DIV_OP_OVERLOAD(hipFloatComplex)
229 COMPLEX_ADD_PREOP_OVERLOAD(hipFloatComplex)
230 COMPLEX_SUB_PREOP_OVERLOAD(hipFloatComplex)
231 COMPLEX_MUL_PREOP_OVERLOAD(hipFloatComplex)
232 COMPLEX_DIV_PREOP_OVERLOAD(hipFloatComplex)
233 COMPLEX_SCALAR_PRODUCT(hipFloatComplex,
unsigned short)
234 COMPLEX_SCALAR_PRODUCT(hipFloatComplex,
signed short)
235 COMPLEX_SCALAR_PRODUCT(hipFloatComplex,
unsigned int)
236 COMPLEX_SCALAR_PRODUCT(hipFloatComplex,
signed int)
237 COMPLEX_SCALAR_PRODUCT(hipFloatComplex,
float)
238 COMPLEX_SCALAR_PRODUCT(hipFloatComplex,
unsigned long)
239 COMPLEX_SCALAR_PRODUCT(hipFloatComplex,
signed long)
240 COMPLEX_SCALAR_PRODUCT(hipFloatComplex,
double)
241 COMPLEX_SCALAR_PRODUCT(hipFloatComplex,
signed long long)
242 COMPLEX_SCALAR_PRODUCT(hipFloatComplex,
unsigned long long)
244 COMPLEX_NEG_OP_OVERLOAD(hipDoubleComplex)
245 COMPLEX_EQ_OP_OVERLOAD(hipDoubleComplex)
246 COMPLEX_NE_OP_OVERLOAD(hipDoubleComplex)
247 COMPLEX_ADD_OP_OVERLOAD(hipDoubleComplex)
248 COMPLEX_SUB_OP_OVERLOAD(hipDoubleComplex)
249 COMPLEX_MUL_OP_OVERLOAD(hipDoubleComplex)
250 COMPLEX_DIV_OP_OVERLOAD(hipDoubleComplex)
251 COMPLEX_ADD_PREOP_OVERLOAD(hipDoubleComplex)
252 COMPLEX_SUB_PREOP_OVERLOAD(hipDoubleComplex)
253 COMPLEX_MUL_PREOP_OVERLOAD(hipDoubleComplex)
254 COMPLEX_DIV_PREOP_OVERLOAD(hipDoubleComplex)
255 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex,
unsigned short)
256 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex,
signed short)
257 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex,
unsigned int)
258 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex,
signed int)
259 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex,
float)
260 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex,
unsigned long)
261 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex,
signed long)
262 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex,
double)
263 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex,
signed long long)
264 COMPLEX_SCALAR_PRODUCT(hipDoubleComplex,
unsigned long long)
269 typedef hipFloatComplex hipComplex;
271 __device__
__host__ static inline hipComplex make_hipComplex(
float x,
float y) {
272 return make_hipFloatComplex(x, y);
275 __device__
__host__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z) {
276 return make_hipFloatComplex((
float)z.x, (
float)z.y);
279 __device__
__host__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z) {
280 return make_hipDoubleComplex((
double)z.x, (
double)z.y);
283 __device__
__host__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r) {
284 float real = (p.x * q.x) + r.x;
285 float imag = (q.x * p.y) + r.y;
287 real = -(p.y * q.y) + real;
288 imag = (p.x * q.y) + imag;
290 return make_hipComplex(real, imag);
293 __device__
__host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q,
294 hipDoubleComplex r) {
295 double real = (p.x * q.x) + r.x;
296 double imag = (q.x * p.y) + r.y;
298 real = -(p.y * q.y) + real;
299 imag = (p.x * q.y) + imag;
301 return make_hipDoubleComplex(real, imag);
304 #endif //HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H #define __host__
Definition: host_defines.h:41
Defines the different newt vector types for HIP runtime.