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

bin/hipconfig

@@ -19,7 +19,7 @@ GetOptions(
 );
 
 if ($p_help) {
-    print "usage: which_hip [OPTIONS]\n";
+    print "usage: hipconfig [OPTIONS]\n";
     print "  --path,  -p        : print HIP_PATH (use env var if set, else determine from hipconfig path)\n";
     print "  --cpp_config, -C   : print C++ compiler options\n";
     print "  --compiler, -c     : print compiler (hcc or nvcc)\n";

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
 
 

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
 

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;
 
 

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__

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

src/hip_hcc.cpp

@@ -383,6 +383,7 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop)
     prop->arch.has3dGrid                   = 1;
     prop->arch.hasDynamicParallelism       = 0;
 
+    prop->concurrentKernels = 1; // All ROCR hardware supports executing multiple kernels concurrently
     return e;
 }
 

tests/src/CMakeLists.txt

@@ -103,7 +103,7 @@ make_hip_executable (hip_popc hip_popc.cpp)
 make_hip_executable (hip_clz hip_clz.cpp) 
 make_hip_executable (hip_brev hip_brev.cpp) 
 make_hip_executable (hip_ffs hip_ffs.cpp) 
-make_hip_executable (hipInfo hipInfo.cpp) 
+make_hip_executable (hipGetDeviceAttribute hipGetDeviceAttribute.cpp) 
 make_hip_executable (hipMemcpy hipMemcpy.cpp) 
 make_hip_executable (hipMemset hipMemset.cpp) 
 make_hip_executable (hipEventRecord hipEventRecord.cpp)