Merge 'master' into 'amd-master'

Change-Id: I3355f0b291d7fff748ea40af473128941a1f792a

hipamd/include/hip/hcc_detail/math_functions.h

@@ -67,7 +67,7 @@ __device__ float fminf(float x, float y);
 __device__ float fmodf(float x, float y);
 __device__ float frexpf(float x, int* nptr);
 __device__ float hypotf(float x, float y);
-__device__ float ilogbf(float x);
+__device__ int ilogbf(float x);
 __device__ int isfinite(float a);
 __device__ unsigned isinf(float a);
 __device__ unsigned isnan(float a);
@@ -153,7 +153,7 @@ __device__ double fmin(double x, double y);
 __device__ double fmod(double x, double y);
 __device__ double frexp(double x, int *nptr);
 __device__ double hypot(double x, double y);
-__device__ double ilogb(double x);
+__device__ int ilogb(double x);
 __device__ int isfinite(double x);
 __device__ unsigned isinf(double x);
 __device__ unsigned isnan(double x);

hipamd/include/hip/nvcc_detail/hip_runtime_api.h

@@ -597,6 +597,10 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t *p_prop, int dev
     p_prop->arch.hasDynamicParallelism       =  (ccVers >= 350);
 
     p_prop->concurrentKernels = cdprop.concurrentKernels;
+    p_prop->pciBusID = cdprop.pciBusID;
+    p_prop->pciDeviceID = cdprop.pciDeviceID;
+    p_prop->maxSharedMemoryPerMultiProcessor = cdprop.sharedMemPerMultiprocessor;
+    p_prop->isMultiGpuBoard = cdprop.isMultiGpuBoard;
 
     return hipCUDAErrorTohipError(cerror);
 }

hipamd/src/math_functions.cpp

@@ -158,7 +158,7 @@ __device__ float hypotf(float x, float y)
 {
     return hc::precise_math::hypotf(x, y);
 }
-__device__ float ilogbf(float x)
+__device__ int ilogbf(float x)
 {
     return hc::precise_math::ilogbf(x);
 }
@@ -539,7 +539,7 @@ __device__ double hypot(double x, double y)
 {
     return hc::precise_math::hypot(x, y);
 }
-__device__ double ilogb(double x)
+__device__ int ilogb(double x)
 {
     return hc::precise_math::ilogb(x);
 }

hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp

@@ -23,7 +23,7 @@ THE SOFTWARE.
 // Also serves as a template for other tests.
 
 /* HIT_START
- * BUILD: %t %s test_common.cpp
+ * BUILD: %t %s ../test_common.cpp
  * RUN: %t EXCLUDE_HIP_PLATFORM all
  * RUN: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM all
  * RUN: %t --mirrorPeers EXCLUDE_HIP_PLATFORM all
@@ -391,7 +391,14 @@ void simpleNegative()
 int main(int argc, char *argv[])
 {
     parseMyArguments(argc, argv);
+    int gpuCount;
+    HIPCHECK(hipGetDeviceCount(&gpuCount));
 
+    if (gpuCount < 2)
+    {
+        printf("P2P application requires atleast 2 gpu devices\n");
+        return 0;
+    }
 
     if (p_tests & 0x100) {
         testPeerHostToDevice(false/*useAsyncCopy*/);

hipamd/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp

@@ -39,14 +39,14 @@ THE SOFTWARE.
     exit(EXIT_FAILURE);\
     }
 
-hipError_t test_hipDeviceGetAttribute(int deviceId, hipDeviceAttribute_t attr, int expectedValue = 0)
+hipError_t test_hipDeviceGetAttribute(int deviceId, hipDeviceAttribute_t attr, int expectedValue = -1)
 {
     int value = 0;
     std::cout << "Test hipDeviceGetAttribute attribute " << attr;
-    if (expectedValue) { std::cout << " expected value " << expectedValue; }
+    if (expectedValue != -1) { std::cout << " expected value " << expectedValue; }
     hipError_t e = hipDeviceGetAttribute(&value, attr, deviceId);
     std::cout << " actual value " << value << std::endl;
-    if (expectedValue && value != expectedValue) {
+    if ((expectedValue != -1) && value != expectedValue) {
         std::cout << "fail" << std::endl;
         return hipErrorInvalidValue;
     }

hipamd/tests/src/test_common.h

@@ -70,7 +70,7 @@ THE SOFTWARE.
 #define HIPCHECK(error) \
 {\
     hipError_t localError = error; \
-    if (localError != hipSuccess) { \
+    if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)){ \
       printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \
           KRED,hipGetErrorString(localError), localError,\
           #error,\