diff --git a/tests/src/hipDeviceBody.cpp b/tests/src/hipDeviceBody.cpp
new file mode 100644
index 0000000000..2e740cf748
--- /dev/null
+++ b/tests/src/hipDeviceBody.cpp
@@ -0,0 +1,5 @@
+#include"hipDeviceHeader.h"
+
+__host__ __device__ int add(int &a, int &b){
+    return a+b;
+}
diff --git a/tests/src/hipDeviceHeader.h b/tests/src/hipDeviceHeader.h
new file mode 100644
index 0000000000..eb7bd05aa5
--- /dev/null
+++ b/tests/src/hipDeviceHeader.h
@@ -0,0 +1,7 @@
+#ifndef HIPDEVICEHEADER_H
+#define HIPDEVICEHEADER_H
+#include"test_common.h"
+
+__host__ __device__ int add(int a, int b);
+
+#endif
diff --git a/tests/src/hipDeviceSrc.cpp b/tests/src/hipDeviceSrc.cpp
new file mode 100644
index 0000000000..862d84f06a
--- /dev/null
+++ b/tests/src/hipDeviceSrc.cpp
@@ -0,0 +1,30 @@
+#include"hipDeviceHeader.h"
+
+#define LEN 1024*1024
+#define SIZE LEN * sizeof(int)
+
+__global__ void deviceAdd(hipLaunchParm lp, int *Ad, int *Bd, int *Cd){
+    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    Cd[tx] = add(Ad[tx], Bd[tx]);
+}
+
+int main(){
+    int *A, *B, *C, *Ad, *Bd, *Cd;
+    A = new int[LEN];
+    B = new int[LEN];
+    C = new int[LEN];
+
+    for(int i=0;i<LEN;i++){
+        A[i] = 1;
+        B[i] = 1;
+    }
+
+    HIPCHECK(hipMalloc((void**)&Ad, SIZE));
+    HIPCHECK(hipMalloc((void**)&Bd, SIZE));
+    HIPCHECK(hipMalloc((void**)&Cd, SIZE));
+
+    HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
+    HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
+    hipLaunchKernel(HIP_KERNEL_NAME(deviceAdd), dim3(LEN/1024), dim3(1024), 0, 0, Ad, Bd, Cd);
+
+}