SWDEV-289409 - Add hiprtc tests

Change-Id: Ib26527b704aed32ae3f3ed38bf6e2fd412462c8e


[ROCm/hip-tests commit: 46fb008ba6]

projects/hip-tests/catch/hipTestMain/CMakeLists.txt

@@ -14,11 +14,11 @@ target_link_libraries(UnitTests PRIVATE UnitDeviceTests
                                         StreamTest
                                         OccupancyTest
                                         DeviceTest
+                                        RTC
                                         stdc++fs)
 
-# Add AMD Only Tests
-if(HIP_PLATFORM MATCHES "amd")
-    # target_link_libraries(UnitTests PRIVATE RTC)
+if(HIP_PLATFORM MATCHES "nvidia")
+    target_link_libraries(UnitTests PRIVATE nvrtc)
 endif()
 
 catch_discover_tests(UnitTests PROPERTIES  SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")

projects/hip-tests/catch/include/hip_test_checkers.hh

@@ -2,13 +2,13 @@
 #include "hip_test_common.hh"
 #include <iostream>
 using namespace std;
-#define guarantee(cond, str)                                                                        \
-   {                                                                                                \
-     if (!(cond)) {                                                                                 \
-       std::cout << str << std::endl;                                                               \
-       abort();                                                                                   \
-     }                                                                                              \
-   }
+#define guarantee(cond, str)                                                                       \
+  {                                                                                                \
+    if (!(cond)) {                                                                                 \
+      std::cout << str << std::endl;                                                               \
+      abort();                                                                                     \
+    }                                                                                              \
+  }
 
 
 namespace HipTest {

projects/hip-tests/catch/include/hip_test_common.hh

@@ -36,6 +36,15 @@ THE SOFTWARE.
     }                                                                                              \
   }
 
