| #include <stdexcept> | |
| #include <cmath> | |
| #include <math.h> | |
| #include <memory> | |
| #include <hip/hip_runtime.h> | |
| #include <hip/hip_fp16.h> | |
| namespace { | |
| template<class T, class F> | |
| __global__ void global_entry_point(F f, T *out) { | |
| *out = f(); | |
| } | |
| template <class T, class F> | |
| bool verify(F f, T expected) | |
| { | |
| std::unique_ptr<T> cpu_T(new T); | |
| T* gpu_T = nullptr; | |
| hipMalloc((void**)&gpu_T, sizeof(T)); | |
| hipLaunchKernelGGL(global_entry_point, 1, 1, 0, 0, f, gpu_T); | |
| hipMemcpy(cpu_T.get(), gpu_T, sizeof(T), hipMemcpyDeviceToHost); | |
| hipFree(gpu_T); | |
| return (*cpu_T == expected); | |
| } | |
| } | |
| int main(int argc, char** argv) | |
| { | |
| bool valid = verify([]__device__(){ return std::round(1.4f); }, 1.0f); | |
| valid &= verify([]__device__(){ return max<_Float16>(1.0f, 2.0f); }, 2.0f); | |
| valid &= verify([]__device__(){ return min<_Float16>(1.0f, 2.0f); }, 1.0f); | |
| if (valid) { | |
| return 0; | |
| } else { | |
| return 1; | |
| } | |
| } |