Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into doc_update

[ROCm/hip commit: 1fb48b0714]

projects/hip/include/hcc_detail/hip_runtime.h

@@ -62,9 +62,6 @@ THE SOFTWARE.
 #if defined(__HCC_ACCELERATOR__) and (__HCC_ACCELERATOR__ != 0)
 // Device compile and not host compile:
 
-
-#define __HIP_DEVICE_COMPILE__ 1
-
 //TODO-HCC enable __HIP_ARCH_HAS_ATOMICS__ when HCC supports these.
     // 32-bit Atomics:
 #define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__       (1)
@@ -95,10 +92,6 @@ THE SOFTWARE.
 #define __HIP_ARCH_HAS_3DGRID__                     (1)
 #define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__           (0)
 
-#else
-// Host compile and not device compile:
-#define __HIP_DEVICE_COMPILE__ 0
-
 #endif
 
 

projects/hip/include/hip_common.h

@@ -32,9 +32,11 @@ THE SOFTWARE.
 #define __HIP_PLATFORM_HCC__
 #define __HIPCC__
 
-# if defined __HCC_ACCELERATOR__
-# define __HIP_DEVICE_COMPILE__ 1
-# endif
+#if defined(__HCC_ACCELERATOR__) and (__HCC_ACCELERATOR__ != 0)
+#define __HIP_DEVICE_COMPILE__ 1
+#else
+#define __HIP_DEVICE_COMPILE__ 0
+#endif
 #endif
 
 // Auto enable __HIP_PLATFORM_NVCC__ if compiling with NVCC
@@ -44,9 +46,11 @@ THE SOFTWARE.
 # define __HIPCC__
 # endif
 
-# ifdef __CUDA_ARCH__
-# define __HIP_DEVICE_COMPILE__ 1
-# endif
+#if defined(__CUDA_ARCH__) and (__CUDA_ARCH__ != 0)
+#define __HIP_DEVICE_COMPILE__ 1
+#else
+#define __HIP_DEVICE_COMPILE__ 0
+#endif
 
 #endif
 

projects/hip/include/hip_runtime_api.h

@@ -93,6 +93,7 @@ typedef struct hipDeviceProp_t {
 	int clockInstructionRate ;   ///< Frequency in khz of the timer used by the device-side "clock*" instructions.  New for HIP.
 
     hipDeviceArch_t arch;  ///< Architectural feature flags.  New for HIP.
+    int concurrentKernels; ///< Device can possibly execute multiple kernels concurrently
  } hipDeviceProp_t;
 
 

projects/hip/include/nvcc_detail/hip_runtime.h

@@ -72,11 +72,6 @@ kernelName<<<numblocks,numthreads,memperblock,streamId>>>(0, __VA_ARGS__);\
 #define __HIP_ARCH_HAS_3DGRID__                     (__CUDA_ARCH__ >= 200)
 #define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__           (__CUDA_ARCH__ >= 350)
 
-
-#else
-
-#define __HIP_DEVICE_COMPILE__ 0
-
 #endif
 
 #ifdef __CUDACC__

projects/hip/include/nvcc_detail/hip_runtime_api.h

@@ -207,6 +207,7 @@ inline static hipError_t hipDeviceGetProperties(hipDeviceProp_t *p_prop, int dev
     p_prop->arch.has3dGrid                   =  (ccVers >= 200);
     p_prop->arch.hasDynamicParallelism       =  (ccVers >= 350);
 
+    p_prop->concurrentKernels = cdprop.concurrentKernels;
 
     return hipCUDAErrorTohipError(cerror);
 }