diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 2a16b5702a..944fc04135 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -266,7 +266,7 @@ foreach $arg (@ARGV) $target_gfx802 = 1; } if($arg eq '--amdgpu-target=gfx803') - { + { $target_gfx803 = 1; } if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0)) @@ -337,18 +337,22 @@ if($HIP_PLATFORM eq "hcc"){ # Handle ROCm target platform if ($target_gfx701 eq 1) { $HIPLDFLAGS .= " --amdgpu-target=gfx701"; + $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX701__=1 "; } if ($target_gfx801 eq 1) { $HIPLDFLAGS .= " --amdgpu-target=gfx801"; + $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX801__=1 "; } if ($target_gfx802 eq 1) { $HIPLDFLAGS .= " --amdgpu-target=gfx802"; + $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX802__=1 "; } if ($target_gfx803 eq 1) { $HIPLDFLAGS .= " --amdgpu-target=gfx803"; + $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX803__=1 "; $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) + 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"; diff --git a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp index 997c445a10..94a3882eea 100644 --- a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp +++ b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp @@ -25,6 +25,8 @@ THE SOFTWARE. #define HALF_SIZE 64*sizeof(__half) #define HALF2_SIZE 64*sizeof(__half2) +#if __HIP_ARCH_GFX803__ > 0 + __global__ void __halfMath(hipLaunchParm lp, __half *A, __half *B, __half *C) { int tx = hipThreadIdx_x; __half a = A[tx]; @@ -56,6 +58,8 @@ __global__ void __half2Math(hipLaunchParm lp, __half2 *A, __half2 *B, __half2 *C c = __hmul2_sat(b, c); } +#endif + int main(){ __half *A, *B, *C; hipMalloc(&A, HALF_SIZE);