[clang][SYCL] Disable force inlining of kernel call operator for FGPA (#8688)

In some cases force inlining can remove necessary FPGA metadata. This
happened with `ivdep` attribute attached to `do .. while(1)` loop.

Author: Fznamznon

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5147,7 +5147,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     // At -O0, disable the inlining for debugging purposes.
     if (!Args.hasFlag(options::OPT_fsycl_force_inline_kernel_lambda,
                       options::OPT_fno_sycl_force_inline_kernel_lambda,
-                      !DisableSYCLForceInlineKernelLambda))
+                      !DisableSYCLForceInlineKernelLambda &&
+                          !IsFPGASYCLOffloadDevice))
       CmdArgs.push_back("-fno-sycl-force-inline-kernel-lambda");
 
     // Add -ffine-grained-bitfield-accesses option. This will be added

clang/test/CodeGenSYCL/fpga-attr-do-while-loops.cpp

@@ -0,0 +1,26 @@
+// RUN: %clang++ -fsycl-device-only -fintelfpga -S %s -o - | FileCheck %s
+
+#include "Inputs/sycl.hpp"
+
+// This test checks that FPGA loop metadata is not lost due to optimizations.
+
+int main() {
+  sycl::queue q;
+  q.submit([&](sycl::handler &cgh) {
+    cgh.single_task<class ivdep>([=](){
+      int m =  16;
+      int i = 0;
+      int b = 1;
+      // CHECK: {!"llvm.loop.ivdep.enable"}
+      [[intel::ivdep]] do {
+        if (i >= m) {
+          break;
+        } else {
+          int b = b *2;
+        }
+        ++i;
+      } while (1);
+    });
+  });
+  return 0;
+}

clang/test/Driver/sycl.c

@@ -79,6 +79,7 @@
 // RUN: %clang_cl -### -fsycl-device-only -Od %s 2>&1 | FileCheck %s --check-prefix=CHECK-NOT-INLINE
 // RUN: %clangxx -### -fsycl-device-only -O1 %s 2>&1 | FileCheck %s --check-prefix=CHECK-INLINE
 // RUN: %clang_cl -### -fsycl-device-only -O2 %s 2>&1 | FileCheck %s --check-prefix=CHECK-INLINE
+// RUN: %clangxx -### -fsycl-device-only -fintelfpga %s 2>&1 | FileCheck %s --check-prefix=CHECK-NOT-INLINE
 // CHECK-NOT-INLINE: "-fno-sycl-force-inline-kernel-lambda"
 // CHECK-INLINE-NOT: "-fno-sycl-force-inline-kernel-lambda"