diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 5dc209e847f5e..dbd49ac454d6c 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1878,16 +1878,18 @@ def SYCLIntelFPGAIVDep : StmtAttr {
   let Documentation = [SYCLIntelFPGAIVDepAttrDocs];
 }
 
-def SYCLIntelFPGAInitiationInterval : StmtAttr {
+def SYCLIntelFPGAInitiationInterval : DeclOrStmtAttr {
   let Spellings = [CXX11<"intelfpga","ii">,
                    CXX11<"intel","ii">,
                    CXX11<"intel", "initiation_interval">];
-  let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
-                             ErrorDiag, "'for', 'while', and 'do' statements">;
+  let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt, Function],
+                              ErrorDiag,
+                              "'for', 'while', 'do' statements, and functions">;
   let Args = [ExprArgument<"IntervalExpr", /*opt*/1>];
   let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let HasCustomTypeTransform = 1;
   let Documentation = [SYCLIntelFPGAInitiationIntervalAttrDocs];
+  let SupportsNonconformingLambdaSyntax = 1;
 }
 
 def SYCLIntelFPGAMaxConcurrency : StmtAttr {
@@ -1912,14 +1914,16 @@ def SYCLIntelFPGALoopCoalesce : StmtAttr {
   let Documentation = [SYCLIntelFPGALoopCoalesceAttrDocs];
 }
 
-def SYCLIntelFPGADisableLoopPipelining : StmtAttr {
+def SYCLIntelFPGADisableLoopPipelining : DeclOrStmtAttr {
   let Spellings = [CXX11<"intelfpga","disable_loop_pipelining">,
                    CXX11<"intel","disable_loop_pipelining">];
-  let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt],
-                             ErrorDiag, "'for', 'while', and 'do' statements">;
+  let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt, Function],
+                              ErrorDiag,
+                              "'for', 'while', 'do' statements, and functions">;
   let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let HasCustomTypeTransform = 1;
   let Documentation = [SYCLIntelFPGADisableLoopPipeliningAttrDocs];
+  let SupportsNonconformingLambdaSyntax = 1;
 }
 
 def SYCLIntelFPGAMaxInterleaving : StmtAttr {
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index bf90470ef4de7..9177e19c54945 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -2832,9 +2832,11 @@ def SYCLIntelFPGAInitiationIntervalAttrDocs : Documentation {
   let Category = DocCatVariable;
   let Heading = "intel::initiation_interval";
   let Content = [{
-This attribute applies to a loop. Indicates that the loop should be pipelined
-with an initiation interval of N. N must be a positive integer. Cannot be
-applied multiple times to the same loop.
+This attribute applies to a loop or a function. Indicates that the loop or
+function should be pipelined with an initiation interval of N. N must be a
+positive integer. Cannot be applied multiple times to the same loop or function.
+Cannot be used on the same loop or function in conjunction with
+disable_loop_pipelining.
 
 The ``[[intel::ii]]`` attribute spelling is a deprecated synonym for
 ``[[intel::initiation_interval]]`` and will be removed in the future.
@@ -2846,11 +2848,16 @@ The ``[[intel::ii]]`` attribute spelling is a deprecated synonym for
     [[intel::initiation_interval(4)]] for (int i = 0; i < 10; ++i) var++;
   }
 
+  [[intel::initiation_interval(4)]] void foo1 { }
+
   template<int N>
   void bar() {
     [[intel::initiation_interval(N)]] for(;;) { }
   }
 
+  template<int N>
+  [[intel::initiation_interval(N)]] void bar1 { }
+
   }];
 }
 
