From 66d76146d2e697c81cf399d81ea6c04332635753 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Fri, 19 Aug 2016 21:48:23 -0500
Subject: [PATCH] Added support for complex device functions - Added complex
 number arithmetic operation for float and double datatypes - TODO: make them
 host functions and support half - Added new function which is not in CUDA,
 hipCsqabs which is square of absolute value

Change-Id: Ib96e194ad45dc64fcba29eb19ad0376542e0591d


[ROCm/hip commit: 78b15bf062073bfc8e152cadc6cf267e4071608a]
---
 projects/hip/include/hcc_detail/hipComplex.h | 116 +++++++++++++++++++
 projects/hip/tests/src/hipComplex.cpp        |  33 ++++++
 2 files changed, 149 insertions(+)
 create mode 100644 projects/hip/include/hcc_detail/hipComplex.h
 create mode 100644 projects/hip/tests/src/hipComplex.cpp

diff --git a/projects/hip/include/hcc_detail/hipComplex.h b/projects/hip/include/hcc_detail/hipComplex.h
new file mode 100644
index 0000000000..c5584047bb
--- /dev/null
+++ b/projects/hip/include/hcc_detail/hipComplex.h
@@ -0,0 +1,116 @@
+#ifndef HIPCOMPLEX_H
+#define HIPCOMPLEX_H
+
+
+typedef struct{
+    float x;
+    float y;
+}hipFloatComplex;
+
+__device__ inline float hipCrealf(hipFloatComplex z){
+    return z.x;
+}
+
+__device__ inline float hipCimagf(hipFloatComplex z){
+    return z.y;
+}
+
+__device__ inline hipFloatComplex make_hipFloatComplex(float a, float b){
+    hipFloatComplex z;
+    z.x = a;
+    z.y = b;
+    return z;
+}
+
+__device__ inline hipFloatComplex hipConjf(hipFloatComplex z){
+    hipFloatComplex ret;
+    ret.x = z.x;
+    ret.y = z.y;
+    return ret;
+}
+
+__device__ inline float hipCsqabsf(hipFloatComplex z){
+    return z.x * z.x + z.y * z.y;
+}
+
+__device__ hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q){
+    return make_hipFloatComplex(p.x + q.x, p.y + q.y);
+}
+
+__device__ hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q){
+    return make_hipFloatComplex(p.x - q.x, p.y - q.y);
+}
+
+__device__ hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q){
+    return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
+}
+
+__device__ hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q){
+    float sqabs = hipCsqabsf(q);
+    hipFloatComplex ret;
+    ret.x = (p.x * q.x + p.y * q.y)/sqabs;
+    ret.y = (p.y * q.x - p.x * q.y)/sqabs;
+    return ret;
+}
+
+__device__ float hipCabsf(hipFloatComplex z){
+    return sqrtf(hipCsqabsf(z));
+}
+
+
+typedef struct{
+    double x;
+    double y;
+}hipDoubleComplex;
+
+__device__ inline double hipCreal(hipDoubleComplex z){
+    return z.x;
+}
+
+__device__ inline double hipCimag(hipDoubleComplex z){
+    return z.y;
+}
+
+__device__ inline hipDoubleComplex make_hipDoubleComplex(double a, double b){
+    hipDoubleComplex z;
+    z.x = a;
+    z.y = b;
+    return z;
+}
+
+__device__ inline hipDoubleComplex hipConj(hipDoubleComplex z){
+    hipDoubleComplex ret;
+    ret.x = z.x;
+    ret.y = z.y;
+    return ret;
+}
+
+__device__ inline double hipCsqabs(hipDoubleComplex z){
+    return z.x * z.x + z.y * z.y;
+}
+
+__device__ hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q){
+    return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
+}
+
+__device__ hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q){
+    return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
+}
+
+__device__ hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q){
+    return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
+}
+
+__device__ hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q){
+    double sqabs = hipCsqabs(q);
+    hipDoubleComplex ret;
+    ret.x = (p.x * q.x + p.y * q.y)/sqabs;
+    ret.y = (p.y * q.x - p.x * q.y)/sqabs;
+    return ret;
+}
+
+__device__ inline double hipCabs(hipDoubleComplex z){
+    return sqrtf(hipCsqabs(z));
+}
+
+#endif
diff --git a/projects/hip/tests/src/hipComplex.cpp b/projects/hip/tests/src/hipComplex.cpp
new file mode 100644
index 0000000000..bfa1a0be15
--- /dev/null
+++ b/projects/hip/tests/src/hipComplex.cpp
@@ -0,0 +1,33 @@
+#include<iostream>
+#include<hip/hip_runtime.h>
+#include<hip/hip_runtime_api.h>
+#include<hip/hcc_detail/hipComplex.h>
+
+#define LEN 64
+#define SIZE 64<<2
+
+__global__  void getSqAbs(hipLaunchParm lp, float *A, float *B, float *C){
+    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    C[tx] = hipCsqabsf(make_hipFloatComplex(A[tx], B[tx]));
+}
+
+int main(){
+    float *A, *Ad, *B, *Bd, *C, *Cd;
+    A = new float[LEN];
+    B = new float[LEN];
+    C = new float[LEN];
+    for(uint32_t i=0;i<LEN;i++){
+        A[i] = i*1.0f;
+        B[i] = i*1.0f;
+        C[i] = i*1.0f;
+    }
+
+    hipMalloc((void**)&Ad, SIZE);
+    hipMalloc((void**)&Bd, SIZE);
+    hipMalloc((void**)&Cd, SIZE);
+    hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+    hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
+    hipLaunchKernel(getSqAbs, dim3(1), dim3(LEN), 0, 0, Ad, Bd, Cd);
+    hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost);
+    std::cout<<A[11]<<" "<<B[11]<<" "<<C[11]<<std::endl;
+}