Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9784,6 +9784,11 @@
 def warn_loop_ctrl_binds_to_inner : Warning<
   "'%0' is bound to current loop, GCC binds it to the enclosing loop">,
   InGroup<GccCompat>;
+def err_omp_bind_required_on_loop : Error<
+  "expected 'bind' clause for loop construct without an enclosing OpenMP "
+  "construct">;
+def err_omp_loop_reduction_clause : Error<
+  "reduction clause not handled with '#pragma omp loop bind(teams)'">;
 def warn_break_binds_to_switch : Warning<
   "'break' is bound to loop, GCC binds it to switch">,
   InGroup<GccCompat>;
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -635,6 +635,11 @@
     const SharingMapTy *Top = getTopOfStackOrNull();
     return Top ? Top->Directive : OMPD_unknown;
   }
+  void setCurrentDirective(OpenMPDirectiveKind newDK) {
+    SharingMapTy *Top = (SharingMapTy *)getTopOfStackOrNull();
+    assert(Top != NULL);
+    Top->Directive = newDK;
+  }
   /// Returns directive kind at specified level.
   OpenMPDirectiveKind getDirective(unsigned Level) const {
     assert(!isStackEmpty() && "No directive at specified level.");
@@ -6120,10 +6125,75 @@
                             isOpenMPTargetDataManagementDirective(Kind)))
     Diag(StartLoc, diag::warn_hip_omp_target_directives);
 
+  llvm::SmallVector<OMPClause *, 8> ClausesWithoutBind;
+  bool UseClausesWithoutBind = false;
+
+  // Restricting to "#pragma omp loop bind"
+  if (Kind == OMPD_loop) {
+    if (BindKind == OMPC_BIND_unknown) {
+      // Setting the enclosing teams or parallel construct for the loop
+      // directive without bind clause.
+      BindKind = OMPC_BIND_thread; // Default bind(thread) if binding is unknown
+
+      const OpenMPDirectiveKind ParentDirective =
+          DSAStack->getParentDirective();
+      if (ParentDirective == OMPD_unknown) {
+        Diag(DSAStack->getDefaultDSALocation(),
+             diag::err_omp_bind_required_on_loop);
+      } else if (ParentDirective == OMPD_parallel ||
+                 ParentDirective == OMPD_target_parallel) {
+        BindKind = OMPC_BIND_parallel;
+      } else if (ParentDirective == OMPD_teams ||
+                 ParentDirective == OMPD_target_teams) {
+        BindKind = OMPC_BIND_teams;
+      }
+    } else {
+      // bind clause is present, so we should set flag indicating to only
+      // use the clauses that aren't the bind clause for the new directive that
+      // loop is lowered to.
+      UseClausesWithoutBind = true;
+    }
+
+    for (OMPClause *C : Clauses) {
+      // Spec restriction : bind(teams) and reduction not permitted.
+      if (BindKind == OMPC_BIND_teams &&
+          C->getClauseKind() == llvm::omp::Clause::OMPC_reduction)
+        Diag(DSAStack->getDefaultDSALocation(),
+             diag::err_omp_loop_reduction_clause);
+
+      // A new Vector ClausesWithoutBind, which does not contain the bind
+      // clause, for passing to new directive.
+      if (C->getClauseKind() != llvm::omp::Clause::OMPC_bind)
+        ClausesWithoutBind.push_back(C);
+    }
+
+    switch (BindKind) {
+    case OMPC_BIND_parallel:
+      Kind = OMPD_for;
+      DSAStack->setCurrentDirective(OMPD_for);
+      break;
+    case OMPC_BIND_teams:
+      Kind = OMPD_distribute;
+      DSAStack->setCurrentDirective(OMPD_distribute);
+      break;
+    case OMPC_BIND_thread:
+      Kind = OMPD_simd;
+      DSAStack->setCurrentDirective(OMPD_simd);
+      break;
+    case OMPC_BIND_unknown:
+      break;
+    }
+  }
+
   llvm::SmallVector<OMPClause *, 8> ClausesWithImplicit;
   VarsWithInheritedDSAType VarsWithInheritedDSA;
   bool ErrorFound = false;
