Merge 'amd-master-next' into 'amd-npi-next'

Change-Id: I935fc8f681fad2df4e932407287a29a6a797351a

tests/performance/dispatch/hipPerfDispatchSpeed.cpp

@@ -7,7 +7,7 @@
 #include "test_common.h"
 
 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp timer.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */

tests/performance/memory/hipPerfBufferCopyRectSpeed.cpp

@@ -7,7 +7,7 @@
 #include "test_common.h"
 
 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp timer.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */

tests/performance/memory/hipPerfBufferCopySpeed.cpp

@@ -7,7 +7,7 @@
 #include "test_common.h"
 
 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp timer.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */

tests/performance/memory/hipPerfDevMemReadSpeed.cpp

@@ -0,0 +1,136 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t
+ * HIT_END
+ */
+
+#include <iostream>
+#include <chrono>
+#include "test_common.h"
+
+using namespace std;
+
+#define arraySize 16
+
+typedef struct d_uint16 {
+  uint data[arraySize];
+} d_uint16;
+
+__global__ void read_kernel(d_uint16 *src, ulong N, uint *dst) {
+
+  size_t idx = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x ;
+
+  uint tmp = 0;
+  for (size_t i = idx; i < N; i += stride) {
+    for (size_t j = 0; j < arraySize; j++) {
+      tmp += src[i].data[j];
+    }
+  }
+
+  atomicAdd(dst, tmp);
+}
+
+int main(int argc, char* argv[]) {
+  d_uint16 *dSrc;
+  d_uint16 *hSrc;
+  uint *dDst;
+  uint *hDst;
+  hipStream_t stream;
+  ulong N = 4 * 1024 * 1024;
+  uint nBytes = N * sizeof(d_uint16);
+
+  int nGpu = 0;
+  HIPCHECK(hipGetDeviceCount(&nGpu));
+  if (nGpu < 1) {
+      cout << "info: didn't find any GPU! skipping the test!\n";
+      passed();
+      return 0;
+  }
+
+  static int device = 0;
+  HIPCHECK(hipSetDevice(device));
+  hipDeviceProp_t props;
+  HIPCHECK(hipGetDeviceProperties(&props, device));
+  cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name <<
+      " with " << props.multiProcessorCount << " CUs" << endl;
+
+  const unsigned threadsPerBlock = 64;
+  const unsigned blocks = props.multiProcessorCount * 4;
+
+  uint inputData = 0x1;
+  int nIter = 1000;
+
+  hSrc =  new d_uint16[nBytes];
+  HIPCHECK(hSrc == 0 ? hipErrorOutOfMemory : hipSuccess);
+  hDst =  new uint;
+  hDst[0] = 0;
+  HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
+  for (size_t i = 0; i < N; i++) {
+    for (int j = 0; j < arraySize; j++) {
+      hSrc[i].data[j] = inputData;
+    }
+  }
+
+  HIPCHECK(hipMalloc(&dSrc, nBytes));
+  HIPCHECK(hipMalloc(&dDst, sizeof(uint)));
+
+  HIPCHECK(hipStreamCreate(&stream));
+
+  HIPCHECK(hipMemcpy(dSrc, hSrc, nBytes, hipMemcpyHostToDevice));
+  HIPCHECK(hipMemcpy(dDst, hDst, sizeof(uint), hipMemcpyHostToDevice));
+
+  hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst);
+  HIPCHECK(hipMemcpy(hDst, dDst, sizeof(uint), hipMemcpyDeviceToHost));
+  hipDeviceSynchronize();
+
+  if (hDst[0] != (nBytes / sizeof(uint))) {
+    cout << "info: Data validation failed for warm up run!" << endl;
+    cout << "info: expected " << nBytes / sizeof(uint) << " got " << hDst[0] << endl;
+    HIPCHECK(hipErrorUnknown);
+  }
+
+  // measure performance based on host time
+  auto all_start = chrono::steady_clock::now();
+
+  for(int i = 0; i < nIter; i++) {
+    hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst);
+  }
+  hipDeviceSynchronize();
+
+  auto all_end = chrono::steady_clock::now();
+  chrono::duration<double> all_kernel_time = all_end - all_start;
+
+  // read speed in GB/s
+  double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count();
+
+  cout << "info: average read speed of " << perf << " GB/s " << "achieved for memory size of " <<
+      nBytes / (1024 * 1024) << " MB" << endl;
+
+  delete [] hSrc;
+  delete hDst;
+  hipFree(dSrc);
+  hipFree(dDst);
+  HIPCHECK(hipStreamDestroy(stream));
+
+  passed();
+}

