From c7462bd5240cae659cd52a480bfa8906c4cc5f52 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Fri, 10 Jun 2016 09:46:31 -0500
Subject: [PATCH] Added more device functions

1. Added copyright for device float test
2. Added device double functions support
3. Added device double functions test
4. Corrected device function signatures in headers

Change-Id: I13c8829682c925992f5cad84062bc9f702fe4048
---
 hipamd/include/hcc_detail/hip_runtime.h  |   8 +-
 hipamd/src/device_util.cpp               |  66 +++
 hipamd/tests/src/CMakeLists.txt          |   1 +
 hipamd/tests/src/hipTestDevice.cpp       |  18 +
 hipamd/tests/src/hipTestDeviceDouble.cpp | 529 +++++++++++++++++++++++
 5 files changed, 618 insertions(+), 4 deletions(-)
 create mode 100644 hipamd/tests/src/hipTestDeviceDouble.cpp

diff --git a/hipamd/include/hcc_detail/hip_runtime.h b/hipamd/include/hcc_detail/hip_runtime.h
index eb518cc88d..fbe1ffce15 100644
--- a/hipamd/include/hcc_detail/hip_runtime.h
+++ b/hipamd/include/hcc_detail/hip_runtime.h
@@ -243,7 +243,7 @@ __device__ double jn(int n, double x);
 __device__ double ldexp(double x, int exp);
 __device__ double lgamma(double x);
 __device__ long long llrint(double x);
-__device__ long llround(double x);
+__device__ long long llround(double x);
 __device__ double log(double x);
 __device__ double log10(double x);
 __device__ double log1p(double x);
@@ -257,7 +257,7 @@ __device__ double nearbyint(double x);
 __device__ double nextafter(double x, double y);
 __device__ double norm(int dim, const double* t);
 __device__ double norm3d(double a, double b, double c);
-__device__ double norm4d(double a, double b, double d);
+__device__ double norm4d(double a, double b, double c, double d);
 __device__ double normcdf(double y);
 __device__ double normcdfinv(double y);
 __device__ double pow(double x, double y);
@@ -275,8 +275,8 @@ __device__ double scalbln(double x, long int n);
 __device__ double scalbn(double x, int n);
 __host__ __device__ unsigned signbit(double a);
 __device__ double sin(double a);
-__device__ double sincos(double x, double *sptr, double *cptr);
-__device__ double sincospi(double x, double *sptr, double *cptr);
+__device__ void sincos(double x, double *sptr, double *cptr);
+__device__ void sincospi(double x, double *sptr, double *cptr);
 __device__ double sinh(double x);
 __host__ __device__ double sinpi(double x);
 __device__ double sqrt(double x);
diff --git a/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp
index 55a394f7b0..cd57e54af3 100644
--- a/hipamd/src/device_util.cpp
+++ b/hipamd/src/device_util.cpp
@@ -486,6 +486,16 @@ __device__ double log(double x)
 {
     return hc::precise_math::log(x);
 }
