diff options
| author |  Deven Desai <deven.desai.amd@gmail.com> | 2020-12-10 14:51:13 +0000 |
|---|---|---|
| committer | Deven Desai <deven.desai.amd@gmail.com> | 2020-12-11 16:14:57 +0000 |
| commit | f3d2ea48f54f1770681de78786bc5b102af43cbd (patch) | |
| tree | 93c116eef7d5942f96ddc706bc72998f483d02ad /Eigen/src/Core/arch/GPU | |
| parent | c7eb3a74cb4d1d2e3d2db5b5c69c1ab27f47d76b (diff) | |
Fix for broken ROCm/HIP Support
The following commit introduced a breakage in ROCm/HIP support for Eigen.
https://gitlab.com/libeigen/eigen/-/commit/5ec4907434742d4555df4aa708b665868b88f3b4#1958e65719641efe5483abc4ce0b61806270f6f3_525_517
```
Building HIPCC object test/CMakeFiles/gpu_basic.dir/gpu_basic_generated_gpu_basic.cu.o
In file included from /home/rocm-user/eigen/test/gpu_basic.cu:20:
In file included from /home/rocm-user/eigen/test/main.h:356:
In file included from /home/rocm-user/eigen/Eigen/QR:11:
In file included from /home/rocm-user/eigen/Eigen/Core:222:
/home/rocm-user/eigen/Eigen/src/Core/arch/GPU/PacketMath.h:556:10: error: use of undeclared identifier 'half2half2'; did you mean '__half2half2'?
return half2half2(from);
^~~~~~~~~~
__half2half2
/opt/rocm/hip/include/hip/hcc_detail/hip_fp16.h:547:21: note: '__half2half2' declared here
__half2 __half2half2(__half x)
^
1 error generated when compiling for gfx900.
```
The cause seems to be a copy-paster error, and the fix is trivial
Diffstat (limited to 'Eigen/src/Core/arch/GPU')
| -rw-r--r-- | Eigen/src/Core/arch/GPU/PacketMath.h | 4 |
1 files changed, 1 insertions, 3 deletions
diff --git a/Eigen/src/Core/arch/GPU/PacketMath.h b/Eigen/src/Core/arch/GPU/PacketMath.h index 83bd551a0..c16f95e7f 100644 --- a/Eigen/src/Core/arch/GPU/PacketMath.h +++ b/Eigen/src/Core/arch/GPU/PacketMath.h @@ -552,9 +552,7 @@ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE __half get_half2_high(const half2& a) { template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pset1<half2>(const Eigen::half& from) { -#if defined(EIGEN_HIP_DEVICE_COMPILE) - return half2half2(from); -#elif defined(EIGEN_CUDA_ARCH) +#if defined(EIGEN_GPU_COMPILE_PHASE) return __half2half2(from); #else const float f = __half2float(from); |