Role of unit tests

[tpkessler]

Hi! I'm the maintainer of ROCm community packages for Arch Linux (rocm-arch). Currently, we're adding unit tests to our package build scripts.

Running the HIP unit test (with HIP_CATCH_TEST=1) many tests fail on my Vega 56 and some, in particular stream- and graph-related tests, even cause the GPU driver to crash. I've encountered similar problems with KFD unit tests (see https://github.com/RadeonOpenCompute/ROCT-Thunk-Interface/issues/76). In contrast to the linked problem which required to blacklist only three tests, filtering all bad HIP tests would be much harder.

Therefore, I would like to know if the unit tests are designed to be compiled and run by an end user? If not, what would be the best way to validate a (basic) HIP installation? I'm only aware of the HIP examples repository that includes checks for parts of the HIP functionalities.

[jatinx]

I expect some tests to fail, vega 56 is quite old and some new tests might not be written with that in mind.

can you share some failed tests names.

Also to verify that HIP is installed correctly, you can run this example. It is very primitive and does not do a thorough check on any major HIP features, it just checks 3 things:

1. can host allocate device memory
2. can host copy memory to and from GPU
3. can host launch kernels on gpu.

#include <iostream>
#include <cassert>

#include <hip/hip_runtime.h>

#define HIP_CHECK(call_)                                                                           \
  {                                                                                                \
    auto res_ = call_;                                                                             \
    if (res_ != hipSuccess) {                                                                      \
      std::cout << "Failed in: " << #call_ << " with error: " << res_ << std::endl;                \
      return -1;                                                                                   \
    }                                                                                              \
  }

__global__ void kernel(int* set) {
  if (*set == 20) {
    *set = 10;
  }
}

int main() {
  int* ptr{nullptr};
  HIP_CHECK(hipMalloc(&ptr, sizeof(int)));
  int val = 20;
  HIP_CHECK(hipMemcpy(ptr, &val, sizeof(int), hipMemcpyHostToDevice));
  kernel<<<1, 1>>>(ptr);
  int res = 0;
  HIP_CHECK(hipMemcpy(&res, ptr, sizeof(int), hipMemcpyDeviceToHost));
  if (res == 10) {
    std::cout << "HIP is working fine!" << std::endl;
  } else {
    std::cout << "HIP is NOT working fine!" << std::endl;
  }
}

[abhimeda]

@tpkessler Hi, is your issue resolved on the latest HIP? If so can we close this ticket?