[SYCL] Fix the case of folded const expr when replacing a call to the fpbuiltin intrinsic (#17244)

This PR adds a sanity check to the FPBuiltinFnSelection pass to ensure
that copyFastMathFlags is not called for a folded constant. This is to
fix the case when ConstantFolder folds original arguments to a constant,
meaning that we have no instruction anymore.

From the LLVM IR perspective the reproducer is as simple as
```
define spir_kernel void @foo() {
entry:
  %r = tail call float @llvm.fpbuiltin.fdiv.f32(float 1.000000e+00, float 2.000000e+00)
  ret void
}
```

The corresponding SYCL code may look like, for instance:

```
#include <sycl/detail/core.hpp>

using namespace sycl;

int main() {
  const unsigned array_size = 4;
  range<1> numOfItems{array_size};
  queue deviceQueue;
  float *a;

  deviceQueue.submit([&](handler& cgh) {
    cgh.parallel_for<class KernelFdiv>(numOfItems,
    [=](id<1> wiID) {
      a[0] = .5f / .9f;
    });
  });

  return 0;
}
```

Having SYCL code, the crash may be reproduced either via LLVM IR
```
./bin/clang++ -S -fsycl -emit-llvm -fno-offload-fp32-prec-div fpbuiltin-issue.cpp -o fpbuiltin-issue.ll
```
or directly using the LLVM SPIR-V backend
```
./bin/clang++ -fsycl -fno-offload-fp32-prec-div fpbuiltin-issue.cpp -o fpbuiltin-issue-exe -fsycl-use-spirv-backend-for-spirv-gen
```

Author: VyacheslavLevytskyy

llvm/lib/Transforms/Scalar/FPBuiltinFnSelection.cpp

@@ -100,7 +100,10 @@ static bool replaceWithLLVMIR(FPBuiltinIntrinsic &BuiltinCall) {
     break;
   }
   BuiltinCall.replaceAllUsesWith(Replacement);
-  cast<Instruction>(Replacement)->copyFastMathFlags(&BuiltinCall);
+  // ConstantFolder may fold original arguments to a constant, meaning we might
+  // have no instruction anymore.
+  if (auto *ReplacementI = dyn_cast<Instruction>(Replacement))
+    ReplacementI->copyFastMathFlags(&BuiltinCall);
   LLVM_DEBUG(dbgs() << DEBUG_TYPE << ": Replaced call to `"
                     << BuiltinCall.getCalledFunction()->getName()
                     << "` with equivalent IR. \n `");

llvm/test/Transforms/FPBuiltin/lower-fpbuiltin-constexpr.ll

@@ -0,0 +1,14 @@
+; RUN: llc -mtriple=spirv64-unknown-unknown -o - %s
+; Test that codegen doesn't crash for fpbuiltin intrinsic lowering.
+
+define float @test() {
+entry:
+  %r = tail call float @llvm.fpbuiltin.fdiv.f32(float 1.000000e+00, float 2.000000e+00)
+  ret float %r
+}
+
+define spir_kernel void @foo() {
+entry:
+  %r = tail call float @test()
+  ret void
+}