-  ClausesWithImplicit.append(Clauses.begin(), Clauses.end());
+  if (UseClausesWithoutBind) {
+    ClausesWithImplicit.append(ClausesWithoutBind.begin(),
+                               ClausesWithoutBind.end());
+  } else {
+    ClausesWithImplicit.append(Clauses.begin(), Clauses.end());
+  }
   if (AStmt && !CurContext->isDependentContext() && Kind != OMPD_atomic &&
       Kind != OMPD_critical && Kind != OMPD_section && Kind != OMPD_master &&
       Kind != OMPD_masked && !isOpenMPLoopTransformationDirective(Kind)) {
Index: clang/test/OpenMP/loop_bind_codegen.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/loop_bind_codegen.cpp
@@ -0,0 +1,151 @@
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+
+
+/*
+#include <stdio.h>
+#include <assert.h>
+#include <pthread.h>
+#include <omp.h>
+*/
+
+#define NNN 50
+int aaa[NNN];
+
+void parallel_loop() {
+  #pragma omp parallel
+  {
+    #pragma omp loop bind(parallel)
+    for (int j = 0 ; j < NNN ; j++) {
+      aaa[j] = j*NNN;
+    }
+  }
+}
+
+void parallel_loop_orphan() {
+  #pragma omp loop bind(parallel)
+  for (int j = 0 ; j < NNN ; j++) {
+     aaa[j] = j*NNN;
+  }
+}
+
+
+void teams_loop() {
+  #pragma omp teams
+  {
+     #pragma omp loop bind(teams)
+     for (int j = 0 ; j < NNN ; j++) {
+       aaa[j] = j*NNN;
+     }
+   }
+}
+
+void thread_loop() {
+  #pragma omp parallel
+  {
+     #pragma omp loop bind(thread)
+     for (int j = 0 ; j < NNN ; j++) {
+       aaa[j] = j*NNN;
+     }
+   }
+}
+
+void thread_loop_orphan() {
+  #pragma omp loop bind(thread)
+  for (int j = 0 ; j < NNN ; j++) {
+    aaa[j] = j*NNN;
+  }
+}
+
+int main() {
+  parallel_loop();
+  parallel_loop_orphan();
+  teams_loop();
+  thread_loop();
+  thread_loop_orphan();
+
+  return 0;
+}
+#endif
+// CHECK-LABEL: define {{[^@]+}}@_Z13parallel_loopv
+// CHECK-NEXT: entry:
+// CHECK-NEXT:     call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @3, i32 0, ptr @.omp_outlined.)
+// CHECK-LABEL: define {{[^@]+}}@.omp_outlined
+// CHECK-NEXT: entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[J:%.*]] = alloca i32, align 4
+// CHECK:       call void @__kmpc_for_static_init_4(ptr @1, i32 %1, i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
+// CHECK-LABEL: cond.true:
+// CHECK-NEXT:    br label [[COND_END:%.*]]
+// CHECK-LABEL: cond.false:
+// CHECK-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
+// CHECK-NEXT:    br label [[COND_END]]
+// CHECK:       omp.inner.for.cond:
+// CHECK-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
+// CHECK:       omp.inner.for.body:
+// CHECK:       omp.loop.exit:
+// CHECK-NEXT:    call void @__kmpc_for_static_fini(ptr @1, i32 %1)
+// CHECK-NEXT:    call void @__kmpc_barrier(ptr @2, i32 %1)
+// CHECK-NEXT:    ret void
+//
+// CHECK-LABEL: define {{[^@]+}}@_Z20parallel_loop_orphanv
+// CHECK-NEXT: entry:
+// CHECK-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK:       call void @__kmpc_for_static_init_4(ptr @1, i32 %0, i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
+// CHECK-LABEL: cond.true:
+// CHECK:       omp.inner.for.cond:
+// CHECK:       omp.loop.exit:
+//
+// CHECK-LABEL: define {{[^@]+}}@_Z10teams_loopv
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:     call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @3, i32 0, ptr @.omp_outlined..1)
+// CHECK-NEXT:     ret void
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_outlined
+// CHECK-NEXT: entry:
+// CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[J:%.*]] = alloca i32, align 4
+// CHECK:       call void @__kmpc_for_static_init_4(ptr @4, i32 %1, i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
+// CHECK-LABEL: cond.true:
+// CHECK-NEXT:    br label [[COND_END:%.*]]
+// CHECK-LABEL: cond.false:
+// CHECK-NEXT:    [[TMP11:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4
+// CHECK-NEXT:    br label [[COND_END]]
+// CHECK:       omp.inner.for.cond:
+// CHECK-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
+// CHECK:       omp.inner.for.body:
+// CHECK:       omp.loop.exit:
+// CHECK-NEXT:    call void @__kmpc_for_static_fini(ptr @4, i32 %1)
+// CHECK-NEXT:    ret void
+//
+// CHECK-LABEL: define {{[^@]+}}@_Z11thread_loopv
+// CHECK-NEXT: entry:
+// CHECK-NEXT:     call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @3, i32 0, ptr @.omp_outlined..2)
+// CHECK-LABEL: define {{[^@]+}}@.omp_outlined
+// CHECK:       omp.inner.for.cond:
+//
+// CHECK-LABEL: define {{[^@]+}}@_Z18thread_loop_orphanv
+// CHECK-NEXT: entry:
+// CHECK-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// CHECK-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// CHECK:       omp.inner.for.cond:
+// CHECK-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4
+//
+// CHECK-LABEL: @main{{.*}}
Index: clang/test/OpenMP/loop_bind_messages.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/loop_bind_messages.cpp
@@ -0,0 +1,76 @@
+#ifndef HEADER
+#define HEADER
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=50 -verify %s
+
+#define NNN 50
+int aaa[NNN];
+
+void parallel_loop() {
+  #pragma omp parallel
+  {
+     #pragma omp loop
+     for (int j = 0 ; j < NNN ; j++) {
+       aaa[j] = j*NNN;
+     }
+   }
+}
+
+void teams_loop() {
+  int var1, var2;
+
+  #pragma omp teams
+  {
+     #pragma omp loop bind(teams)
+     for (int j = 0 ; j < NNN ; j++) {
+       aaa[j] = j*NNN;
+     }
+
+     #pragma omp loop bind(teams) collapse(2) private(var1)
+     for (int i = 0 ; i < 3 ; i++) {
+       for (int j = 0 ; j < NNN ; j++) {
+         var1 += aaa[j];
+       }
+     }
+   }
+}
+
+void orphan_loop_with_bind() {
+  #pragma omp loop bind(parallel) 
+  for (int j = 0 ; j < NNN ; j++) {
+    aaa[j] = j*NNN;
+  }
+}
+
+void orphan_loop_no_bind() {
+  #pragma omp loop  // expected-error{{expected 'bind' clause for loop construct without an enclosing OpenMP construct}}
+  for (int j = 0 ; j < NNN ; j++) {
+    aaa[j] = j*NNN;
+  }
+}
+
+void teams_loop_reduction() {
+  int total = 0;
+
+  #pragma omp teams
+  {
+     #pragma omp loop bind(teams)
+     for (int j = 0 ; j < NNN ; j++) {
+       aaa[j] = j*NNN;
+     }
+
+     #pragma omp loop bind(teams) reduction(+:total) // expected-error{{reduction clause not handled with '#pragma omp loop bind(teams)'}}
+     for (int j = 0 ; j < NNN ; j++) {
+       total+=aaa[j];
+     }
+   }
+}
+
+int main(int argc, char *argv[]) {
+  parallel_loop();
+  teams_loop();
+  orphan_loop_with_bind();
+  orphan_loop_no_bind();
+  teams_loop_reduction();
+}
+
+#endif