@@ -2921,10 +2928,11 @@ def SYCLIntelFPGADisableLoopPipeliningAttrDocs : Documentation {
   let Category = DocCatVariable;
   let Heading = "intel::disable_loop_pipelining";
   let Content = [{
-This attribute applies to a loop. Disables pipelining of the loop data path,
-causing the loop to be executed serially. Cannot be used on the same loop in
-conjunction with max_interleaving, speculated_iterations, max_concurrency, ii
-or ivdep.
+This attribute applies to a loop or a function. Takes no arguments and
+disables pipelining of the loop or function data path, causing the loop
+or function to be executed serially. Cannot be used on the same loop or
+function in conjunction with max_interleaving, speculated_iterations,
+max_concurrency, initiation_interval, or ivdep.
 
 .. code-block:: c++
 
@@ -2933,6 +2941,8 @@ or ivdep.
     [[intel::disable_loop_pipelining] for (int i = 0; i < 10; ++i) var++;
   }
 
+  [[intel::disable_loop_pipelining] void foo1() { }
+
   }];
 }
 
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index f53fd9f5621b5..856121e0f6b48 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10263,6 +10263,11 @@ class Sema final {
   IntelFPGAForcePow2DepthAttr *
   MergeIntelFPGAForcePow2DepthAttr(Decl *D,
                                    const IntelFPGAForcePow2DepthAttr &A);
+  void AddSYCLIntelFPGAInitiationIntervalAttr(Decl *D,
+                                              const AttributeCommonInfo &CI,
+                                              Expr *E);
+  SYCLIntelFPGAInitiationIntervalAttr *MergeSYCLIntelFPGAInitiationIntervalAttr(
+      Decl *D, const SYCLIntelFPGAInitiationIntervalAttr &A);
 
   /// AddAlignedAttr - Adds an aligned attribute to a particular declaration.
   void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E,
diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index 1bc87b53c94f4..839fb556dcdba 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -738,6 +738,22 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
         llvm::ConstantAsMetadata::get(Builder.getInt32(1))};
     Fn->setMetadata("stall_enable", llvm::MDNode::get(Context, AttrMDArgs));
   }
+
+  if (FD->hasAttr<SYCLIntelFPGADisableLoopPipeliningAttr>()) {
+    llvm::Metadata *AttrMDArgs[] = {
+        llvm::ConstantAsMetadata::get(Builder.getInt32(1))};
+    Fn->setMetadata("disable_loop_pipelining",
+                    llvm::MDNode::get(Context, AttrMDArgs));
+  }
+
+  if (const auto *A = FD->getAttr<SYCLIntelFPGAInitiationIntervalAttr>()) {
+    const auto *CE = cast<ConstantExpr>(A->getIntervalExpr());
+    llvm::APSInt ArgVal = CE->getResultAsAPSInt();
+    llvm::Metadata *AttrMDArgs[] = {
+        llvm::ConstantAsMetadata::get(Builder.getInt32(ArgVal.getSExtValue()))};
+    Fn->setMetadata("initiation_interval",
+                    llvm::MDNode::get(Context, AttrMDArgs));
+  }
 }
 
 /// Determine whether the function F ends with a return stmt.
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index d4d937c249b32..da65779bb63de 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -2630,6 +2630,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
     NewAttr = S.MergeIntelFPGAMaxReplicatesAttr(D, *A);
   else if (const auto *A = dyn_cast<IntelFPGAForcePow2DepthAttr>(Attr))
     NewAttr = S.MergeIntelFPGAForcePow2DepthAttr(D, *A);
+  else if (const auto *A = dyn_cast<SYCLIntelFPGAInitiationIntervalAttr>(Attr))
+    NewAttr = S.MergeSYCLIntelFPGAInitiationIntervalAttr(D, *A);
   else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr))
     NewAttr = cast<InheritableAttr>(Attr->clone(S.Context));
 
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 3d61bd65f4d72..f2185137e6d4b 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3330,6 +3330,105 @@ static void handleUseStallEnableClustersAttr(Sema &S, Decl *D,
   handleSimpleAttribute<SYCLIntelUseStallEnableClustersAttr>(S, D, Attr);
 }
 
