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

@@ -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);

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);

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 " " )

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"

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();
+}
+}