From d06509f680d83e86c32fc17fbea101302adebcf2 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 13 Oct 2016 14:16:48 -0500 Subject: [PATCH] added compiler flag for polaris Change-Id: Ib14c14c0618982ac7b48f5bc704c04b54ff40ed9 [ROCm/hip commit: 90a71c4be476b60508f7d2d83370e528fb65ac58] --- projects/hip/bin/hipcc | 3 +++ projects/hip/include/hip/hcc_detail/hip_fp16.h | 7 ------- projects/hip/src/hip_fp16.cpp | 8 +++++++- 3 files changed, 10 insertions(+), 8 deletions(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 9de4cdf732..a3d70c46d7 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -119,6 +119,9 @@ if ($HIP_PLATFORM eq "hcc") { if ($ROCM_TARGET eq "hawaii") { $HIPLDFLAGS .= " -amdgpu-target=AMD:AMDGPU:7:0:1"; } + if ($ROCM_TARGET eq "polaris") { + $HIPLDFLAGS .= " -amdgpu-target=AMD:AMDGPU:8:0:3"; + } # Add trace marker library: # TODO - once we cleanly separate the HIP API headers from HIP library headers this logic should move to CMakebuild option - apps do not need to see the marker library. diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index 9c7b3a6646..7558fd348d 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h @@ -34,13 +34,6 @@ typedef struct __attribute__((aligned(4))){ typedef __half half; typedef __half2 half2; -typedef struct{ - union{ - float f; - unsigned u; - }; -} struct_float; - /* Arithmetic functions */ diff --git a/projects/hip/src/hip_fp16.cpp b/projects/hip/src/hip_fp16.cpp index 5b179b9ba5..a1257b2bfb 100644 --- a/projects/hip/src/hip_fp16.cpp +++ b/projects/hip/src/hip_fp16.cpp @@ -25,6 +25,13 @@ static const __half __half_value_zero_float = {0x0}; static const unsigned __half_pos_inf = 0x7C00; static const unsigned __half_neg_inf = 0xFC00; +typedef struct{ + union{ + float f; + unsigned u; + }; +} struct_float; + static __device__ float cvt_half_to_float(__half a){ struct_float ret = {0}; if(a.x == 0){ @@ -362,4 +369,3 @@ __device__ __half2 __lowhigh2highlow(const __half2 a){ __device__ __half2 __low2half2(const __half2 a, const __half2 b){ return {a.q, b.q}; } -