diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index 7df674f57b..2a16b5702a 100755 --- a/hipamd/bin/hipcc +++ b/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"; } diff --git a/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/hipamd/tests/src/deviceLib/hipTestHalf.cpp index 2284148bb5..997c445a10 100644 --- a/hipamd/tests/src/deviceLib/hipTestHalf.cpp +++ b/hipamd/tests/src/deviceLib/hipTestHalf.cpp @@ -21,12 +21,51 @@ THE SOFTWARE. #include #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); }