+#define HIPRTC_CHECK(error)                                                                        \
+  {                                                                                                \
+    auto localError = error;                                                                       \
+    if (localError != HIPRTC_SUCCESS) {                                                            \
+      INFO("Error: " << hiprtcGetErrorString(localError) << " Code: " << localError << " Str: "    \
+                     << #error << " In File: " << __FILE__ << " At line: " << __LINE__);           \
+      REQUIRE(false);                                                                              \
+    }                                                                                              \
+  }
 // Although its assert, it will be evaluated at runtime
 #define HIP_ASSERT(x)                                                                              \
   { REQUIRE((x)); }

projects/hip-tests/catch/unit/CMakeLists.txt

@@ -3,6 +3,4 @@ add_subdirectory(deviceLib)
 add_subdirectory(stream)
 add_subdirectory(occupancy)
 add_subdirectory(device)
-
-# Disable Saxpy test temporarily to see if CI Passes
-# add_subdirectory(rtc)
+add_subdirectory(rtc)

projects/hip-tests/catch/unit/rtc/CMakeLists.txt

@@ -1,12 +1,10 @@
 # AMD Tests
-set(AMD_TEST_SRC
+set(TEST_SRC
     saxpy.cc
 )
 
-if(HIP_PLATFORM MATCHES "amd")
-    # Create shared lib of all tests
-    add_library(RTC SHARED EXCLUDE_FROM_ALL ${AMD_TEST_SRC})
+# Create shared lib of all tests
+add_library(RTC SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
 
-    # Add dependency on build_tests to build it on this custom target
-    add_dependencies(build_tests RTC)
-endif()
+# Add dependency on build_tests to build it on this custom target
+add_dependencies(build_tests RTC)

projects/hip-tests/catch/unit/rtc/saxpy.cc

@@ -15,20 +15,19 @@ static constexpr auto NUM_THREADS{128};
 static constexpr auto NUM_BLOCKS{32};
 
 static constexpr auto saxpy{
-R"(
+    R"(
 extern "C"
 __global__
 void saxpy(float a, float* x, float* y, float* out, size_t n)
 {
     size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
-        if (tid < n) {
-               out[tid] = a * x[tid] + y[tid];
-        }
-        
+    if (tid < n) {
+           out[tid] = a * x[tid] + y[tid];
+    }
 }
 )"};
 
-TEST_CASE("saxpy", "[hiprtc][saxpy]") {
+TEST_CASE("Unit_hiprtc_saxpy") {
   using namespace std;
   hiprtcProgram prog;
   hiprtcCreateProgram(&prog,       // prog
@@ -37,34 +36,43 @@ TEST_CASE("saxpy", "[hiprtc][saxpy]") {
                       0, nullptr, nullptr);
   hipDeviceProp_t props;
   int device = 0;
-  hipGetDeviceProperties(&props, device);
+  HIP_CHECK(hipGetDeviceProperties(&props, device));
+#ifdef __HIP_PLATFORM_AMD__
   std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
+#else
+  std::string sarg = std::string("--fmad=false");
+#endif
   const char* options[] = {sarg.c_str()};
   hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
   size_t logSize;
-  hiprtcGetProgramLogSize(prog, &logSize);
+  HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
   if (logSize) {
     string log(logSize, '\0');
-    hiprtcGetProgramLog(prog, &log[0]);
-
+    HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
     std::cout << log << '\n';
   }
   REQUIRE(compileResult == HIPRTC_SUCCESS);
   size_t codeSize;
-  hiprtcGetCodeSize(prog, &codeSize);
+  HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
 
   vector<char> code(codeSize);
-  hiprtcGetCode(prog, code.data());
+  HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
 
-  hiprtcDestroyProgram(&prog);
+  HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
+
+  // Do hip malloc first so that we donot need to do a cuInit manually before calling hipModule APIs
+  size_t n = NUM_THREADS * NUM_BLOCKS;
+  size_t bufferSize = n * sizeof(float);
+
+  float *dX, *dY, *dOut;
+  HIP_CHECK(hipMalloc(&dX, bufferSize));
+  HIP_CHECK(hipMalloc(&dY, bufferSize));
+  HIP_CHECK(hipMalloc(&dOut, bufferSize));
 
   hipModule_t module;
   hipFunction_t kernel;
-  hipModuleLoadData(&module, code.data());
-  hipModuleGetFunction(&kernel, module, "saxpy");
-
-  size_t n = NUM_THREADS * NUM_BLOCKS;
-  size_t bufferSize = n * sizeof(float);
+  HIP_CHECK(hipModuleLoadData(&module, code.data()));
+  HIP_CHECK(hipModuleGetFunction(&kernel, module, "saxpy"));
 
   float a = 5.1f;
   unique_ptr<float[]> hX{new float[n]};
@@ -75,18 +83,14 @@ TEST_CASE("saxpy", "[hiprtc][saxpy]") {
     hY[i] = static_cast<float>(i * 2);
   }
 
-  hipDeviceptr_t dX, dY, dOut;
-  hipMalloc(&dX, bufferSize);
-  hipMalloc(&dY, bufferSize);
-  hipMalloc(&dOut, bufferSize);
-  hipMemcpyHtoD(dX, hX.get(), bufferSize);
-  hipMemcpyHtoD(dY, hY.get(), bufferSize);
+  HIP_CHECK(hipMemcpy(dX, hX.get(), bufferSize, hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(dY, hY.get(), bufferSize, hipMemcpyHostToDevice));
 
   struct {
     float a_;
-    hipDeviceptr_t b_;
-    hipDeviceptr_t c_;
-    hipDeviceptr_t d_;
+    float* b_;
+    float* c_;
+    float* d_;
     size_t e_;
   } args{a, dX, dY, dOut, n};
 
@@ -95,17 +99,18 @@ TEST_CASE("saxpy", "[hiprtc][saxpy]") {
                     HIP_LAUNCH_PARAM_END};
 
   hipModuleLaunchKernel(kernel, NUM_BLOCKS, 1, 1, NUM_THREADS, 1, 1, 0, nullptr, nullptr, config);
-  hipMemcpyDtoH(hOut.get(), dOut, bufferSize);
+
+  HIP_CHECK(hipMemcpy(hOut.get(), dOut, bufferSize, hipMemcpyDeviceToHost));
+
+  hipFree(dX);
+  hipFree(dY);
+  hipFree(dOut);
+
+  HIP_CHECK(hipModuleUnload(module));
 
   for (size_t i = 0; i < n; ++i) {
     INFO("For " << i << " Value: " << fabs(a * hX[i] + hY[i] - hOut[i])
                 << " with: " << (fabs(hOut[i] * 1.0f) * 1e-6));
     REQUIRE(fabs(a * hX[i] + hY[i] - hOut[i]) <= fabs(hOut[i]) * 1e-6);
   }
-
-  hipFree(dX);
-  hipFree(dY);
-  hipFree(dOut);
-
-  hipModuleUnload(module);
 }