tests/performance/memory/hipPerfDevMemWriteSpeed.cpp

@@ -0,0 +1,126 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t
+ * HIT_END
+ */
+
+#include <iostream>
+#include <chrono>
+#include "test_common.h"
+
+using namespace std;
+
+#define arraySize 16
+
+typedef struct d_uint16 {
+  uint data[arraySize];
+} d_uint16;
+
+__global__ void write_kernel(d_uint16 *dst, ulong N, d_uint16 pval) {
+    size_t idx = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x;
+    for (size_t i = idx; i < N; i += stride) {
+      dst[i] = pval;
+    }
+};
+
+int main(int argc, char* argv[]) {
+  d_uint16 *dDst;
+  d_uint16 *hDst;
+  hipStream_t stream;
+  ulong N = 4 * 1024 * 1024;
+  uint nBytes = N * sizeof(d_uint16);
+  d_uint16 pval;
+
+  for (int i = 0; i < arraySize; i++) {
+    pval.data[i] = 0xabababab;
+  }
+
+  int nGpu = 0;
+  HIPCHECK(hipGetDeviceCount(&nGpu));
+  if (nGpu < 1) {
+      cout << "info: didn't find any GPU! skipping the test!\n";
+      passed();
+      return 0;
+  }
+
+  static int device = 0;
+  HIPCHECK(hipSetDevice(device));
+  hipDeviceProp_t props;
+  HIPCHECK(hipGetDeviceProperties(&props, device));
+  cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name <<
+      " with " << props.multiProcessorCount << " CUs" << endl;
+
+  size_t threadsPerBlock = 64;
+  size_t blocks = props.multiProcessorCount * 4;
+
+  uint inputData = 0xabababab;
+  int nIter = 1000;
+
+  hDst =  new d_uint16[nBytes];
+  HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
+  for (size_t i = 0; i < N; i++) {
+    for (size_t j = 0; j < arraySize; j++) {
+      hDst[i].data[j] = 0;
+    }
+  }
+
+  HIPCHECK(hipMalloc(&dDst, nBytes));
+
+  HIPCHECK(hipStreamCreate(&stream));
+
+  hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval);
+  HIPCHECK(hipMemcpy(hDst, dDst, nBytes , hipMemcpyDeviceToHost));
+  hipDeviceSynchronize();
+
+  for (uint i = 0; i < N; i++) {
+    for (uint j = 0; j < arraySize; j++) {
+      if (hDst[i].data[j] != inputData) {
+        cout << "info: Data validation failed for warm up run! " << endl;
+        cout << "at index i: " << i << " element j: " << j << endl;
+        cout << hex << "expected 0x" << inputData << " but got 0x" << hDst[i].data[j] << endl;
+        HIPCHECK(hipErrorUnknown);
+      }
+    }
+  }
+
+  auto all_start = chrono::steady_clock::now();
+  for(int i = 0; i < nIter; i++) {
+    hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval);
+  }
+  hipDeviceSynchronize();
+  auto all_end = chrono::steady_clock::now();
+  chrono::duration<double> all_kernel_time = all_end - all_start;
+
+  // read speed in GB/s
+  double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count();
+
+  cout << "info: average write speed of " << perf << " GB/s " << "achieved for memory size of " <<
+      nBytes / (1024 * 1024) << " MB" << endl;
+
+
+  delete [] hDst;
+  hipFree(dDst);
+  HIPCHECK(hipStreamDestroy(stream));
+
+  passed();
+}

tests/performance/memory/hipPerfHostNumaAlloc.cpp

@@ -34,12 +34,13 @@ THE SOFTWARE.
 #include <array>
 #include "hip/hip_runtime.h"
 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD_CMD: hipPerfHostNumaAlloc %hc -I%S/../../src %S/%s %S/../../src/test_common.cpp -lnuma -o %T/%t EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
 
-// To run it correctly, we must not export HIP_VISIBLE_DEVICES
+// To run it correctly, we must not export HIP_VISIBLE_DEVICES.
+// And we must explicitly link libnuma because of numa api move_pages().
 #define NUM_PAGES 4
 char *h = nullptr;
 char *d_h = nullptr;
