diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index 2f11a2642f..0a5e2675dc 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -24,7 +24,8 @@ set(CMAKE_C_COMPILER   "${HCC_DIR}/bin/hcc")
 set(CMAKE_CXX_FLAGS " -hc -I${HCC_DIR}/include -I${HSA_DIR}/include -stdlib=libc++ ")
 set(CMAKE_C_FLAGS   " -hc -I${HCC_DIR}/include -I${HSA_DIR}/include -stdlib=libc++ ")
 
-set(SOURCE_FILES src/hip_hcc.cpp 
+set(SOURCE_FILES src/device_util.cpp
+src/hip_hcc.cpp 
 src/hip_device.cpp
 src/hip_error.cpp
 src/hip_event.cpp
diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h
index 01a278ecf8..7c34754355 100644
--- a/projects/hip/include/hcc_detail/hip_runtime.h
+++ b/projects/hip/include/hcc_detail/hip_runtime.h
@@ -273,9 +273,11 @@ __device__ inline unsigned long long int atomicXor(unsigned long long int* addre
 // integer intrinsic function __poc __clz __ffs __brev
 __device__ inline unsigned int __popc( unsigned int input)
 {
-	return hc::__popcount_u32_b32( input);
+    return hc::__popcount_u32_b32(input);
 }
 
+__device__ unsigned int test__popc(unsigned int input);
+
 __device__ inline unsigned int __popcll( unsigned long long int input)
 {
 	return hc::__popcount_u32_b64(input);
diff --git a/projects/hip/src/device_util.cpp b/projects/hip/src/device_util.cpp
new file mode 100644
index 0000000000..ef1ed0d6ed
--- /dev/null
+++ b/projects/hip/src/device_util.cpp
@@ -0,0 +1,9 @@
+#include"hip_runtime.h"
+#include<hc.hpp>
+#include<grid_launch.h>
+
+__device__ unsigned int test__popc(unsigned int input)
+{
+    return hc::__popcount_u32_b32(input);
+}
+
diff --git a/projects/hip/tests/src/hipDeviceBody.cpp b/projects/hip/tests/src/hipDeviceBody.cpp
deleted file mode 100644
index 2e740cf748..0000000000
--- a/projects/hip/tests/src/hipDeviceBody.cpp
+++ /dev/null
@@ -1,5 +0,0 @@
-#include"hipDeviceHeader.h"
-
-__host__ __device__ int add(int &a, int &b){
-    return a+b;
-}
diff --git a/projects/hip/tests/src/hipDeviceHeader.h b/projects/hip/tests/src/hipDeviceHeader.h
deleted file mode 100644
index eb7bd05aa5..0000000000
--- a/projects/hip/tests/src/hipDeviceHeader.h
+++ /dev/null
@@ -1,7 +0,0 @@
-#ifndef HIPDEVICEHEADER_H
-#define HIPDEVICEHEADER_H
-#include"test_common.h"
-
-__host__ __device__ int add(int a, int b);
-
-#endif
diff --git a/projects/hip/tests/src/hipDeviceSrc.cpp b/projects/hip/tests/src/hipDeviceSrc.cpp
deleted file mode 100644
index 862d84f06a..0000000000
--- a/projects/hip/tests/src/hipDeviceSrc.cpp
+++ /dev/null
@@ -1,30 +0,0 @@
-#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);
-
-}
diff --git a/projects/hip/tests/src/hipTestDevice.cpp b/projects/hip/tests/src/hipTestDevice.cpp
new file mode 100644
index 0000000000..9e6450171b
--- /dev/null
+++ b/projects/hip/tests/src/hipTestDevice.cpp
@@ -0,0 +1,12 @@
+#include"test_common.h"
+#include"hip_runtime.h"
+
+__device__ void test(){
+    test__popc(1);
+}
+
+__global__ void Test(hipLaunchParm lp, int a){
+    test();
+}
+
+int main(){}