main.hip 1.1 KB

12345678910111213141516171819202122232425262728293031323334353637383940414243444546
  1. #include <cmath>
  2. #include <memory>
  3. #include <stdexcept>
  4. #include <hip/hip_fp16.h>
  5. #include <hip/hip_runtime.h>
  6. #include <math.h>
  7. namespace {
  8. template <class T, class F>
  9. __global__ void global_entry_point(F f, T* out)
  10. {
  11. *out = f();
  12. }
  13. template <class T, class F>
  14. bool verify(F f, T expected)
  15. {
  16. std::unique_ptr<T> cpu_T(new T);
  17. T* gpu_T = nullptr;
  18. if (hipMalloc((void**)&gpu_T, sizeof(T)) != hipSuccess) {
  19. return false;
  20. }
  21. bool result = true;
  22. hipLaunchKernelGGL(global_entry_point, 1, 1, 0, 0, f, gpu_T);
  23. result = hipMemcpy(cpu_T.get(), gpu_T, sizeof(T), hipMemcpyDeviceToHost) ==
  24. hipSuccess &&
  25. result;
  26. result = hipFree(gpu_T) == hipSuccess && result;
  27. result = *cpu_T == expected && result;
  28. return result;
  29. }
  30. }
  31. int main(int argc, char** argv)
  32. {
  33. bool valid = verify([] __device__() { return std::round(1.4f); }, 1.0f);
  34. valid &= verify([] __device__() { return max<_Float16>(1.0f, 2.0f); }, 2.0f);
  35. valid &= verify([] __device__() { return min<_Float16>(1.0f, 2.0f); }, 1.0f);
  36. if (valid) {
  37. return 0;
  38. } else {
  39. return 1;
  40. }
  41. }