From 880cd966de36c8fd3f571c065988c3f57b1e955c Mon Sep 17 00:00:00 2001 From: Aaron En Ye Shi Date: Tue, 5 Dec 2017 23:14:10 +0000 Subject: [PATCH 1/3] Fix ilogb/ilogbf functions to return int This patch will fix hipDoublePrecisionMathDevice test on ThinLTO, which uncovered that hip math_function's ilogb/ilogbf should return type int instead of double. This will match rocdl. [ROCm/clr commit: b439b45641ec2a3d57d2bd5eaba72257e9924995] --- projects/clr/hipamd/include/hip/hcc_detail/math_functions.h | 4 ++-- projects/clr/hipamd/src/math_functions.cpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h index efc15ea70c..7272639d70 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h +++ b/projects/clr/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); diff --git a/projects/clr/hipamd/src/math_functions.cpp b/projects/clr/hipamd/src/math_functions.cpp index 80ccece1a3..d75ca9005a 100644 --- a/projects/clr/hipamd/src/math_functions.cpp +++ b/projects/clr/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); } @@ -544,7 +544,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); } From 156e35cfe900470654462d80cee0ec86a5f1fbbe Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 6 Dec 2017 15:49:06 +0530 Subject: [PATCH 2/3] Fix hipGetDeviceAttribute dtest for HIP/NVCC [ROCm/clr commit: a62ef42c09b80b033680a84c13b41aa064bf7ae2] --- .../clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 4 ++++ .../tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp | 6 +++--- 2 files changed, 7 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 870dcc5dae..dbd6d8b300 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/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); } diff --git a/projects/clr/hipamd/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/projects/clr/hipamd/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index 8e55e2f699..2919939694 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/projects/clr/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; } From 56754062fc370ac9cb6f4bb38af5c599e697b148 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 11 Dec 2017 15:21:39 +0530 Subject: [PATCH 3/3] Fixed hipPeerToPeer_simple test - Moved test inside p2p dir - Updated HIPCHECK to ignore hipErrorPeerAccessAlreadyEnabled - Added check for mGPUs [ROCm/clr commit: 2de0f1cafd7e94d9aa271149454c5f0b0c792eb3] --- .../hipamd/tests/src/{ => p2p}/hipPeerToPeer_simple.cpp | 9 ++++++++- projects/clr/hipamd/tests/src/test_common.h | 2 +- 2 files changed, 9 insertions(+), 2 deletions(-) rename projects/clr/hipamd/tests/src/{ => p2p}/hipPeerToPeer_simple.cpp (98%) diff --git a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp similarity index 98% rename from projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp rename to projects/clr/hipamd/tests/src/p2p/hipPeerToPeer_simple.cpp index 1ea594f4bb..32cab371b8 100644 --- a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp +++ b/projects/clr/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*/); diff --git a/projects/clr/hipamd/tests/src/test_common.h b/projects/clr/hipamd/tests/src/test_common.h index 1b9744647b..c2eb0d4f4e 100644 --- a/projects/clr/hipamd/tests/src/test_common.h +++ b/projects/clr/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,\