@@ -127,6 +128,7 @@ bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) {
   printf("\n");
 
   HIPCHECK(hipHostFree((void* )h));
+  hipHostUnregister(m);
   free(m);
 
   if (cpuId >= 0 && (numaMode == MPOL_BIND || numaMode == MPOL_PREFERRED)) {
@@ -149,8 +151,7 @@ bool runTest(const int &cpuCount, const int &gpuCount,
 
     for (int i = 0; i < cpuCount; i++) {
       for (int j = 0; j < gpuCount; j++) {
-        if (!test(i, j, mode[m],
-                  hipHostMallocDefault | hipHostMallocNumaUser)) {
+        if (!test(i, j, mode[m], hostMallocflags)) {
           return false;
         }
       }

tests/performance/memory/hipPerfSharedMemReadSpeed.cpp

@@ -0,0 +1,250 @@
+/*
+ Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights
+ to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ copies of the Software, and to permit persons to whom the Software is
+ furnished to do so, subject to the following conditions:
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE.
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t
+ * HIT_END
+ */
+
+#include <iostream>
+#include <chrono>
+#include "test_common.h"
+
+using namespace std;
+
+#define sharedMemSize1 2048
+#define sharedMemSize2 256
+
+__global__ void sharedMemReadSpeed1(float *outBuf, ulong N) {
+
+  size_t gid = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t lid = threadIdx.x;
+  __shared__ float local[sharedMemSize1];
+
+  float val1 = 0;
+  float val2 = 0;
+  float val3 = 0;
+  float val4 = 0;
+
+  for (int i = 0; i < (sharedMemSize1 / 64); i++) {
+    local[lid + i * 64] = lid;
+  }
+
+  __syncthreads();
+
+  val1 += local[lid];
+  val2 += local[lid + 64];
+  val3 += local[lid + 128];
+  val4 += local[lid + 192];
+  val1 += local[lid + 256];
+  val2 += local[lid + 320];
+  val3 += local[lid + 384];
+  val4 += local[lid + 448];
+  val1 += local[lid + 512];
+  val2 += local[lid + 576];
+  val3 += local[lid + 640];
+  val4 += local[lid + 704];
+  val1 += local[lid + 768];
+  val2 += local[lid + 832];
+  val3 += local[lid + 896];
+  val4 += local[lid + 960];
+  val1 += local[lid + 1024];
+  val2 += local[lid + 1088];
+  val3 += local[lid + 1152];
+  val4 += local[lid + 1216];
+  val1 += local[lid + 1280];
+  val2 += local[lid + 1344];
+  val3 += local[lid + 1408];
+  val4 += local[lid + 1472];
+  val1 += local[lid + 1536];
+  val2 += local[lid + 1600];
+  val3 += local[lid + 1664];
+  val4 += local[lid + 1728];
+  val1 += local[lid + 1792];
+  val2 += local[lid + 1856];
+  val3 += local[lid + 1920];
+  val4 += local[lid + 1984];
+
+  if (gid < N) {
+    outBuf[gid] = val1 + val2 + val3 + val4;
+  }
+};
+
+__global__ void sharedMemReadSpeed2(float *outBuf, ulong N) {
+  size_t gid = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t lid = threadIdx.x;
+  __shared__ float local[sharedMemSize2];
+
+  float val0 = 0.0f;
+  float val1 = 0.0f;
+
+  for (int i = 0; i < (sharedMemSize2 / 64); i++) {
+    local[lid + i * 64] = lid;
+  }
+
+  __syncthreads();
+
+#pragma nounroll
+  for (uint i = 0; i < 32; i++) {
+    val0 += local[8 * i + 0];
+    val1 += local[8 * i + 1];
+    val0 += local[8 * i + 2];
+    val1 += local[8 * i + 3];
+    val0 += local[8 * i + 4];
+    val1 += local[8 * i + 5];
+    val0 += local[8 * i + 6];
+    val1 += local[8 * i + 7];
+  }
+
+  if (gid < N) {
+    outBuf[gid] = val0 + val1;
+  }
+};
+
+int main(int argc, char *argv[]) {
+  float *dDst;
+  float *hDst;
+  hipStream_t stream;
+  constexpr uint numSizes = 4;
+  constexpr uint Sizes[numSizes] = {262144, 1048576, 4194304, 16777216};
+  uint numReads1 = 32;
+  uint numReads2 = 256;
+  uint sharedMemSizeBytes1 = sharedMemSize1 * sizeof(float);
+  uint sharedMemSizeBytes2 = sharedMemSize2 * sizeof(float);
+  int nIter = 1000;
+  const unsigned threadsPerBlock = 64;
+
+  int nGpu = 0;
+  HIPCHECK(hipGetDeviceCount(&nGpu));
+  if (nGpu < 1) {
+    cout << "info: didn't find any GPU! skipping the test!\n";
+    passed();
+    return 0;
+  }
+
+  static int device = 0;
+  HIPCHECK(hipSetDevice(device));
+  hipDeviceProp_t props;
+  HIPCHECK(hipGetDeviceProperties(&props, device));
+  cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name
+      << " with " << props.multiProcessorCount << " CUs" << endl;
+
+  HIPCHECK(hipStreamCreate(&stream));
+
+  for (int nTest = 0; nTest < numSizes; nTest++) {
+    uint nBytes = Sizes[nTest % numSizes];
+    ulong N = nBytes / sizeof(float);
+    const unsigned blocks = N / threadsPerBlock;
+
+    hDst = new float[nBytes];
+    HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
+    memset(hDst, 0, nBytes);
+
+    HIPCHECK(hipMalloc(&dDst, nBytes));
+    HIPCHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice));
+
+    hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), dim3(threadsPerBlock),
+        0, stream, dDst, N);
+    HIPCHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost));
+    hipDeviceSynchronize();
+
+    int tmp = 0;
+    for (int i = 0; i < N; i++) {
+      if (i % threadsPerBlock == 0) {
+        tmp = 0;
+      }
+      if (hDst[i] != tmp) {
+        cout << "info: Data validation failed for warm up run!" << endl;
+        cout << "info: expected " << tmp << " got " << hDst[i] << endl;
+        HIPCHECK (hipErrorUnknown);
+      }
+      tmp += threadsPerBlock / 2;
+    }
+
+    auto all_start = chrono::steady_clock::now();
+    for (int i = 0; i < nIter; i++) {
+      hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks),
+          dim3(threadsPerBlock), 0, stream, dDst, N);
+    }
+    hipDeviceSynchronize();
+
+    auto all_end = chrono::steady_clock::now();
+    chrono::duration<double> all_kernel_time = all_end - all_start;
+
+    // read speed in GB/s
+    double perf = ((double) blocks * threadsPerBlock
+        * (numReads1 * sizeof(float) + sharedMemSizeBytes1 / 64) * nIter
+        * (double) (1e-09)) / all_kernel_time.count();
+
+    cout << "info: read speed = " << setw(8) << perf << " GB/s for "
+        << sharedMemSizeBytes1 / 1024 << " KB shared memory"
+            " with " << setw(8) << blocks * threadsPerBlock << " threads, "
+        << setw(4) << numReads1 << " reads in sharedMemReadSpeed1 kernel" << endl;
+
+    delete[] hDst;
+    hipFree(dDst);
+  }
+
+
+  for (int nTest = 0; nTest < numSizes; nTest++) {
+    uint nBytes = Sizes[nTest % numSizes];
+    ulong N = nBytes / sizeof(float);
+    const unsigned blocks = N / threadsPerBlock;
+
+    hDst = new float[nBytes];
+    HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
+    memset(hDst, 0, nBytes);
+
+    HIPCHECK(hipMalloc(&dDst, nBytes));
+    HIPCHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice));
+
+    hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), dim3(threadsPerBlock),
+        0, stream, dDst, N);
+    HIPCHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost));
+    hipDeviceSynchronize();
+
+    auto all_start = chrono::steady_clock::now();
+    for (int i = 0; i < nIter; i++) {
+      hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks),
+          dim3(threadsPerBlock), 0, stream, dDst, N);
+    }
+    hipDeviceSynchronize();
+
+    auto all_end = chrono::steady_clock::now();
+    chrono::duration<double> all_kernel_time = all_end - all_start;
+
+    // read speed in GB/s
+    double perf = ((double) blocks * threadsPerBlock
+        * (numReads2 * sizeof(float) + sharedMemSizeBytes2 / 64) * nIter
+        * (double) (1e-09)) / all_kernel_time.count();
+
+    cout << "info: read speed = " << setw(8) << perf << " GB/s for "
+        << sharedMemSizeBytes2 / 1024 << " KB shared memory"
+            " with " << setw(8) << blocks * threadsPerBlock << " threads, "
+        << setw(4) << numReads2 << " reads in sharedMemReadSpeed2 kernel" << endl;
+
+    delete[] hDst;
+    hipFree(dDst);
+  }
+
+  HIPCHECK(hipStreamDestroy(stream));
+
+  passed();
+}