+__device__ long long int llrint(double x)
+{
+    long long int y = hc::precise_math::round(x);
+    return y;
+}
+__device__ long long int llround(double x)
+{
+    long long int y = hc::precise_math::round(x);
+    return y;
+}
 __device__ double log10(double x)
 {
     return hc::precise_math::log10(x);
@@ -502,6 +512,16 @@ __device__ double logb(double x)
 {
     return hc::precise_math::logb(x);
 }
+__device__ long int lrint(double x)
+{
+    long int y = hc::precise_math::round(x);
+    return y;
+}
+__device__ long int lround(double x)
+{
+    long int y = hc::precise_math::round(x);
+    return y;
+}
 __device__ double nan(const char *tagp)
 {
     return hc::precise_math::nan((int)*tagp);
@@ -510,6 +530,17 @@ __device__ double nearbyint(double x)
 {
     return hc::precise_math::nearbyint(x);
 }
+__device__ double norm3d(double a, double b, double c)
+{
+    double x = a*a + b*b + c*c;
+    return hc::precise_math::sqrt(x);
+}
+__device__ double norm4d(double a, double b, double c, double d)
+{
+    double x = a*a + b*b;
+    double y = c*c + d*d;
+    return hc::precise_math::sqrt(x+y);
+}
 __device__ double pow(double x, double y)
 {
     return hc::precise_math::pow(x, y);
@@ -518,6 +549,31 @@ __device__ double remainder(double x, double y)
 {
     return hc::precise_math::remainder(x, y);
 }
+__device__ double rhypot(double x, double y)
+{
+    return 1/hc::precise_math::sqrt(x*x + y*y);
+}
+__device__ double rint(double x)
+{
+    return hc::precise_math::round(x);
+}
+__device__ double rnorm3d(double a, double b, double c)
+{
+    return hc::precise_math::rsqrt(a*a + b*b + c*c);
+}
+__device__ double rnorm4d(double a, double b, double c, double d)
+{
+    return hc::precise_math::rsqrt(a*a + b*b + c*c + d*d);
+}
+__device__ double rnorm(int dim, const double* t)
+{
+    double x = 0.0;
+    for(int i=0;i<dim;i++)
+    {
+        x = hc::precise_math::fma(t[i], t[i], x);
+    }
+    return 1/x;
+}
 __device__ double round(double x)
 {
     return hc::precise_math::round(x);
@@ -538,6 +594,16 @@ __device__ double sin(double x)
 {
     return hc::precise_math::sin(x);
 }
+__device__ void sincos(double x, double *sptr, double *cptr)
+{
+    *sptr = hc::precise_math::sin(x);
+    *cptr = hc::precise_math::cos(x);
+}
+__device__ void sincospi(double x, double *sptr, double *cptr)
+{
+    *sptr = hc::precise_math::sinpi(x);
+    *cptr = hc::precise_math::cospi(x);
+}
 __device__ double sinh(double x)
 {
     return hc::precise_math::sinh(x);
diff --git a/hipamd/tests/src/CMakeLists.txt b/hipamd/tests/src/CMakeLists.txt
index 749b362297..eb2fd603c2 100644
--- a/hipamd/tests/src/CMakeLists.txt
+++ b/hipamd/tests/src/CMakeLists.txt
@@ -199,6 +199,7 @@ make_hip_executable (hipMultiThreadDevice hipMultiThreadDevice.cpp)
 make_hip_executable (hipTestMemcpyPin hipTestMemcpyPin.cpp)
 make_hip_executable (hipDynamicShared hipDynamicShared.cpp)
 make_hip_executable (hipTestDevice hipTestDevice.cpp)
+make_hip_executable (hipTestDeviceDouble hipTestDeviceDouble.cpp)
 make_test(hip_ballot " " )
 make_test(hip_anyall " " )
 make_test(hip_popc " " )
diff --git a/hipamd/tests/src/hipTestDevice.cpp b/hipamd/tests/src/hipTestDevice.cpp
index d36164aed5..4391e559ba 100644
--- a/hipamd/tests/src/hipTestDevice.cpp
+++ b/hipamd/tests/src/hipTestDevice.cpp
@@ -1,3 +1,21 @@
+/*
+Copyright (c) 2015-2016 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.
+*/
 #include"test_common.h"
 #include"hip_runtime.h"
 #include"hip_runtime_api.h"
diff --git a/hipamd/tests/src/hipTestDeviceDouble.cpp b/hipamd/tests/src/hipTestDeviceDouble.cpp
new file mode 100644
index 0000000000..d8890f0596
--- /dev/null
+++ b/hipamd/tests/src/hipTestDeviceDouble.cpp
@@ -0,0 +1,529 @@
+/*
+Copyright (c) 2015-2016 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.
+*/
+#include"test_common.h"
+#include"hip_runtime.h"
+#include"hip_runtime_api.h"
+
+#define N 512
+#define SIZE N*sizeof(double)
+
+__global__ void test_sincos(hipLaunchParm lp, double* a, double* b, double *c){
+    int tid = hipThreadIdx_x;
+    sincos(a[tid], b+tid, c+tid);
+}
+
+__global__ void test_sincospi(hipLaunchParm lp, double* a, double* b, double *c){
+    int tid = hipThreadIdx_x;
+    sincospi(a[tid], b+tid, c+tid);
+}
+
+__global__ void test_llrint(hipLaunchParm lp, double *a, long long int *b){
+    int tid = hipThreadIdx_x;
+    b[tid] = llrint(a[tid]);
+}
+
+__global__ void test_lrint(hipLaunchParm lp, double *a, long int *b){
+    int tid = hipThreadIdx_x;
+    b[tid] = lrint(a[tid]);
+}
+
+__global__ void test_rint(hipLaunchParm lp, double *a, double *b){
+    int tid = hipThreadIdx_x;
+    b[tid] = rint(a[tid]);
+}
+
+__global__ void test_llround(hipLaunchParm lp, double *a, long long int *b){
+    int tid = hipThreadIdx_x;
+    b[tid] = llround(a[tid]);
+}
+
+__global__ void test_lround(hipLaunchParm lp, double *a, long int *b){
+    int tid = hipThreadIdx_x;
+    b[tid] = lround(a[tid]);
+}
+
+__global__ void test_rhypot(hipLaunchParm lp, double *a, double* b, double *c){
+    int tid = hipThreadIdx_x;
+    c[tid] = rhypot(a[tid], b[tid]);
+}
+
+__global__ void test_norm3d(hipLaunchParm lp, double *a, double* b, double *c, double *d){
+    int tid = hipThreadIdx_x;
+    d[tid] = norm3d(a[tid], b[tid], c[tid]);
+}
+
+__global__ void test_norm4d(hipLaunchParm lp, double *a, double* b, double *c, double *d, double *e){
+    int tid = hipThreadIdx_x;
+    e[tid] = norm4d(a[tid], b[tid], c[tid], d[tid]);
+}
+
+__global__ void test_rnorm3d(hipLaunchParm lp, double *a, double* b, double *c, double *d){
+    int tid = hipThreadIdx_x;
+    d[tid] = rnorm3d(a[tid], b[tid], c[tid]);
+}
+
+__global__ void test_rnorm4d(hipLaunchParm lp, double *a, double* b, double *c, double *d, double *e){
+    int tid = hipThreadIdx_x;
+    e[tid] = rnorm4d(a[tid], b[tid], c[tid], d[tid]);
+}
+
+__global__ void test_rnorm(hipLaunchParm lp, double *a, double *b){
+    int tid = hipThreadIdx_x;
+    b[tid] = rnorm(N, a);
+}
+
+bool run_sincos(){
+double *A, *Ad, *B, *C, *Bd, *Cd;
+A = new double[N];
+B = new double[N];
+C = new double[N];
+for(int i=0;i<N;i++){
+A[i] = 1.0;
+}
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, SIZE);
+hipMalloc((void**)&Cd, SIZE);
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_sincos, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
+hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
+hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    if(B[i] == sin(1.0)){
+        passed = 1;
+    }
+}
+passed = 0;
+for(int i=0;i<512;i++){
+    if(C[i] == cos(1.0)){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+bool run_sincospi(){
+double *A, *Ad, *B, *C, *Bd, *Cd;
+A = new double[N];
+B = new double[N];
+C = new double[N];
+for(int i=0;i<N;i++){
+A[i] = 1.0;
+}
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, SIZE);
+hipMalloc((void**)&Cd, SIZE);
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_sincospi, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
+hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
+hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    if(B[i] - sinpi(1.0) < 0.1){
+        passed = 1;
+    }
+}
+passed = 0;
+for(int i=0;i<512;i++){
+    if(C[i] - cospi(1.0) < 0.1){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+
+bool run_llrint(){
+double *A, *Ad;
+long long int *B, *Bd;
+A = new double[N];
+B = new long long int[N];
+for(int i=0;i<N;i++){
+A[i] = 1.345;
+}
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, N*sizeof(long long int));
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_llrint, dim3(1), dim3(N), 0, 0, Ad, Bd);
+hipMemcpy(B, Bd, N*sizeof(long long int), hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    int x = round(A[i]);
+    long long int y = x;
+    if(B[i] == x){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+bool run_lrint(){
+double *A, *Ad;
+long int *B, *Bd;
+A = new double[N];
+B = new long int[N];
+for(int i=0;i<N;i++){
+A[i] = 1.345;
+}
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, N*sizeof(long int));
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_lrint, dim3(1), dim3(N), 0, 0, Ad, Bd);
+hipMemcpy(B, Bd, N*sizeof(long int), hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    long int x = round(A[i]);
+    if(B[i] == x){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+bool run_rint(){
+double *A, *Ad;
+double *B, *Bd;
+A = new double[N];
+B = new double[N];
+for(int i=0;i<N;i++){
+A[i] = 1.345;
+}
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, SIZE);
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_rint, dim3(1), dim3(N), 0, 0, Ad, Bd);
+hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    double x = round(A[i]);
+    if(B[i] == x){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+
+bool run_llround(){
+double *A, *Ad;
+long long int *B, *Bd;
+A = new double[N];
+B = new long long int[N];
+for(int i=0;i<N;i++){
+A[i] = 1.345;
+}
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, N*sizeof(long long int));
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_llround, dim3(1), dim3(N), 0, 0, Ad, Bd);
+hipMemcpy(B, Bd, N*sizeof(long long int), hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    long long int x = round(A[i]);
+    if(B[i] == x){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+bool run_lround(){
+double *A, *Ad;
+long int *B, *Bd;
+A = new double[N];
+B = new long int[N];
+for(int i=0;i<N;i++){
+A[i] = 1.345;
+}
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, N*sizeof(long int));
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_lround, dim3(1), dim3(N), 0, 0, Ad, Bd);
+hipMemcpy(B, Bd, N*sizeof(long int), hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    long int x = round(A[i]);
+    if(B[i] == x){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+
+bool run_norm3d(){
+double *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd;
+A = new double[N];
+B = new double[N];
+C = new double[N];
+D = new double[N];
+double val = 0.0;
+for(int i=0;i<N;i++){
+A[i] = 1.0;
+B[i] = 2.0;
+C[i] = 3.0;
+}
+val = sqrt(1.0 + 4.0 + 9.0);
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, SIZE);
+hipMalloc((void**)&Cd, SIZE);
+hipMalloc((void**)&Dd, SIZE);
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_norm3d, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd);
+hipMemcpy(D, Dd, SIZE, hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    if(D[i] - val < 0.000001){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+bool run_norm4d(){
+double *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd, *E, *Ed;
+A = new double[N];
+B = new double[N];
+C = new double[N];
+D = new double[N];
+E = new double[N];
+double val = 0.0;
+for(int i=0;i<N;i++){
+A[i] = 1.0;
+B[i] = 2.0;
+C[i] = 3.0;
+D[i] = 4.0;
+}
+val = sqrt(1.0 + 4.0 + 9.0 + 16.0);
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, SIZE);
+hipMalloc((void**)&Cd, SIZE);
+hipMalloc((void**)&Dd, SIZE);
+hipMalloc((void**)&Ed, SIZE);
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Dd, D, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_norm4d, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd, Ed);
+hipMemcpy(E, Ed, SIZE, hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    if(E[i] - val < 0.000001){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+
+bool run_rhypot(){
+double *A, *Ad, *B, *Bd, *C, *Cd;
+A = new double[N];
+B = new double[N];
+C = new double[N];
+double val = 0.0;
+for(int i=0;i<N;i++){
+A[i] = 1.0;
+B[i] = 2.0;
+}
+val = 1/sqrt(1.0 + 4.0);
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, SIZE);
+hipMalloc((void**)&Cd, SIZE);
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_rhypot, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
+hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    if(C[i] - val < 0.000001){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+bool run_rnorm3d(){
+double *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd;
+A = new double[N];
+B = new double[N];
+C = new double[N];
+D = new double[N];
+double val = 0.0;
+for(int i=0;i<N;i++){
+A[i] = 1.0;
+B[i] = 2.0;
+C[i] = 3.0;
+}
+val = 1/sqrt(1.0 + 4.0 + 9.0);
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, SIZE);
+hipMalloc((void**)&Cd, SIZE);
+hipMalloc((void**)&Dd, SIZE);
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_rnorm3d, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd);
+hipMemcpy(D, Dd, SIZE, hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    if(D[i] - val < 0.000001){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+bool run_rnorm4d(){
+double *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd, *E, *Ed;
+A = new double[N];
+B = new double[N];
+C = new double[N];
+D = new double[N];
+E = new double[N];
+double val = 0.0;
+for(int i=0;i<N;i++){
+A[i] = 1.0;
+B[i] = 2.0;
+C[i] = 3.0;
+D[i] = 4.0;
+}
+val = 1/sqrt(1.0 + 4.0 + 9.0 + 16.0);
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, SIZE);
+hipMalloc((void**)&Cd, SIZE);
+hipMalloc((void**)&Dd, SIZE);
+hipMalloc((void**)&Ed, SIZE);
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice);
+hipMemcpy(Dd, D, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_rnorm4d, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd, Ed);
+hipMemcpy(E, Ed, SIZE, hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    if(E[i] - val < 0.000001){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+bool run_rnorm(){
+double *A, *Ad, *B, *Bd;
+A = new double[N];
+B = new double[N];
+double val = 0.0;
+for(int i=0;i<N;i++){
+A[i] = 1.0;
+B[i] = 0.0;
+val += 1.0;
+}
+val = 1/sqrt(val);
+hipMalloc((void**)&Ad, SIZE);
+hipMalloc((void**)&Bd, SIZE);
+hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
+hipLaunchKernel(test_rnorm, dim3(1), dim3(N), 0, 0, Ad, Bd);
+hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
+int passed = 0;
+for(int i=0;i<512;i++){
+    if(B[0] - val < 0.000001){
+        passed = 1;
+    }
+}
+free(A);
+if(passed == 1){
+    return true;
+}
+assert(passed == 1);
+return false;
+}
+
+
+int main(){
+if(run_sincos() && run_sincospi() && run_llrint() && run_norm3d() && run_norm4d() &&
+   run_rnorm3d() && run_rnorm4d() &&
+   run_rnorm() && run_lround() && run_llround() &&
+   run_rint() && run_rhypot()
+){
+passed();
+}
+}