+// Handles disable_loop_pipelining attribute.
+static void handleSYCLIntelFPGADisableLoopPipeliningAttr(Sema &S, Decl *D,
+                                                         const ParsedAttr &A) {
+  S.CheckDeprecatedSYCLAttributeSpelling(A);
+
+  // [[intel::disable_loop_pipelining] and [[intel::initiation_interval()]]
+  // attributes are incompatible.
+  if (checkAttrMutualExclusion<SYCLIntelFPGAInitiationIntervalAttr>(S, D, A))
+    return;
+
+  D->addAttr(::new (S.Context)
+                 SYCLIntelFPGADisableLoopPipeliningAttr(S.Context, A));
+}
+
+// Handles initiation_interval attribute.
+void Sema::AddSYCLIntelFPGAInitiationIntervalAttr(Decl *D,
+                                                  const AttributeCommonInfo &CI,
+                                                  Expr *E) {
+  if (!E->isValueDependent()) {
+    // Validate that we have an integer constant expression and then store the
+    // converted constant expression into the semantic attribute so that we
+    // don't have to evaluate it again later.
+    llvm::APSInt ArgVal;
+    ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal);
+    if (Res.isInvalid())
+      return;
+    E = Res.get();
+    // This attribute requires a strictly positive value.
+    if (ArgVal <= 0) {
+      Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
+          << CI << /*positive*/ 0;
+      return;
+    }
+    // Check to see if there's a duplicate attribute with different values
+    // already applied to the declaration.
+    if (const auto *DeclAttr =
+            D->getAttr<SYCLIntelFPGAInitiationIntervalAttr>()) {
+      // If the other attribute argument is instantiation dependent, we won't
+      // have converted it to a constant expression yet and thus we test
+      // whether this is a null pointer.
+      if (const auto *DeclExpr =
+              dyn_cast<ConstantExpr>(DeclAttr->getIntervalExpr())) {
+        if (ArgVal != DeclExpr->getResultAsAPSInt()) {
+          Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
+          Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
+        }
+        // Drop the duplicate attribute.
+        return;
+      }
+    }
+  }
+
+  // [[intel::disable_loop_pipelining] and [[intel::initiation_interval()]]
+  // attributes are incompatible.
+  if (checkAttrMutualExclusion<SYCLIntelFPGADisableLoopPipeliningAttr>(*this, D,
+                                                                       CI))
+    return;
+
+  D->addAttr(::new (Context)
+                 SYCLIntelFPGAInitiationIntervalAttr(Context, CI, E));
+}
+
+SYCLIntelFPGAInitiationIntervalAttr *
+Sema::MergeSYCLIntelFPGAInitiationIntervalAttr(
+    Decl *D, const SYCLIntelFPGAInitiationIntervalAttr &A) {
+  // Check to see if there's a duplicate attribute with different values
+  // already applied to the declaration.
+  if (const auto *DeclAttr =
+          D->getAttr<SYCLIntelFPGAInitiationIntervalAttr>()) {
+    if (const auto *DeclExpr =
+            dyn_cast<ConstantExpr>(DeclAttr->getIntervalExpr())) {
+      if (const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getIntervalExpr())) {
+        if (DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
+          Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
+          Diag(A.getLoc(), diag::note_previous_attribute);
+        }
+        // Do not add a duplicate attribute.
+        return nullptr;
+      }
+    }
+  }
+
+  // [[intel::initiation_interval()]] and [[intel::disable_loop_pipelining]
+  // attributes are incompatible.
+  if (checkAttrMutualExclusion<SYCLIntelFPGADisableLoopPipeliningAttr>(*this, D,
+                                                                       A))
+    return nullptr;
+
+  return ::new (Context)
+      SYCLIntelFPGAInitiationIntervalAttr(Context, A, A.getIntervalExpr());
+}
+
+static void handleSYCLIntelFPGAInitiationIntervalAttr(Sema &S, Decl *D,
+                                                      const ParsedAttr &A) {
+  S.CheckDeprecatedSYCLAttributeSpelling(A);
+
+  S.AddSYCLIntelFPGAInitiationIntervalAttr(D, A, A.getArgAsExpr(0));
+}
+
 // Handle scheduler_target_fmax_mhz
 void Sema::AddSYCLIntelSchedulerTargetFmaxMhzAttr(Decl *D,
                                                   const AttributeCommonInfo &CI,
@@ -9275,6 +9374,12 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_SYCLIntelLoopFuse:
     handleSYCLIntelLoopFuseAttr(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLIntelFPGADisableLoopPipelining:
+    handleSYCLIntelFPGADisableLoopPipeliningAttr(S, D, AL);
+    break;
+  case ParsedAttr::AT_SYCLIntelFPGAInitiationInterval:
+    handleSYCLIntelFPGAInitiationIntervalAttr(S, D, AL);
+    break;
   case ParsedAttr::AT_VecTypeHint:
     handleVecTypeHint(S, D, AL);
     break;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 9ffa0ce900698..4a89906acae1c 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -565,6 +565,24 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
         }
       }
 
