From c3ba3095626ca593285bf17b201664c2cb1d0c3e Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Wed, 15 Nov 2017 18:52:59 +0530
Subject: [PATCH] Fixed half2 issue on gfx900

[ROCm/clr commit: 20947f80fba711bc051742a06960131d46653628]
---
 projects/clr/hipamd/bin/hipcc                 |  1 +
 .../tests/src/deviceLib/hipTestHalf.cpp       | 33 +++++++++++--------
 2 files changed, 21 insertions(+), 13 deletions(-)

diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc
index bbe079db41..b991bd40aa 100755
--- a/projects/clr/hipamd/bin/hipcc
+++ b/projects/clr/hipamd/bin/hipcc
@@ -439,6 +439,7 @@ if($HIP_PLATFORM eq "hcc"){
     if ($target_gfx900 eq 1) {
         $HIPLDFLAGS .= " --amdgpu-target=gfx900";
         $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX900__=1 ";
+        $ENV{HCC_EXTRA_LIBRARIES_GFX900}="$HIP_PATH/lib/hip_hc_gfx803.ll\n";
     }
 }
 
diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestHalf.cpp
index 94a3882eea..873dd23805 100644
--- a/projects/clr/hipamd/tests/src/deviceLib/hipTestHalf.cpp
+++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestHalf.cpp
@@ -19,13 +19,14 @@ THE SOFTWARE.
 
 #include <iostream>
 #include <hip/hip_fp16.h>
-#include "hip/hip_runtime_api.h"
+#include "hip/hip_runtime.h"
+#include "test_common.h"
 
 #define LEN 64
 #define HALF_SIZE 64*sizeof(__half)
 #define HALF2_SIZE 64*sizeof(__half2)
 
-#if __HIP_ARCH_GFX803__ > 0
+#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX803__
 
 __global__ void __halfMath(hipLaunchParm lp, __half *A, __half *B, __half *C) {
   int tx = hipThreadIdx_x;
@@ -61,15 +62,21 @@ __global__ void __half2Math(hipLaunchParm lp, __half2 *A, __half2 *B, __half2 *C
 #endif
 
 int main(){
-  __half *A, *B, *C;
-  hipMalloc(&A, HALF_SIZE);
-  hipMalloc(&B, HALF_SIZE);
-  hipMalloc(&C, HALF_SIZE);
-  hipLaunchKernel(__halfMath, dim3(1,1,1), dim3(LEN,1,1), 0, 0, A, B, C);
-  __half2 *A2, *B2, *C2;
-  hipMalloc(&A, HALF2_SIZE);
-  hipMalloc(&B, HALF2_SIZE);
-  hipMalloc(&C, HALF2_SIZE);
-  hipLaunchKernel(__half2Math, dim3(1,1,1), dim3(LEN,1,1), 0, 0, A2, B2, C2);
-
+    __half *A, *B, *C;
+    hipMalloc(&A, HALF_SIZE);
+    hipMalloc(&B, HALF_SIZE);
+    hipMalloc(&C, HALF_SIZE);
+    hipLaunchKernel(__halfMath, dim3(1,1,1), dim3(LEN,1,1), 0, 0, A, B, C);
+    hipFree(A);
+    hipFree(B);
+    hipFree(C);
+    __half2 *A2, *B2, *C2;
+    hipMalloc(&A2, HALF2_SIZE);
+    hipMalloc(&B2, HALF2_SIZE);
+    hipMalloc(&C2, HALF2_SIZE);
+    hipLaunchKernel(__half2Math, dim3(1,1,1), dim3(LEN,1,1), 0, 0, A2, B2, C2);
+    hipFree(A2);
+    hipFree(B2);
+    hipFree(C2);
+    passed();
 }