added C guard to hip_runtime.h

[ROCm/hip commit: d5feabfa1c]

projects/hip/bin/hipcc

@@ -66,13 +66,13 @@ if ($HIP_PLATFORM eq "hcc") {
 
     # HCC* may be used to compile src/hip_hcc.o (and also feed the HIPCXXFLAGS below)
     $HCC = "$HCC_HOME/bin/hcc";
-    $HCCFLAGS = " -hc -I$HCC_HOME/include ";
+    $HCCFLAGS = " -I$HCC_HOME/include ";
 
     $HIPCC=$HCC;
     $HIPCXXFLAGS = $HCCFLAGS;
     $HIPCXXFLAGS .= " -I$HIP_PATH/include/hcc_detail/cuda";
     $HIPCXXFLAGS .= " -I$HSA_PATH/include";
-	$HIPLDFLAGS =  "-hc -L$HCC_HOME/lib -Wl,--rpath=$HCC_HOME/lib -lc++ -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive";
+	$HIPLDFLAGS =  " -L$HCC_HOME/lib -Wl,--rpath=$HCC_HOME/lib -lc++ -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive";
     # Suppress linker warnings in case HCC distribution contains OpenCL/SPIR symbols
     $HIPLDFLAGS .= " -Wl,--defsym=_binary_kernel_spir_end=0 -Wl,--defsym=_binary_kernel_spir_start=0 -Wl,--defsym=_binary_kernel_cl_start=0 -Wl,--defsym=_binary_kernel_cl_end=0";
     # Satisfy HCC dependencies

projects/hip/include/hcc_detail/hip_runtime.h

@@ -43,17 +43,18 @@ THE SOFTWARE.
 //---
 // Remainder of this file only compiles with HCC
 #ifdef __HCC__
-//#if __cplusplus
+#if __cplusplus
 #include <hc.hpp>
-//#endif
+#endif
 #include <grid_launch.h>
 extern int HIP_TRACE_API;
 
 //TODO-HCC-GL - change this to typedef.
 //typedef grid_launch_parm hipLaunchParm ;
 #define hipLaunchParm grid_launch_parm
-
+#ifdef __cplusplus
 #include <hcc_detail/hip_texture.h>
+#endif
 #include <hcc_detail/host_defines.h>
 // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
 #if defined (__KALMAR_ACCELERATOR__) && !defined (__HCC_ACCELERATOR__)
@@ -107,6 +108,7 @@ extern int HIP_TRACE_API;
 #define __HCC_C__
 #endif
 
+#if __cplusplus
 __device__ float acosf(float x);
 __device__ float acoshf(float x);
 __device__ float asinf(float x);
@@ -286,7 +288,7 @@ __device__ double trunc(double x);
 __device__ double y0(double x);
 __device__ double y1(double y);
 __device__ double yn(int n, double x);
-
+#endif
 
 // TODO - hipify-clang - change to use the function call.
 //#define warpSize hc::__wavesize()
@@ -296,7 +298,7 @@ extern const int warpSize;
 #define clock_t long long int
 __device__ long long int clock64();
 __device__ clock_t clock();
-
+#if __cplusplus
 //atomicAdd()
 __device__ int atomicAdd(int* address, int val);
 __device__ unsigned int atomicAdd(unsigned int* address,
@@ -441,7 +443,7 @@ __device__ float __dsqrt_rd(double x);
 __device__ float __dsqrt_rn(double x);
 __device__ float __dsqrt_ru(double x);
 __device__ float __dsqrt_rz(double x);
-
+#endif
 /**
  * Kernel launching
  */

projects/hip/tests/src/hipC.c

@@ -3,7 +3,7 @@
 
 #define ITER 1<<20
 #define SIZE 1024*1024*sizeof(int)
-/*
+
 __global__ void Iter(hipLaunchParm lp, int *Ad){
     int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
     if(tx == 0){
@@ -12,11 +12,14 @@ __global__ void Iter(hipLaunchParm lp, int *Ad){
         }
     }
 }
-*/
+
 int main(){
     int A=0, *Ad;
     hipMalloc((void**)&Ad, SIZE);
     hipMemcpy(Ad, &A, SIZE, hipMemcpyHostToDevice);
-    hipLaunchKernel(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, 0, Ad);
+    dim3 dimGrid, dimBlock;
+    dimGrid.x = 1, dimGrid.y =1, dimGrid.z = 1;
+    dimBlock.x = 1, dimBlock.y = 1, dimGrid.z = 1;
+    hipLaunchKernel(HIP_KERNEL_NAME(Iter), dimGrid, dimBlock, 0, 0, Ad);
     hipMemcpy(&A, Ad, SIZE, hipMemcpyDeviceToHost);
 }