tests/src/cg/hipCGThreadBlockType.cpp

@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp
+ * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */

tests/src/deviceLib/hipHalf2Comparision.cpp

@@ -19,12 +19,13 @@ THE SOFTWARE.
 
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
+ * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
 
 #include "test_common.h"
+#include <hip/hip_runtime.h>
 #include "hip/hip_fp16.h"
 
 #define test_passed(test_name) \

tests/src/g++/hipMalloc.cpp

@@ -18,7 +18,7 @@
  * */
 
 /* HIT_START
- * BUILD_CMD: hipMalloc %cxx -D__HIP_PLATFORM_HCC__ -I%hip-path/include -I/opt/rocm/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t -std=c++11 EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD_CMD: hipMalloc %cxx -D__HIP_PLATFORM_HCC__ -I%hip-path/include -I/opt/rocm/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libamdhip64.so -o %T/%t -std=c++11 EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */

tests/src/gcc/LaunchKernel.c

@@ -21,7 +21,7 @@
 /* HIT_START
  * BUILD_CMD: gpu.o %hc -I%hip-path/include -g -c %S/gpu.cpp -o %T/gpu.o EXCLUDE_HIP_PLATFORM nvcc rocclr
  * BUILD_CMD: launchkernel.o %hc -D__HIP_PLATFORM_HCC__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvcc rocclr
- * BUILD_CMD: LaunchKernel %hc %T/launchkernel.o %T/gpu.o -g -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t DEPENDS gpu.o launchkernel.o EXCLUDE_HIP_PLATFORM nvcc rocclr
+ * BUILD_CMD: LaunchKernel %hc %T/launchkernel.o %T/gpu.o -g -Wl,--rpath=%hip-path/lib %hip-path/lib/libamdhip64.so -o %T/%t DEPENDS gpu.o launchkernel.o EXCLUDE_HIP_PLATFORM nvcc rocclr
  * TEST: %t EXCLUDE_HIP_PLATFORM nvcc rocclr
  * HIT_END
  */

