fixed hipcc for new compiler flags

Change-Id: I49ec059be20ff26b7482c84d91ab7a43826c6a8d

hipamd/bin/hipcc

@@ -134,28 +134,6 @@ if ($HIP_PLATFORM eq "hcc") {
     }
     $HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt";
 
-    $ENV{HCC_EXTRA_LIBRARIES}="$HIP_PATH/lib/hip_hc.ll\n";
-
-    # Handle ROCm target platform
-    if ($target_gfx701 eq 1) {
-        $HIPLDFLAGS .= " --amdgpu-target=gfx701";
-    }
-    if ($target_gfx801 eq 1) {
-        $HIPLDFLAGS .= " --amdgpu-target=gfx801";
-    }
-    if ($target_gfx802 eq 1) {
-        $HIPLDFLAGS .= " --amdgpu-target=gfx802";
-    }
-    if ($target_gfx803 eq 1) {
-        $HIPLDFLAGS .= " --amdgpu-target=gfx803";
-        $ENV{HIP_HC_IR_GFX803}="$HIP_PATH/lib/hip_hc_gfx803.ll\n";
-    }
-    if ($target_gfx701 eq 0 and $target_gfx801 eq 0 and $target_gfx802 eq 0 and $target_gfx803 eq 0) 
-    {
-        $HIPLDFLAGS .= " --amdgpu-target=gfx701 --amdgpu-target=gfx801 --amdgpu-target=gfx802 --amdgpu-target=gfx803";
-        $ENV{HIP_HC_IR_GFX803}="$HIP_PATH/lib/hip_hc_gfx803.ll\n";
-    }
-
     # 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.
     if ($HIP_ATP_MARKER) {
@@ -352,6 +330,32 @@ foreach $arg (@ARGV)
     $toolArgs .= " $arg" unless $swallowArg;
 }
 
+if($HIP_PLATFORM eq "hcc"){
+
+    $ENV{HCC_EXTRA_LIBRARIES}="$HIP_PATH/lib/hip_hc.ll\n";
+
+    # Handle ROCm target platform
+    if ($target_gfx701 eq 1) {
+        $HIPLDFLAGS .= " --amdgpu-target=gfx701";
+    }
+    if ($target_gfx801 eq 1) {
+        $HIPLDFLAGS .= " --amdgpu-target=gfx801";
+    }
+    if ($target_gfx802 eq 1) {
+        $HIPLDFLAGS .= " --amdgpu-target=gfx802";
+    }
+    if ($target_gfx803 eq 1) {
+        $HIPLDFLAGS .= " --amdgpu-target=gfx803";
+        $ENV{HIP_HC_IR_GFX803}="$HIP_PATH/lib/hip_hc_gfx803.ll\n";
+    }
+    if ($target_gfx701 eq 0 and $target_gfx801 eq 0 and $target_gfx802 eq 0 and $target_gfx803 eq 0) 
+    {
+        $HIPLDFLAGS .= " --amdgpu-target=gfx701 --amdgpu-target=gfx801 --amdgpu-target=gfx802 --amdgpu-target=gfx803";
+        $ENV{HIP_HC_IR_GFX803}="$HIP_PATH/lib/hip_hc_gfx803.ll\n";
+    }
+
+}
+
 if ($hasC and $HIP_PLATFORM eq 'nvcc') {
     $HIPCXXFLAGS .= " -x cu";
 }

hipamd/tests/src/deviceLib/hipTestHalf.cpp

@@ -21,12 +21,51 @@ THE SOFTWARE.
 #include <hip/hip_fp16.h>
 #include "hip/hip_runtime_api.h"
 
-__global__ void halfMath(hipLaunchParm lp, half *A, half *B, half *C) {
+#define LEN 64
+#define HALF_SIZE 64*sizeof(__half)
+#define HALF2_SIZE 64*sizeof(__half2)
+
+__global__ void __halfMath(hipLaunchParm lp, __half *A, __half *B, __half *C) {
   int tx = hipThreadIdx_x;
   __half a = A[tx];
-  
+  __half b = B[tx];
+  __half c = C[tx];
+  c = __hadd(a, c);
+  c = __hadd_sat(b, c);
+  c = __hfma(a, c, b);
+  c = __hfma_sat(b, c, a);
+  c = __hsub(a, c);
+  c = __hsub_sat(b, c);
+  c = __hmul(a, c);
+  c = __hmul_sat(b, c);
+  c = hdiv(a, c);
+}
+
+__global__ void __half2Math(hipLaunchParm lp, __half2 *A, __half2 *B, __half2 *C) {
+  int tx = hipThreadIdx_x;
+  __half2 a = A[tx];
+  __half2 b = B[tx];
+  __half2 c = C[tx];
+  c = __hadd2(a, c);
+  c = __hadd2_sat(b, c);
+  c = __hfma2(a, c, b);
+  c = __hfma2_sat(b, c, a);
+  c = __hsub2(a, c);
+  c = __hsub2_sat(b, c);
+  c = __hmul2(a, c);
+  c = __hmul2_sat(b, c);
 }
 
 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);
 
 }