Fixing/Disabling tests on nvcc path

Change-Id: I7ed4a1246943867482ff595d5f98bb25bd79074b


[ROCm/clr commit: d251c2db57]

projects/clr/hipamd/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
  */

projects/clr/hipamd/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) \

projects/clr/hipamd/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

projects/clr/hipamd/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) {

projects/clr/hipamd/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();

projects/clr/hipamd/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));
 

projects/clr/hipamd/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));
 

projects/clr/hipamd/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
  */
 

projects/clr/hipamd/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.");

projects/clr/hipamd/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