diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 6565eaacb8..dbeedb77b1 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -167,9 +167,10 @@ if ($HIP_PLATFORM eq "hcc") { $CUDA_PATH=$ENV{'CUDA_PATH'} // '/usr/local/cuda'; $HIPCC="$CUDA_PATH/bin/nvcc"; + $HIPCXXFLAGS .= " -Wno-deprecated-gpu-targets "; $HIPCXXFLAGS .= " -I$CUDA_PATH/include"; - $HIPLDFLAGS = "-lcuda -lcudart -L$CUDA_PATH/lib64"; + $HIPLDFLAGS = " -Wno-deprecated-gpu-targets -lcuda -lcudart -L$CUDA_PATH/lib64"; } else { printf ("error: unknown HIP_PLATFORM = '$HIP_PLATFORM'"); exit (-1); diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index efc15ea70c..7272639d70 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/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/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 870dcc5dae..dbd6d8b300 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/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/hip/src/math_functions.cpp b/projects/hip/src/math_functions.cpp index 9118318469..01fe28b84e 100644 --- a/projects/hip/src/math_functions.cpp +++ b/projects/hip/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); } diff --git a/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp b/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp index e845ae8f2f..2dd5d42b0f 100644 --- a/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp +++ b/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp @@ -29,7 +29,7 @@ __global__ void set(hipLaunchParm lp, uint32_t *ptr, uint8_t val, size_t size) int main() { uint32_t *A, *Ad, *B, *Bd; - uint32_t *Val, *Vald; + uint32_t *Val; A = new uint32_t[LEN]; B = new uint32_t[LEN]; Val = new uint32_t; diff --git a/projects/hip/tests/src/deviceLib/hipFloatMath.cpp b/projects/hip/tests/src/deviceLib/hipFloatMath.cpp index 9c97e50aa6..38c8db3398 100644 --- a/projects/hip/tests/src/deviceLib/hipFloatMath.cpp +++ b/projects/hip/tests/src/deviceLib/hipFloatMath.cpp @@ -55,7 +55,7 @@ __global__ void floatMath(hipLaunchParm lp, float *In, float *Out) { } int main(){ - float *Inh, *Outh, *Ind, *Outd; + float *Ind, *Outd; hipMalloc((void**)&Ind, SIZE); hipMalloc((void**)&Outd, SIZE); hipLaunchKernel(floatMath, dim3(LEN,1,1), dim3(1,1,1), 0, 0, Ind, Outd); diff --git a/projects/hip/tests/src/deviceLib/hipFloatMathPrecise.cpp b/projects/hip/tests/src/deviceLib/hipFloatMathPrecise.cpp index 78fe47cfe8..d1e2f782b7 100644 --- a/projects/hip/tests/src/deviceLib/hipFloatMathPrecise.cpp +++ b/projects/hip/tests/src/deviceLib/hipFloatMathPrecise.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. __global__ void FloatMathPrecise(hipLaunchParm lp) { - int iX; + //int iX; //uncomment this when remqouf() is enabled again float fX, fY; acosf(1.0f); diff --git a/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp index 5485e4afc9..090a2f5697 100644 --- a/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp +++ b/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp @@ -287,7 +287,7 @@ void runTest(int argc, char **argv) "SM %d.%d compute capabilities\n\n", deviceProp.multiProcessorCount, deviceProp.major, deviceProp.minor); - int version = (deviceProp.major * 0x10 + deviceProp.minor); + unsigned int numThreads = 256; unsigned int numBlocks = 64; diff --git a/projects/hip/tests/src/deviceLib/hipTestDevice.cpp b/projects/hip/tests/src/deviceLib/hipTestDevice.cpp index 068866021b..990462d0cd 100644 --- a/projects/hip/tests/src/deviceLib/hipTestDevice.cpp +++ b/projects/hip/tests/src/deviceLib/hipTestDevice.cpp @@ -249,7 +249,6 @@ hipMemcpy(B, Bd, N*sizeof(long long int), hipMemcpyDeviceToHost); int passed = 0; for(int i=0;i<512;i++){ int x = roundf(A[i]); - long long int y = x; if(B[i] == x){ passed = 1; } @@ -283,7 +282,6 @@ hipMemcpy(B, Bd, N*sizeof(long int), hipMemcpyDeviceToHost); int passed = 0; for(int i=0;i<512;i++){ int x = roundf(A[i]); - long int y = x; if(B[i] == x){ passed = 1; } @@ -351,7 +349,6 @@ hipMemcpy(B, Bd, N*sizeof(long long int), hipMemcpyDeviceToHost); int passed = 0; for(int i=0;i<512;i++){ int x = roundf(A[i]); - long long int y = x; if(B[i] == x){ passed = 1; } @@ -385,7 +382,6 @@ hipMemcpy(B, Bd, N*sizeof(long int), hipMemcpyDeviceToHost); int passed = 0; for(int i=0;i<512;i++){ int x = roundf(A[i]); - long int y = x; if(B[i] == x){ passed = 1; } diff --git a/projects/hip/tests/src/deviceLib/hip_anyall.cpp b/projects/hip/tests/src/deviceLib/hip_anyall.cpp index f0b314ce8d..06354383df 100644 --- a/projects/hip/tests/src/deviceLib/hip_anyall.cpp +++ b/projects/hip/tests/src/deviceLib/hip_anyall.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp * RUN: %t * HIT_END */ diff --git a/projects/hip/tests/src/kernel/hipLanguageExtensions.cpp b/projects/hip/tests/src/kernel/hipLanguageExtensions.cpp index 5c5397675c..af5ec7cfa3 100644 --- a/projects/hip/tests/src/kernel/hipLanguageExtensions.cpp +++ b/projects/hip/tests/src/kernel/hipLanguageExtensions.cpp @@ -90,14 +90,6 @@ vectorADD(const hipLaunchParm lp, size_t N) { // KERNELBEGIN; - int ws = warpSize; - - - int zuzu = deviceVar + 1; - - - int b = threadIdx.x; - int c; #ifdef NOT_YET int a = __shfl_up(x, 1); #endif @@ -109,6 +101,10 @@ vectorADD(const hipLaunchParm lp, #endif #ifdef __HCC__ + + int b = threadIdx.x; + int c; + // TODO - move to HIP atomics when ready. concurrency :: atomic_fetch_add(&c, b); //Concurrency::atomic_add_unsigned (&x, a); diff --git a/projects/hip/tests/src/hipPeerToPeer_simple.cpp b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp similarity index 98% rename from projects/hip/tests/src/hipPeerToPeer_simple.cpp rename to projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp index 1ea594f4bb..32cab371b8 100644 --- a/projects/hip/tests/src/hipPeerToPeer_simple.cpp +++ b/projects/hip/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/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index 8e55e2f699..2919939694 100644 --- a/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/projects/hip/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; } diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp index b3f25658fc..7b47576e5d 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp @@ -192,8 +192,6 @@ HostMemory::~HostMemory () free(A_hh); free(B_hh); } - T *A_hh = NULL; - T *B_hh = NULL; }; diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp index 95b19c1090..a1ec69f9be 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp @@ -32,7 +32,6 @@ THE SOFTWARE. int main() { - hipDevice_t device; size_t Nbytes = N*sizeof(int); int numDevices = 0; int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d; diff --git a/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp b/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp index 81450f1fba..b1f9d45bf1 100644 --- a/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp @@ -188,7 +188,6 @@ int main(int argc, char *argv[]) { HipTest::parseStandardArguments(argc, argv, true); - int numElements = N; int dev0 = 0; int dev1 = 1; diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp index 07a64a194d..32a2793479 100644 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp @@ -27,7 +27,10 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "test_common.h" + +#ifdef __HIP_PLATFORM_HCC__ #define HIPRT_CB +#endif class CallbackClass { diff --git a/projects/hip/tests/src/test_common.h b/projects/hip/tests/src/test_common.h index 1b9744647b..c2eb0d4f4e 100644 --- a/projects/hip/tests/src/test_common.h +++ b/projects/hip/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,\