tests/src/gcc/hipMalloc.c

@@ -19,7 +19,7 @@
 
 /* HIT_START
  * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_NVCC__ -I%hip-path/include -I/usr/local/cuda/include %S/%s -o %T/hipMalloc_nv -L/usr/local/cuda/lib64 -lcudart EXCLUDE_HIP_PLATFORM hcc rocclr
- * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc rocclr
+ * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libamdhip64.so -o %T/hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc rocclr
  * TEST: hipMalloc_nv EXCLUDE_HIP_PLATFORM hcc rocclr
  * TEST: hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc rocclr
  * HIT_END

tests/src/ipc/hipSimpleIpc.cpp

@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp
+ * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
@@ -43,6 +43,7 @@ void single_process() {
 
   // Negative, Make sure we return error when an offset of original ptr is passed
   ipc_offset_dptr = ipc_dptr + (OFFSET * sizeof(int));
+  // HIP API return value differs from CUDA's return type
   assert(hipErrorInvalidDevicePointer == hipIpcGetMemHandle(&ipc_offset_handle, ipc_offset_dptr));
 
   // Get handle for the device_ptr

tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp

@@ -26,7 +26,7 @@
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
  * TEST_NAMED: %t  hipDeviceGetPCIBusId-vs-hipDeviceGetAttribute --tests 0x1
- * TEST_NAMED: %t  hipDeviceGetPCIBusId-vs-lspci --tests 0x2
+ * TEST_NAMED: %t  hipDeviceGetPCIBusId-vs-lspci --tests 0x2 EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */
 
