Calling intrinsics from global kernel

[ROCm/hip commit: 61cd2bb399]

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

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

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

projects/hip/tests/src/hipDeviceBody.cpp

@@ -1,5 +0,0 @@
-#include"hipDeviceHeader.h"
-
-__host__ __device__ int add(int &a, int &b){
-    return a+b;
-}

projects/hip/tests/src/hipDeviceHeader.h

@@ -1,7 +0,0 @@
-#ifndef HIPDEVICEHEADER_H
-#define HIPDEVICEHEADER_H
-#include"test_common.h"
-
-__host__ __device__ int add(int a, int b);
-
-#endif

projects/hip/tests/src/hipDeviceSrc.cpp

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

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(){}