From 0b445d68de8fb005e9fef31ba538ec5ee0846308 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 13 Dec 2016 20:06:56 -0600 Subject: [PATCH] disabled compiler flag hcc 4.0 for half support Change-Id: I32175113f4c05d43310b3a05c2a14e12f6d48b09 [ROCm/hip commit: ed39a7f43bcc284997593e4acf12faa2a602b617] --- .../hip/include/hip/hcc_detail/hip_fp16.h | 22 +++++++++---------- projects/hip/src/hip_fp16.cpp | 3 --- 2 files changed, 10 insertions(+), 15 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index 445df78eb4..735f915bd2 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h @@ -27,7 +27,16 @@ THE SOFTWARE. #define __CLANG_VERSION__ __clang_major__ * 10 + __clang_minor__ -#if __CLANG_VERSION__ == 35 +#ifdef HIP_HALF_HW_SUPPORT + +typedef __fp16 __half; +extern "C" __half __hip_hadd_clang40_gfx803(__half a, __half b); + +__device__ inline __half __hadd(__half a, __half b){ + return __hip_hadd_clang40_gfx803(a, b); +} + +#else typedef struct{ unsigned x: 16; @@ -178,16 +187,5 @@ __device__ __half2 __lowhigh2highlow(const __half2 a); __device__ __half2 __low2half2(const __half2 a, const __half2 b); -#endif - -#if __CLANG_VERSION__ == 40 - -typedef __fp16 __half; -extern "C" __half __hip_hadd_clang40_gfx803(__half a, __half b); - -__device__ inline __half __hadd(__half a, __half b){ - return __hip_hadd_clang40_gfx803(a, b); -} - #endif #endif diff --git a/projects/hip/src/hip_fp16.cpp b/projects/hip/src/hip_fp16.cpp index 5d01f73cf7..63d91eb107 100644 --- a/projects/hip/src/hip_fp16.cpp +++ b/projects/hip/src/hip_fp16.cpp @@ -22,7 +22,6 @@ THE SOFTWARE. #include"hip/hip_fp16.h" -#if __CLANG_VERSION__ == 35 static const unsigned sign_val = 0x8000; static const __half __half_value_one_float = {0x3C00}; @@ -375,5 +374,3 @@ __device__ __half2 __low2half2(const __half2 a, const __half2 b){ return {a.q, b.q}; } -#endif -