@@ -106,8 +106,13 @@ bool compareHipDeviceGetPCIBusIdWithLspci() {
   getPciBusId(deviceCount, hipDeviceList);
 
   // Get lspci device list and compare with hip device list
+#if defined(__CUDA_ARCH__)
+  char const *command = "lspci -D | grep controller | grep NVIDIA | "
+                        "cut -d ' ' -f 1";
+#else
   char const *command = "lspci -D | grep controller | grep AMD/ATI | "
                         "cut -d ' ' -f 1";
+#endif
   fpipe = popen(command, "r");
 
   if (fpipe == nullptr) {

tests/src/runtimeApi/memory/hipMemcpyWithStreamMultiThread.cpp

@@ -467,6 +467,7 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(void) {
                       &A_h[0], &B_h[0], &C_h[0], N, false);
 
   for (int i=1; i < numDevices; ++i) {
+    HIPCHECK(hipSetDevice(i));
     HIPCHECK(hipMalloc(&A_d[i], Nbytes));
     HIPCHECK(hipMalloc(&B_d[i], Nbytes));
     HIPCHECK(hipMalloc(&C_d[i], Nbytes));
@@ -476,6 +477,7 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindDefaultForDtoD(void) {
 
   hipStream_t stream[numDevices];
   for (int i=0; i < numDevices; ++i) {
+    HIPCHECK(hipSetDevice(i));
     HIPCHECK(hipStreamCreate(&stream[i]));
   }
 

tests/src/runtimeApi/memory/hipMemset.cpp

@@ -166,11 +166,11 @@ bool testhipMemset2AsyncOps() {
   hipStream_t s;
   hipStreamCreate(&s);
   hipMemsetAsync(p2, 0, 32*32*4, s);
-  hipMemsetD32Async(p3, 0x3fe00000, 32*32, s );
+  hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s );
   hipStreamSynchronize(s);
   for (int i = 0; i < 256; ++i) {
     hipMemsetAsync(p2, 0, 32*32*4, s);
-    hipMemsetD32Async(p3, 0x3fe00000, 32*32, s );
+    hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s );
   }
   hipStreamSynchronize(s);
   hipDeviceSynchronize();

tests/src/runtimeApi/module/hipModuleLoadDataMultThreadOnMultGPU.cpp

@@ -60,8 +60,6 @@ void run(const std::vector<char>& buffer, int deviceNo) {
   hipSetDevice(deviceNo);
   hipModule_t Module;
   hipFunction_t Function;
-  HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
-  HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
 
   float *A, *B, *Ad, *Bd;
   A = new float[LEN];
@@ -78,6 +76,9 @@ void run(const std::vector<char>& buffer, int deviceNo) {
   HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
   HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
 
+  HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
+  HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
+
   hipStream_t stream;
   HIPCHECK(hipStreamCreate(&stream));
 

tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp

@@ -56,8 +56,6 @@ std::vector<char> load_file() {
 void run(const std::vector<char>& buffer) {
   hipModule_t Module;
   hipFunction_t Function;
-  HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
-  HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
 
   float *A, *B, *Ad, *Bd;
   A = new float[LEN];
@@ -74,6 +72,9 @@ void run(const std::vector<char>& buffer) {
   HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
   HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
 
+  HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
+  HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
+
   hipStream_t stream;
   HIPCHECK(hipStreamCreate(&stream));
 

tests/src/runtimeApi/module/tex2d_kernel.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD_CMD: tex2d_kernel.code %hc --genco %S/tex2d_kernel.cpp -o tex2d_kernel.code EXCLUDE_HIP_PLATFORM rocclr
+ * BUILD_CMD: tex2d_kernel.code %hc --genco %S/tex2d_kernel.cpp -o tex2d_kernel.code EXCLUDE_HIP_PLATFORM rocclr nvcc
  * HIT_END
  */
 

tests/src/runtimeApi/stream/hipStreamACb_StrmSyncTiming.cpp

@@ -81,6 +81,13 @@ static void HIPRT_CB Callback1(hipStream_t stream, hipError_t status,
   sleep(SECONDS_TO_WAIT);
 }
 
+bool rangedCompare(long a, long b) {
+  auto diff = b - a;
+  if (diff < 0) diff *= -1;
+  if (diff < 500) return true;
+  return false;
+}
+
 
 int main(int argc, char* argv[]) {
   float *A_d, *C_d;
@@ -139,7 +146,8 @@ int main(int argc, char* argv[]) {
   // completes the execution. Therefore the hipStreamSynchronize() in the
   // main thread should hardly take any time to complete.
 
-  if (duration.count() < SECONDS_TO_WAIT * TO_MICROSECONDS) {
+  if ((duration.count() < (SECONDS_TO_WAIT * TO_MICROSECONDS)) ||
+      (rangedCompare(duration.count(), SECONDS_TO_WAIT * TO_MICROSECONDS))) {
     passed();
   } else {
     failed("hipStreamSynchronize is waiting untill Callback() completes.");

tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp

@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all
  * TEST: %t
  * HIT_END
  */

tests/src/runtimeApi/stream/hipStreamGetPriority.cpp

@@ -45,7 +45,7 @@ int main(int argc, char *argv[]) {
 
     // Check if priorities are indeed supported
     if ((priority_low + priority_high) != 0) {
-      failed("Priorities are not supported");
+      passed(); // exit the test since priorities are not supported
     }
 
     // Checking Priority of default stream