Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_purge_genco

[ROCm/hip commit: e14ac87474]
2017-12-14 13:50:49 +00:00
commit 044ae14b92
@@ -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);
@@ -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);
@@ -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);
 }
@@ -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);
 }
@@ -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;
@@ -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);
@@ -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);
@@ -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;
@@ -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;
    }
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../test_common.cpp
+ * BUILD: %t %s ../test_common.cpp 
 * RUN: %t
 * HIT_END
 */
@@ -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);
@@ -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*/);
@@ -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;
    }
@@ -192,8 +192,6 @@ HostMemory<T>::~HostMemory ()
        free(A_hh);
        free(B_hh);
    }
-    T *A_hh = NULL;
-    T *B_hh = NULL;

 };

@@ -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;
@@ -188,7 +188,6 @@ int main(int argc, char *argv[])
 {
    HipTest::parseStandardArguments(argc, argv, true);

-    int numElements = N;

    int dev0 = 0;
    int dev1 = 1;
@@ -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
 {
@@ -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,\