+      // Attribute "disable_loop_pipelining" can be applied explicitly on
+      // kernel function. Attribute should not be propagated from device
+      // functions to kernel.
+      if (auto *A = FD->getAttr<SYCLIntelFPGADisableLoopPipeliningAttr>()) {
+        if (ParentFD == SYCLKernel) {
+          Attrs.push_back(A);
+        }
+      }
+
+      // Attribute "initiation_interval" can be applied explicitly on
+      // kernel function. Attribute should not be propagated from device
+      // functions to kernel.
+      if (auto *A = FD->getAttr<SYCLIntelFPGAInitiationIntervalAttr>()) {
+        if (ParentFD == SYCLKernel) {
+          Attrs.push_back(A);
+        }
+      }
+
       // TODO: vec_len_hint should be handled here
 
       CallGraphNode *N = SYCLCG.getNode(FD);
@@ -3517,6 +3535,8 @@ void Sema::MarkDevice(void) {
         case attr::Kind::SYCLIntelNoGlobalWorkOffset:
         case attr::Kind::SYCLIntelUseStallEnableClusters:
         case attr::Kind::SYCLIntelLoopFuse:
+        case attr::Kind::SYCLIntelFPGADisableLoopPipelining:
+        case attr::Kind::SYCLIntelFPGAInitiationInterval:
         case attr::Kind::SYCLSimd: {
           if ((A->getKind() == attr::Kind::SYCLSimd) && KernelBody &&
               !KernelBody->getAttr<SYCLSimdAttr>()) {
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index d69282ccd6fa6..0b7bf057a685f 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -713,6 +713,16 @@ static void instantiateIntelFPGAMaxReplicatesAttr(
     S.AddIntelFPGAMaxReplicatesAttr(New, *A, Result.getAs<Expr>());
 }
 
+static void instantiateSYCLIntelFPGAInitiationIntervalAttr(
+    Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+    const SYCLIntelFPGAInitiationIntervalAttr *A, Decl *New) {
+  EnterExpressionEvaluationContext Unevaluated(
+      S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+  ExprResult Result = S.SubstExpr(A->getIntervalExpr(), TemplateArgs);
+  if (!Result.isInvalid())
+    S.AddSYCLIntelFPGAInitiationIntervalAttr(New, *A, Result.getAs<Expr>());
+}
+
 /// Determine whether the attribute A might be relevent to the declaration D.
 /// If not, we can skip instantiating it. The attribute may or may not have
 /// been instantiated yet.
@@ -939,6 +949,12 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
           *this, TemplateArgs, SYCLIntelMaxWorkGroupSize, New);
       continue;
     }
+    if (const auto *SYCLIntelFPGAInitiationInterval =
+            dyn_cast<SYCLIntelFPGAInitiationIntervalAttr>(TmplAttr)) {
+      instantiateSYCLIntelFPGAInitiationIntervalAttr(
+          *this, TemplateArgs, SYCLIntelFPGAInitiationInterval, New);
+      continue;
+    }
     // Existing DLL attribute on the instantiation takes precedence.
     if (TmplAttr->getKind() == attr::DLLExport ||
         TmplAttr->getKind() == attr::DLLImport) {
diff --git a/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp b/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp
new file mode 100644
index 0000000000000..33f83f6b1961f
--- /dev/null
+++ b/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp
@@ -0,0 +1,36 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+queue q;
+
+class Foo {
+public:
+  [[intel::disable_loop_pipelining]] void operator()() const {}
+};
+
+[[intel::disable_loop_pipelining]] void foo() {}
+
+int main() {
+  q.submit([&](handler &h) {
+    // Test attribute is presented on function metadata.
+    Foo f;
+    h.single_task<class test_kernel1>(f);
+
+    // Test attribute is not propagated.
+    h.single_task<class test_kernel2>(
+        []() { foo(); });
+
+    // Test attribute is applied on lambda.
+    h.single_task<class test_kernel3>(
+        []() [[intel::disable_loop_pipelining]]{});
+  });
+  return 0;
+}
+
+// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 !kernel_arg_buffer_location ![[NUM4:[0-9]+]] !disable_loop_pipelining ![[NUM5:[0-9]+]]
+// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 !kernel_arg_buffer_location ![[NUM4]]
+// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 !kernel_arg_buffer_location ![[NUM4]] !disable_loop_pipelining ![[NUM5]]
+// CHECK: ![[NUM4]] = !{}
+// CHECK: ![[NUM5]] = !{i32 1}
diff --git a/clang/test/CodeGenSYCL/initiation_interval.cpp b/clang/test/CodeGenSYCL/initiation_interval.cpp
new file mode 100644
index 0000000000000..95ef9ce4cde50
--- /dev/null
+++ b/clang/test/CodeGenSYCL/initiation_interval.cpp
@@ -0,0 +1,49 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+queue q;
+
+class Foo {
+public:
+  [[intel::initiation_interval(1)]] void operator()() const {}
+};
+
+template <int SIZE>
+class Functor {
+public:
+  [[intel::initiation_interval(SIZE)]] void operator()() const {}
+};
+
+[[intel::initiation_interval(5)]] void foo() {}
+
+int main() {
+  q.submit([&](handler &h) {
+    // Test attribute argument size.
+    Foo boo;
+    h.single_task<class kernel_name1>(boo);
+
+    // Test attribute is applied on lambda.
+    h.single_task<class kernel_name2>(
+        []() [[intel::initiation_interval(42)]]{});
+
+    // Test template argument.
+    Functor<2> f;
+    h.single_task<class kernel_name3>(f);
+
+    // Test attribute is not propagated.
+    h.single_task<class kernel_name4>(
+        []() { foo(); });
+  });
+  return 0;
+}
+
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 !kernel_arg_buffer_location ![[NUM0:[0-9]+]] !initiation_interval ![[NUM1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM42:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM2:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 !kernel_arg_buffer_location ![[NUM0]]
+// CHECK: ![[NUM0]] = !{}
+// CHECK: ![[NUM1]] = !{i32 1}
+// CHECK: ![[NUM42]] = !{i32 42}
+// CHECK: ![[NUM2]] = !{i32 2}
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index c9c3caabcd492..ec4848a24277c 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -152,6 +152,8 @@
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDevice (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLIntelFPGADisableLoopPipelining (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLIntelFPGAInitiationInterval (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelKernelArgsRestrict (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelLoopFuse (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLIntelMaxGlobalWorkDim (SubjectMatchRule_function)
diff --git a/clang/test/SemaSYCL/disable_loop_pipelining.cpp b/clang/test/SemaSYCL/disable_loop_pipelining.cpp
new file mode 100644
index 0000000000000..c36f0f6d29cdb
--- /dev/null
+++ b/clang/test/SemaSYCL/disable_loop_pipelining.cpp
@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 %s -fsyntax-only -internal-isystem %S/Inputs -fsycl-is-device -sycl-std=2020  -DTRIGGER_ERROR -verify
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -ast-dump -sycl-std=2020 %s | FileCheck %s
+
+// Test that checks disable_loop_pipelining attribute support on Function.
+
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
+
+#ifdef TRIGGER_ERROR
+[[intel::disable_loop_pipelining(1)]] void bar1() {} // expected-error{{'disable_loop_pipelining' attribute takes no arguments}}
+[[intel::disable_loop_pipelining]] int N;            // expected-error{{'disable_loop_pipelining' attribute only applies to 'for', 'while', 'do' statements, and functions}}
+#endif
+
+struct FuncObj {
+  [[intel::disable_loop_pipelining]] void operator()() const {}
+};
+
+int main() {
+  deviceQueue.submit([&](sycl::handler &h) {
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1
+    // CHECK:       SYCLIntelFPGADisableLoopPipeliningAttr {{.*}}
+    h.single_task<class test_kernel1>(
+        FuncObj());
+
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2
+    // CHECK:       SYCLIntelFPGADisableLoopPipeliningAttr {{.*}}
+    h.single_task<class test_kernel2>(
+        []() [[intel::disable_loop_pipelining]]{});
+  });
+  return 0;
+}
diff --git a/clang/test/SemaSYCL/initiation_interval.cpp b/clang/test/SemaSYCL/initiation_interval.cpp
new file mode 100644
index 0000000000000..556e9671831ee
--- /dev/null
+++ b/clang/test/SemaSYCL/initiation_interval.cpp
@@ -0,0 +1,84 @@
+// RUN: %clang_cc1 -fsycl-is-device -verify %s
+
+// Test that checks disable_loop_pipelining attribute support on Function.
+
+// Tests for incorrect argument values for Intel FPGA initiation_interval function attribute.
+[[intel::initiation_interval(5)]] int a; // expected-error{{'initiation_interval' attribute only applies to 'for', 'while', 'do' statements, and functions}}
+
+[[intel::initiation_interval("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char [4]'}}
+
+[[intel::initiation_interval(-1)]] void func1() {} // expected-error{{'initiation_interval' attribute requires a positive integral compile time constant expression}}
+
+[[intel::initiation_interval(0, 1)]] void func2() {} // expected-error{{'initiation_interval' attribute takes no more than 1 argument}}
+
+// Tests for Intel FPGA initiation_interval function attribute duplication.
+// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
+[[intel::initiation_interval(2)]]
+[[intel::initiation_interval(2)]] void func3() {}
+
+// No diagnostic is emitted because the arguments match.
+[[intel::initiation_interval(4)]] void func4();
+[[intel::initiation_interval(4)]] void func4(); // OK
+
+// Diagnostic is emitted because the arguments mismatch.
+[[intel::initiation_interval(2)]]                  // expected-note {{previous attribute is here}}
+[[intel::initiation_interval(4)]] void func5() {}  // expected-warning {{attribute 'initiation_interval' is already applied with different arguments}}
+
+[[intel::initiation_interval(1)]] void func6(); // expected-note {{previous attribute is here}}
+[[intel::initiation_interval(3)]] void func6(); // expected-warning {{attribute 'initiation_interval' is already applied with different arguments}}
+
+// Tests for Intel FPGA loop fusion function attributes compatibility
+// expected-error@+2 {{'initiation_interval' and 'disable_loop_pipelining' attributes are not compatible}}
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::disable_loop_pipelining]] [[intel::initiation_interval(2)]] void func7();
+
+// expected-error@+2 {{'disable_loop_pipelining' and 'initiation_interval' attributes are not compatible}}
+// expected-note@+1 {{conflicting attribute is here}}
+[[intel::initiation_interval(4)]] [[intel::disable_loop_pipelining]] void func8();
+
+// expected-error@+2 {{'initiation_interval' and 'disable_loop_pipelining' attributes are not compatible}}
+// expected-note@+2 {{conflicting attribute is here}}
+[[intel::initiation_interval(4)]] void func9();
+[[intel::disable_loop_pipelining]] void func9();
+
+// Tests that check template parameter support for Intel FPGA initiation_interval function attributes
+template <int N>
+[[intel::initiation_interval(N)]] void func10(); // expected-error 2{{'initiation_interval' attribute requires a positive integral compile time constant expression}}
+
+template <int size>
+[[intel::initiation_interval(10)]] void func11();     // expected-note {{previous attribute is here}}
+template <int size>
+[[intel::initiation_interval(size)]] void func11() {} // expected-warning {{attribute 'initiation_interval' is already applied with different arguments}}
+
+void checkTemplates() {
+  func10<4>();  // OK
+  func10<-1>(); // expected-note {{in instantiation of function template specialization 'func10<-1>' requested here}}
+  func10<0>();  // expected-note {{in instantiation of function template specialization 'func10<0>' requested here}}
+  func11<20>(); // expected-note {{in instantiation of function template specialization 'func11<20>' requested here}}
+}
+
+// Test that checks expression is not a constant expression.
+// expected-note@+1{{declared here}}
+int baz();
+// expected-error@+2{{expression is not an integral constant expression}}
+// expected-note@+1{{non-constexpr function 'baz' cannot be used in a constant expression}}
+[[intel::initiation_interval(baz() + 1)]] void func12();
+
+// Test that checks expression is a constant expression.
+constexpr int bar() { return 0; }
+[[intel::initiation_interval(bar() + 2)]] void func13(); // OK
+
+// Test that checks wrong function template instantiation and ensures that the type
+// is checked properly when instantiating from the template definition.
+template <typename Ty>
+// expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
+// expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
+[[intel::initiation_interval(Ty{})]] void func14() {}
+
+struct S {};
+void test() {
+  //expected-note@+1{{in instantiation of function template specialization 'func14<S>' requested here}}
+  func14<S>();
+  //expected-note@+1{{in instantiation of function template specialization 'func14<float>' requested here}}
+  func14<float>();
+}
diff --git a/clang/test/SemaSYCL/initiation_interval_ast.cpp b/clang/test/SemaSYCL/initiation_interval_ast.cpp
new file mode 100644
index 0000000000000..777230e224572
--- /dev/null
+++ b/clang/test/SemaSYCL/initiation_interval_ast.cpp
@@ -0,0 +1,97 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s
+
+// Tests for AST of Intel FPGA initiation_interval function attributes.
+#include "sycl.hpp"
+
+sycl::queue deviceQueue;
+
+// CHECK: FunctionDecl {{.*}} func1 'void ()'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: SYCLIntelFPGAInitiationIntervalAttr {{.*}} initiation_interval
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 4
+// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4
+[[intel::initiation_interval(4)]] void func1() {}
+
+// Test that checks template parameter support on function.
+// CHECK: FunctionTemplateDecl {{.*}} func2
+// CHECK: FunctionDecl {{.*}} func2 'void ()'
+// CHECK-NEXT: CompoundStmt
+// CHECK_NEXT: SYCLIntelFPGAInitiationIntervalAttr {{.*}} initiation_interval
+// CHECK_NEXT: DeclRefExpr {{.*}} 'int' NonTypeTemplateParm {{.*}} 'N' 'int'
+// CHECK: FunctionDecl {{.*}} func2 'void ()'
+// CHECK-NEXT: TemplateArgument integral 6
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: SYCLIntelFPGAInitiationIntervalAttr {{.*}} initiation_interval
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 6
+// CHECK-NEXT: SubstNonTypeTemplateParmExpr
+// CHECK-NEXT: NonTypeTemplateParmDecl
+// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 6
+template <int N>
+[[intel::initiation_interval(N)]] void func2() {}
+
+// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
+// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()'
+// CHECK: SYCLIntelFPGAInitiationIntervalAttr {{.*}} initiation_interval
+// CHECK-NEXT: ConstantExpr {{.*}} 'int'
+// CHECK-NEXT: value: Int 10
+// CHECK-NEXT: IntegerLiteral{{.*}}10{{$}}
+[[intel::initiation_interval(10)]]
+[[intel::initiation_interval(10)]] void func3() {}
+
+class KernelFunctor {
+public:
+  void operator()() const {
+    func1();
+  }
+};
+
+// Test that checks template parameter support on class member function.
+template <int N>
+class KernelFunctor2 {
+public:
+  [[intel::initiation_interval(N)]] void operator()() const {
+  }
+};
+
+int main() {
+  deviceQueue.submit([&](sycl::handler &h) {
+    // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1
+    // CHECK-NOT: SYCLIntelFPGAInitiationIntervalAttr
+    KernelFunctor f1;
+    h.single_task<class kernel_name_1>(f1);
+
+    // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_2
+    // CHECK: SYCLIntelFPGAInitiationIntervalAttr {{.*}} initiation_interval
+    // CHECK-NEXT: ConstantExpr{{.*}}'int'
+    // CHECK-NEXT: value: Int 3
+    // CHECK-NEXT: SubstNonTypeTemplateParmExpr
+    // CHECK-NEXT: NonTypeTemplateParmDecl
+    // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 3
+    KernelFunctor2<3> f2;
+    h.single_task<class kernel_name_2>(f2);
+    
+    // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_3
+    // CHECK: SYCLIntelFPGAInitiationIntervalAttr {{.*}} initiation_interval
+    // CHECK-NEXT: ConstantExpr{{.*}}'int'
+    // CHECK-NEXT: value: Int 4
+    // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4
+    h.single_task<class kernel_name_3>(
+        []() [[intel::initiation_interval(4)]]{});
+
+    // Ignore duplicate attribute.
+    h.single_task<class kernel_name_4>(
+    // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_4
+    // CHECK: SYCLIntelFPGAInitiationIntervalAttr {{.*}} initiation_interval
+    // CHECK-NEXT: ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT: value: Int 6
+    // CHECK-NEXT: IntegerLiteral{{.*}}6{{$}}
+        []() [[intel::initiation_interval(6),
+               intel::initiation_interval(6)]]{});
+  });
+
+  func2<6>();
+
+  return 0;
+}
diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp
index 77a8ac100a7f7..274f47fb6aac7 100644
--- a/clang/test/SemaSYCL/intel-fpga-loops.cpp
+++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp
@@ -8,11 +8,11 @@ sycl::queue deviceQueue;
 void foo() {
   // expected-error@+1 {{'ivdep' attribute cannot be applied to a declaration}}
   [[intel::ivdep]] int a[10];
-  // expected-error@+1 {{'initiation_interval' attribute cannot be applied to a declaration}}
+  // expected-error@+1 {{'initiation_interval' attribute only applies to 'for', 'while', 'do' statements, and functions}}
   [[intel::initiation_interval(2)]] int c[10];
   // expected-error@+1 {{'max_concurrency' attribute cannot be applied to a declaration}}
   [[intel::max_concurrency(2)]] int d[10];
-  // expected-error@+1 {{'disable_loop_pipelining' attribute cannot be applied to a declaration}}
+  // expected-error@+1 {{'disable_loop_pipelining' attribute only applies to 'for', 'while', 'do' statements, and functions}}
   [[intel::disable_loop_pipelining]] int g[10];
   // expected-error@+1 {{'loop_coalesce' attribute cannot be applied to a declaration}}
   [[intel::loop_coalesce(2)]] int h[10];