AMDGPU: InstCombine amdgcn.sqrt.f16 to sqrt.f16

[arsenm]

Summary

There's nothing special about f16 sqrt handling.

[arsenm]

ping

[rampitec]

Matt, please check failed clang test.

[arsenm]

Fix clang test

[rampitec]

LGTM

[arsenm]

61c8af67924b02c8f2cf871439c24650a0207f29

clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

 // CHECK-LABEL: @test_rcp_f16
 // CHECK: call half @llvm.amdgcn.rcp.f16
 void test_rcp_f16(global half* out, half a)
 {
   *out = __builtin_amdgcn_rcph(a);
 }
 
 // CHECK-LABEL: @test_sqrt_f16
-// CHECK: call half @llvm.amdgcn.sqrt.f16
+// CHECK: call half @llvm.sqrt.f16
 void test_sqrt_f16(global half* out, half a)
 {
   *out = __builtin_amdgcn_sqrth(a);
 }
 
 // CHECK-LABEL: @test_rsq_f16
 // CHECK: call half @llvm.amdgcn.rsq.f16
 void test_rsq_f16(global half* out, half a)

llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp

case Intrinsic::amdgcn_rsq: {
 
     // TODO: Move to ConstantFolding/InstSimplify?
     if (isa<UndefValue>(Src)) {
       Type *Ty = II.getType();
       auto *QNaN = ConstantFP::get(Ty, APFloat::getQNaN(Ty->getFltSemantics()));
       return IC.replaceInstUsesWith(II, QNaN);
     }
 
+    // f16 amdgcn.sqrt is identical to regular sqrt.
+    if (IID == Intrinsic::amdgcn_sqrt && Src->getType()->isHalfTy()) {
+      Function *NewDecl = Intrinsic::getDeclaration(
+          II.getModule(), Intrinsic::sqrt, {II.getType()});
+      II.setCalledFunction(NewDecl);
+      return &II;
+    }
+
     break;
   }
   case Intrinsic::amdgcn_log:
   case Intrinsic::amdgcn_exp2: {
     const bool IsLog = IID == Intrinsic::amdgcn_log;
     const bool IsExp = IID == Intrinsic::amdgcn_exp2;
     Value *Src = II.getArgOperand(0);
     Type *Ty = II.getType();

llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll

 ; CHECK-NEXT:    ret double 0x7FF8000000000000
 ;
   %val = call double @llvm.amdgcn.sqrt.f64(double undef) nounwind readnone
   ret double %val
 }
 
 define half @test_constant_fold_sqrt_f16_0() nounwind {
 ; CHECK-LABEL: @test_constant_fold_sqrt_f16_0(
-; CHECK-NEXT:    [[VAL:%.*]] = call half @llvm.amdgcn.sqrt.f16(half 0xH0000) #[[ATTR15:[0-9]+]]
-; CHECK-NEXT:    ret half [[VAL]]
+; CHECK-NEXT:    ret half 0xH0000
 ;
   %val = call half @llvm.amdgcn.sqrt.f16(half 0.0) nounwind readnone
   ret half %val
 }
 
 define float @test_constant_fold_sqrt_f32_0() nounwind {
 ; CHECK-LABEL: @test_constant_fold_sqrt_f32_0(
-; CHECK-NEXT:    [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float 0.000000e+00) #[[ATTR15]]
+; CHECK-NEXT:    [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float 0.000000e+00) #[[ATTR15:[0-9]+]]
 ; CHECK-NEXT:    ret float [[VAL]]
 ;
   %val = call float @llvm.amdgcn.sqrt.f32(float 0.0) nounwind readnone
   ret float %val
 }
 
 define double @test_constant_fold_sqrt_f64_0() nounwind {
 ; CHECK-LABEL: @test_constant_fold_sqrt_f64_0(
 ; CHECK-NEXT:    [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double 0.000000e+00) #[[ATTR15]]
 ; CHECK-NEXT:    ret double [[VAL]]
 ;
   %val = call double @llvm.amdgcn.sqrt.f64(double 0.0) nounwind readnone
   ret double %val
 }
 
 define half @test_constant_fold_sqrt_f16_neg0() nounwind {
 ; CHECK-LABEL: @test_constant_fold_sqrt_f16_neg0(
-; CHECK-NEXT:    [[VAL:%.*]] = call half @llvm.amdgcn.sqrt.f16(half 0xH8000) #[[ATTR15]]
-; CHECK-NEXT:    ret half [[VAL]]
+; CHECK-NEXT:    ret half 0xH8000
 ;
   %val = call half @llvm.amdgcn.sqrt.f16(half -0.0) nounwind readnone
   ret half %val
 }
 
 define float @test_constant_fold_sqrt_f32_neg0() nounwind {
 ; CHECK-LABEL: @test_constant_fold_sqrt_f32_neg0(
 ; CHECK-NEXT:    [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float -0.000000e+00) #[[ATTR15]]

 ; CHECK-LABEL: @test_constant_fold_sqrt_neg1(
 ; CHECK-NEXT:    [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double -1.000000e+00)
 ; CHECK-NEXT:    ret double [[VAL]]
 ;
   %val = call double @llvm.amdgcn.sqrt.f64(double -1.0)
   ret double %val
 }
 
+define half @test_amdgcn_sqrt_f16(half %arg) {
+; CHECK-LABEL: @test_amdgcn_sqrt_f16(
+; CHECK-NEXT:    [[VAL:%.*]] = call half @llvm.sqrt.f16(half [[ARG:%.*]])
+; CHECK-NEXT:    ret half [[VAL]]
+;
+  %val = call half @llvm.amdgcn.sqrt.f16(half %arg)
+  ret half %val
+}
+
+define half @test_amdgcn_sqrt_f16_flags(half %arg) {
+; CHECK-LABEL: @test_amdgcn_sqrt_f16_flags(
+; CHECK-NEXT:    [[VAL:%.*]] = call nnan half @llvm.sqrt.f16(half [[ARG:%.*]])
+; CHECK-NEXT:    ret half [[VAL]]
+;
+  %val = call nnan half @llvm.amdgcn.sqrt.f16(half %arg)
+  ret half %val
+}
+
+define float @test_amdgcn_sqrt_f32(float %arg) {
+; CHECK-LABEL: @test_amdgcn_sqrt_f32(
+; CHECK-NEXT:    [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float [[ARG:%.*]])
+; CHECK-NEXT:    ret float [[VAL]]
+;
+  %val = call float @llvm.amdgcn.sqrt.f32(float %arg)
+  ret float %val
+}
+
+define double @test_amdgcn_sqrt_f64(double %arg) {
+; CHECK-LABEL: @test_amdgcn_sqrt_f64(
+; CHECK-NEXT:    [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double [[ARG:%.*]])
+; CHECK-NEXT:    ret double [[VAL]]
+;
+  %val = call double @llvm.amdgcn.sqrt.f64(double %arg)
+  ret double %val
+}
+
 ; --------------------------------------------------------------------
 ; llvm.amdgcn.rsq
 ; --------------------------------------------------------------------
 
 declare float @llvm.amdgcn.rsq.f32(float) nounwind readnone
 
 define float @test_constant_fold_rsq_f32_undef() nounwind {
 ; CHECK-LABEL: @test_constant_fold_rsq_f32_undef(

llvm/test/Transforms/InstCombine/AMDGPU/rcp-contract-rsq.ll

;
   %rsq = call contract half @llvm.amdgcn.rcp.f16(half %sqrt)
   ret half %rsq
 }
 
 ; contract required on both calls
 define half @amdgcn_rcp_amdgcn_sqrt_f16_missing_contract0(half %x) {
 ; CHECK-LABEL: define half @amdgcn_rcp_amdgcn_sqrt_f16_missing_contract0
 ; CHECK-SAME: (half [[X:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT:    [[SQRT:%.*]] = call half @llvm.amdgcn.sqrt.f16(half [[X]])
+; CHECK-NEXT:    [[SQRT:%.*]] = call half @llvm.sqrt.f16(half [[X]])
 ; CHECK-NEXT:    [[RSQ:%.*]] = call contract half @llvm.amdgcn.rcp.f16(half [[SQRT]])
 ; CHECK-NEXT:    ret half [[RSQ]]
 ;
   %sqrt = call half @llvm.amdgcn.sqrt.f16(half %x)
   %rsq = call contract half @llvm.amdgcn.rcp.f16(half %sqrt)
   ret half %rsq
 }
 
 ; contract required on both calls
 define half @amdgcn_rcp_amdgcn_sqrt_f16_missing_contract1(half %x) {
 ; CHECK-LABEL: define half @amdgcn_rcp_amdgcn_sqrt_f16_missing_contract1
 ; CHECK-SAME: (half [[X:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT:    [[SQRT:%.*]] = call contract half @llvm.amdgcn.sqrt.f16(half [[X]])
+; CHECK-NEXT:    [[SQRT:%.*]] = call contract half @llvm.sqrt.f16(half [[X]])
 ; CHECK-NEXT:    [[RSQ:%.*]] = call half @llvm.amdgcn.rcp.f16(half [[SQRT]])
 ; CHECK-NEXT:    ret half [[RSQ]]
 ;
   %sqrt = call contract half @llvm.amdgcn.sqrt.f16(half %x)
   %rsq = call half @llvm.amdgcn.rcp.f16(half %sqrt)
   ret half %rsq
 }
 
 ; Reject from multiple uses of sqrt
 define half @amdgcn_rcp_amdgcn_sqrt_f16_contract_multi_use(half %x, ptr %ptr) {
 ; CHECK-LABEL: define half @amdgcn_rcp_amdgcn_sqrt_f16_contract_multi_use
 ; CHECK-SAME: (half [[X:%.*]], ptr [[PTR:%.*]]) #[[ATTR1]] {
-; CHECK-NEXT:    [[SQRT:%.*]] = call contract half @llvm.amdgcn.sqrt.f16(half [[X]])
+; CHECK-NEXT:    [[SQRT:%.*]] = call contract half @llvm.sqrt.f16(half [[X]])
 ; CHECK-NEXT:    store half [[SQRT]], ptr [[PTR]], align 2
 ; CHECK-NEXT:    [[RSQ:%.*]] = call contract half @llvm.amdgcn.rcp.f16(half [[SQRT]])
 ; CHECK-NEXT:    ret half [[RSQ]]
 ;
   %sqrt = call contract half @llvm.amdgcn.sqrt.f16(half %x)
   store half %sqrt, ptr %ptr
   %rsq = call contract half @llvm.amdgcn.rcp.f16(half %sqrt)
   ret half %rsq