[PATCH] D99622: [OpenMP51] Accept `primary` as proc bind affinity policy in Clang

Chi Chun Chen via Phabricator via cfe-commits Wed, 31 Mar 2021 10:36:55 -0700

cchen updated this revision to Diff 334482.
cchen added a comment.

Add some codegen tests



Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D99622/new/

https://reviews.llvm.org/D99622

Files:
  clang/lib/Sema/SemaOpenMP.cpp
  clang/test/OpenMP/distribute_parallel_for_proc_bind_codegen.cpp
  clang/test/OpenMP/nvptx_target_parallel_proc_bind_codegen.cpp
  clang/test/OpenMP/parallel_ast_print.cpp
  clang/test/OpenMP/parallel_master_codegen.cpp
  clang/test/OpenMP/parallel_proc_bind_codegen.cpp
  clang/test/OpenMP/parallel_proc_bind_messages.cpp
  clang/test/OpenMP/target_teams_distribute_parallel_for_proc_bind_codegen.cpp
  
clang/test/OpenMP/target_teams_distribute_parallel_for_simd_proc_bind_codegen.cpp
  clang/test/OpenMP/teams_distribute_parallel_for_proc_bind_codegen.cpp
  clang/test/OpenMP/teams_distribute_parallel_for_simd_proc_bind_codegen.cpp
  llvm/include/llvm/Frontend/OpenMP/OMP.td
  llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

Index: llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -967,6 +967,7 @@
 __OMP_PROC_BIND_KIND(master, 2)
 __OMP_PROC_BIND_KIND(close, 3)
 __OMP_PROC_BIND_KIND(spread, 4)
+__OMP_PROC_BIND_KIND(primary, 5)
 __OMP_PROC_BIND_KIND(default, 6)
 __OMP_PROC_BIND_KIND(unknown, 7)
 
Index: llvm/include/llvm/Frontend/OpenMP/OMP.td
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMP.td
+++ llvm/include/llvm/Frontend/OpenMP/OMP.td
@@ -103,13 +103,15 @@
 def OMP_PROC_BIND_master : ClauseVal<"master",2,1> {}
 def OMP_PROC_BIND_close : ClauseVal<"close",3,1> {}
 def OMP_PROC_BIND_spread : ClauseVal<"spread",4,1> {}
-def OMP_PROC_BIND_default : ClauseVal<"default",5,0> {}
-def OMP_PROC_BIND_unknown : ClauseVal<"unknown",6,0> { let isDefault = true; }
+def OMP_PROC_BIND_primary : ClauseVal<"primary",5,1> {}
+def OMP_PROC_BIND_default : ClauseVal<"default",6,0> {}
+def OMP_PROC_BIND_unknown : ClauseVal<"unknown",7,0> { let isDefault = true; }
 def OMPC_ProcBind : Clause<"proc_bind"> {
   let clangClass = "OMPProcBindClause";
   let flangClass = "OmpProcBindClause";
   let enumClauseValue = "ProcBindKind";
   let allowedClauseValues = [
+    OMP_PROC_BIND_primary,
     OMP_PROC_BIND_master,
     OMP_PROC_BIND_close,
     OMP_PROC_BIND_spread,
Index: clang/test/OpenMP/teams_distribute_parallel_for_simd_proc_bind_codegen.cpp
===================================================================
--- clang/test/OpenMP/teams_distribute_parallel_for_simd_proc_bind_codegen.cpp
+++ clang/test/OpenMP/teams_distribute_parallel_for_simd_proc_bind_codegen.cpp
@@ -1,22 +1,23 @@
 // add -fopenmp-targets
-
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
-
-// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK1
 typedef __INTPTR_TYPE__ intptr_t;
 
-// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
-// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
-// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+// CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
 
 void foo();
 
@@ -36,7 +37,7 @@
 }
 
 int main() {
-  // CHECK-LABEL: @main
+  // CK1-LABEL: @main
 #pragma omp target
 #pragma omp teams distribute parallel for simd proc_bind(spread)
   for(int i = 0; i < 1000; i++) {}
@@ -46,53 +47,123 @@
   return tmain<int>();
 }
 
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL1:@.+]]()
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL2:@.+]]()
-// CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
-// CHECK: ret i32 [[CALL_RET]]
-
-// CHECK: define{{.+}} void [[OFFL1]](
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: define{{.+}} [[OFFL2]]()
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: define{{.+}} [[TMAIN]]()
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL3:@.+]]()
-
-// CHECK: define{{.+}} [[OFFL3]]()
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: !{!"llvm.loop.vectorize.enable", i1 true}
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL1:@.+]]()
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL2:@.+]]()
+// CK1: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
+// CK1: ret i32 [[CALL_RET]]
+
+// CK1: define{{.+}} void [[OFFL1]](
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: define{{.+}} [[OFFL2]]()
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: define{{.+}} [[TMAIN]]()
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL3:@.+]]()
+
+// CK1: define{{.+}} [[OFFL3]]()
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: !{!"llvm.loop.vectorize.enable", i1 true}
+
+#endif
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK2
+typedef __INTPTR_TYPE__ intptr_t;
+
+// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+
+void foo();
 
+template <typename T>
+T tmain() {
+#pragma omp target
+#pragma omp teams distribute parallel for simd proc_bind(primary)
+  for(int i = 0; i < 1000; i++) {}
+  return T();
+}
+
+int main() {
+  // CK2-LABEL: @main
+#pragma omp target
+#pragma omp teams distribute parallel for simd proc_bind(primary)
+  for(int i = 0; i < 1000; i++) {}
+  return tmain<int>();
+}
+
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL1:@.+]]()
+// CK2: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
+// CK2: ret i32 [[CALL_RET]]
+
+// CK2: define{{.+}} void [[OFFL1]](
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+
+// CK2: define{{.+}} [[TMAIN]]()
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL3:@.+]]()
+
+// CK2: define{{.+}} [[OFFL2]]()
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+
+// CK2: !{!"llvm.loop.vectorize.enable", i1 true}
+#endif
 #endif
Index: clang/test/OpenMP/teams_distribute_parallel_for_proc_bind_codegen.cpp
===================================================================
--- clang/test/OpenMP/teams_distribute_parallel_for_proc_bind_codegen.cpp
+++ clang/test/OpenMP/teams_distribute_parallel_for_proc_bind_codegen.cpp
@@ -1,22 +1,23 @@
 // add -fopenmp-targets
-
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
-
-// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK1
 typedef __INTPTR_TYPE__ intptr_t;
 
-// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
-// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
-// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+// CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
 
 void foo();
 
@@ -36,7 +37,7 @@
 }
 
 int main() {
-  // CHECK-LABEL: @main
+  // CK1-LABEL: @main
 #pragma omp target
 #pragma omp teams distribute parallel for proc_bind(spread)
   for(int i = 0; i < 1000; i++) {}
@@ -46,50 +47,118 @@
   return tmain<int>();
 }
 
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL1:@.+]]()
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL2:@.+]]()
-// CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
-// CHECK: ret i32 [[CALL_RET]]
-
-// CHECK: define{{.+}} void [[OFFL1]](
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: define{{.+}} [[OFFL2]]()
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: define{{.+}} [[TMAIN]]()
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL3:@.+]]()
-
-// CHECK: define{{.+}} [[OFFL3]]()
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL1:@.+]]()
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL2:@.+]]()
+// CK1: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
+// CK1: ret i32 [[CALL_RET]]
+
+// CK1: define{{.+}} void [[OFFL1]](
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: define{{.+}} [[OFFL2]]()
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: define{{.+}} [[TMAIN]]()
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL3:@.+]]()
+
+// CK1: define{{.+}} [[OFFL3]]()
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+#endif
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK2
+typedef __INTPTR_TYPE__ intptr_t;
+
+// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+
+void foo();
+
+template <typename T>
+T tmain() {
+#pragma omp target
+#pragma omp teams distribute parallel for proc_bind(primary)
+  for(int i = 0; i < 1000; i++) {}
+  return T();
+}
+
+int main() {
+  // CK2-LABEL: @main
+#pragma omp target
+#pragma omp teams distribute parallel for proc_bind(primary)
+  for(int i = 0; i < 1000; i++) {}
+  return tmain<int>();
+}
+
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL1:@.+]]()
+// CK2: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
+// CK2: ret i32 [[CALL_RET]]
+
+// CK2: define{{.+}} void [[OFFL1]](
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+
+// CK2: define{{.+}} [[TMAIN]]()
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL3:@.+]]()
+
+// CK2: define{{.+}} [[OFFL2]]()
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+#endif
 #endif
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_proc_bind_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_proc_bind_codegen.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_proc_bind_codegen.cpp
@@ -1,23 +1,23 @@
 // add -fopenmp-targets
-
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
-
-// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
-// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
-
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
+// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK1
 typedef __INTPTR_TYPE__ intptr_t;
 
-// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
-// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
-// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+// CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
 
 void foo();
 
@@ -36,7 +36,7 @@
 }
 
 int main() {
-  // CHECK-LABEL: @main
+  // CK1-LABEL: @main
 #pragma omp target teams distribute parallel for simd proc_bind(spread)
   for(int i = 0; i < 1000; i++) {}
 #pragma omp target teams distribute parallel for simd proc_bind(close)
@@ -44,50 +44,116 @@
   return tmain<int>();
 }
 
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL1:@.+]]()
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL2:@.+]]()
-// CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
-// CHECK: ret i32 [[CALL_RET]]
-
-// CHECK: define{{.+}} void [[OFFL1]](
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: define{{.+}} [[OFFL2]]()
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: define{{.+}} [[TMAIN]]()
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL3:@.+]]()
-
-// CHECK: define{{.+}} [[OFFL3]]()
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL1:@.+]]()
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL2:@.+]]()
+// CK1: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
+// CK1: ret i32 [[CALL_RET]]
+
+// CK1: define{{.+}} void [[OFFL1]](
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: define{{.+}} [[OFFL2]]()
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: define{{.+}} [[TMAIN]]()
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL3:@.+]]()
+
+// CK1: define{{.+}} [[OFFL3]]()
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+#endif
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
+// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK2
+typedef __INTPTR_TYPE__ intptr_t;
+
+// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+
+void foo();
+
+template <typename T>
+T tmain() {
+#pragma omp target teams distribute parallel for simd proc_bind(primary)
+  for(int i = 0; i < 1000; i++) {}
+  return T();
+}
+
+int main() {
+  // CK2-LABEL: @main
+#pragma omp target teams distribute parallel for simd proc_bind(primary)
+  for(int i = 0; i < 1000; i++) {}
+  return tmain<int>();
+}
+
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL1:@.+]]()
+// CK2: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
+// CK2: ret i32 [[CALL_RET]]
+
+// CK2: define{{.+}} void [[OFFL1]](
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+
+// CK2: define{{.+}} [[TMAIN]]()
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL2:@.+]]()
+
+// CK2: define{{.+}} [[OFFL2]]()
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+#endif
 #endif
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_proc_bind_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_proc_bind_codegen.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_proc_bind_codegen.cpp
@@ -1,23 +1,23 @@
 // add -fopenmp-targets
-
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
-
-// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
-// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
-
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
+// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK1
 typedef __INTPTR_TYPE__ intptr_t;
 
-// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
-// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
-// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+// CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
 
 void foo();
 
@@ -36,7 +36,7 @@
 }
 
 int main() {
-  // CHECK-LABEL: @main
+  // CK1-LABEL: @main
 #pragma omp target teams distribute parallel for proc_bind(spread)
   for(int i = 0; i < 1000; i++) {}
 #pragma omp target teams distribute parallel for proc_bind(close)
@@ -44,50 +44,117 @@
   return tmain<int>();
 }
 
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL1:@.+]]()
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL2:@.+]]()
-// CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
-// CHECK: ret i32 [[CALL_RET]]
-
-// CHECK: define{{.+}} void [[OFFL1]](
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: define{{.+}} [[OFFL2]]()
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: define{{.+}} [[TMAIN]]()
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL3:@.+]]()
-
-// CHECK: define{{.+}} [[OFFL3]]()
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL1:@.+]]()
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL2:@.+]]()
+// CK1: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
+// CK1: ret i32 [[CALL_RET]]
+
+// CK1: define{{.+}} void [[OFFL1]](
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: define{{.+}} [[OFFL2]]()
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: define{{.+}} [[TMAIN]]()
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL3:@.+]]()
+
+// CK1: define{{.+}} [[OFFL3]]()
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+#endif
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
+// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK2
+typedef __INTPTR_TYPE__ intptr_t;
+
+// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+
+void foo();
+
+template <typename T>
+T tmain() {
+#pragma omp target teams distribute parallel for proc_bind(primary)
+  for(int i = 0; i < 1000; i++) {}
+  return T();
+}
+
+int main() {
+  // CK2-LABEL: @main
+#pragma omp target teams distribute parallel for proc_bind(primary)
+  for(int i = 0; i < 1000; i++) {}
+  return tmain<int>();
+}
+
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL1:@.+]]()
+// CK2: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
+// CK2: ret i32 [[CALL_RET]]
+
+// CK2: define{{.+}} void [[OFFL1]](
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+
+// CK2: define{{.+}} [[TMAIN]]()
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL2:@.+]]()
+
+// CK2: define{{.+}} [[OFFL2]]()
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+#endif
 #endif
Index: clang/test/OpenMP/parallel_proc_bind_messages.cpp
===================================================================
--- clang/test/OpenMP/parallel_proc_bind_messages.cpp
+++ clang/test/OpenMP/parallel_proc_bind_messages.cpp
@@ -1,16 +1,20 @@
-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
 
 void foo();
 
 int main(int argc, char **argv) {
   #pragma omp parallel proc_bind // expected-error {{expected '(' after 'proc_bind'}}
-  #pragma omp parallel proc_bind ( // expected-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
-  #pragma omp parallel proc_bind () // expected-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}}
+  #pragma omp parallel proc_bind ( // ge51-error {{expected 'master', 'close', 'spread' or 'primary' in OpenMP clause 'proc_bind'}} lt51-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp parallel proc_bind () // ge51-error {{expected 'master', 'close', 'spread' or 'primary' in OpenMP clause 'proc_bind'}} lt51-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}}
   #pragma omp parallel proc_bind (master // expected-error {{expected ')'}} expected-note {{to match this '('}}
   #pragma omp parallel proc_bind (close), proc_bind(spread) // expected-error {{directive '#pragma omp parallel' cannot contain more than one 'proc_bind' clause}}
-  #pragma omp parallel proc_bind (x) // expected-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}}
+  #pragma omp parallel proc_bind (x) // ge51-error {{expected 'master', 'close', 'spread' or 'primary' in OpenMP clause 'proc_bind'}} lt51-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}} 
   foo();
 
   #pragma omp parallel proc_bind(master)
@@ -19,5 +23,8 @@
   #pragma omp parallel proc_bind(close)
   #pragma omp parallel proc_bind(spread)
   ++argc;
+
+  #pragma omp parallel proc_bind(primary) // lt51-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}}
+  ++argc;
   return 0;
 }
Index: clang/test/OpenMP/parallel_proc_bind_codegen.cpp
===================================================================
--- clang/test/OpenMP/parallel_proc_bind_codegen.cpp
+++ clang/test/OpenMP/parallel_proc_bind_codegen.cpp
@@ -1,20 +1,22 @@
-// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
-
-// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s --check-prefix CK1
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK1
+
 typedef __INTPTR_TYPE__ intptr_t;
 
-// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
-// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
-// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+// CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
 
 void foo();
 
@@ -40,18 +42,70 @@
   return tmain<int>();
 }
 
-// CHECK-LABEL: @main
-// CHECK:       [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
-// CHECK:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 4)
-// CHECK:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 3)
-// CHECK:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1-LABEL: @main
+// CK1:       [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
+// CK1:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 4)
+// CK1:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 3)
+// CK1:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+
+// CK1-LABEL: @{{.+}}tmain
+// CK1:       [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
+// CK1:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 2)
+// CK1:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1:       ret i32 0
+// CK1-NEXT:  }
+
+#endif
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s --check-prefix CK2
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK2
+
+typedef __INTPTR_TYPE__ intptr_t;
 
-// CHECK-LABEL: @{{.+}}tmain
-// CHECK:       [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
-// CHECK:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 2)
-// CHECK:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK:       ret i32 0
-// CHECK-NEXT:  }
+// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+
+void foo();
 
+struct S {
+  intptr_t a, b, c;
+  S(intptr_t a) : a(a) {}
+  operator char() { return a; }
+  ~S() {}
+};
+
+template <typename T>
+T tmain() {
+#pragma omp parallel proc_bind(primary)
+  foo();
+  return T();
+}
+
+int main() {
+#pragma omp parallel proc_bind(primary)
+  foo();
+  return tmain<int>();
+}
+
+// CK2-LABEL: @main
+// CK2:       [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
+// CK2:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 5)
+// CK2:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+
+// CK2-LABEL: @{{.+}}tmain
+// CK2:       [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
+// CK2:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 5)
+// CK2:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2:       ret i32 0
+// CK2-NEXT:  }
+
+#endif
 #endif
Index: clang/test/OpenMP/parallel_master_codegen.cpp
===================================================================
--- clang/test/OpenMP/parallel_master_codegen.cpp
+++ clang/test/OpenMP/parallel_master_codegen.cpp
@@ -647,5 +647,64 @@
 // CK9:       store i32 [[INC]], i32* [[A_FP_ADDR]]
 // CK9-NOT:   __kmpc_global_thread_num
 // CK9:       call void @__kmpc_end_master({{.+}})
+#endif
+#ifdef CK10
+///==========================================================================///
+// RUN: %clang_cc1 -DCK10 -verify -fopenmp -fopenmp-version=51 -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s --check-prefix CK10
+// RUN: %clang_cc1 -DCK10 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK10 -fopenmp -fopenmp-version=51 -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK10
+
+// RUN: %clang_cc1 -DCK10 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK10 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK10 -fopenmp-simd -fopenmp-version=51 -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+typedef __INTPTR_TYPE__ intptr_t;
+
+// CK10-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK10-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK10-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+
+void foo();
+
+struct S {
+  intptr_t a, b, c;
+  S(intptr_t a) : a(a) {}
+  operator char() { return a; }
+  ~S() {}
+};
+
+template <typename T>
+T tmain() {
+#pragma omp parallel master proc_bind(primary)
+  foo();
+  return T();
+}
+
+int main() {
+#pragma omp parallel master proc_bind(spread)
+  foo();
+#pragma omp parallel master proc_bind(close)
+  foo();
+  return tmain<int>();
+}
+
+// CK10-LABEL: @main
+// CK10:       [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
+// CK10:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 4)
+// CK10:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK10:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 3)
+// CK10:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+
+// CK10-LABEL: @{{.+}}tmain
+// CK10:       [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
+// CK10:       call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 5)
+// CK10:       call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK10:       ret i32 0
+// CK10-NEXT:  }
+
+// CK10:       call i32 @__kmpc_master(%struct.ident_t* [[DEF_LOC_2]], i32 [[ONE:%.+]])
+// CK10:       call void @__kmpc_end_master(%struct.ident_t* [[DEF_LOC_2]], i32 [[ONE]])
+
 #endif
 #endif
Index: clang/test/OpenMP/parallel_ast_print.cpp
===================================================================
--- clang/test/OpenMP/parallel_ast_print.cpp
+++ clang/test/OpenMP/parallel_ast_print.cpp
@@ -5,6 +5,14 @@
 // RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
+
+// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
 // expected-no-diagnostics
 
 #ifndef HEADER
@@ -152,6 +160,10 @@
   a=2;
 #pragma omp parallel default(none), private(argc,b) firstprivate(argv) shared (d) if (parallel:argc > 0) num_threads(C) copyin(S<T>::TS, thrp) proc_bind(master) reduction(+:c, arr1[argc]) reduction(max:e, arr[:C][0:10])
   foo();
+#ifdef OMP51
+#pragma omp parallel default(none), private(argc,b) firstprivate(argv) shared (d) if (parallel:argc > 0) num_threads(C) copyin(S<T>::TS, thrp) proc_bind(primary) reduction(+:c, arr1[argc]) reduction(max:e, arr[:C][0:10])
+  foo();
+#endif
 #pragma omp parallel if (C) num_threads(s) proc_bind(close) reduction(^:e, f, arr[0:C][:argc]) reduction(default, && : g) reduction(task,+:argc)
   foo();
   return 0;
@@ -166,6 +178,8 @@
 // CHECK-NEXT: a = 2;
 // CHECK-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(C) copyin(S<T>::TS,thrp) proc_bind(master) reduction(+: c,arr1[argc]) reduction(max: e,arr[:C][0:10])
 // CHECK-NEXT: foo()
+// OMP51-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(C) copyin(S<T>::TS,thrp) proc_bind(primary) reduction(+: c,arr1[argc]) reduction(max: e,arr[:C][0:10])
+// OMP51-NEXT: foo()
 // CHECK-NEXT: #pragma omp parallel if(C) num_threads(s) proc_bind(close) reduction(^: e,f,arr[0:C][:argc]) reduction(default, &&: g) reduction(task, +: argc)
 // CHECK-NEXT: foo()
 // CHECK: template<> int tmain<int, 5>(int argc, int *argv) {
@@ -177,6 +191,8 @@
 // CHECK-NEXT: a = 2;
 // CHECK-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(5) copyin(S<int>::TS,thrp) proc_bind(master) reduction(+: c,arr1[argc]) reduction(max: e,arr[:5][0:10])
 // CHECK-NEXT: foo()
+// OMP51-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(5) copyin(S<int>::TS,thrp) proc_bind(primary) reduction(+: c,arr1[argc]) reduction(max: e,arr[:5][0:10])
+// OMP51-NEXT: foo()
 // CHECK-NEXT: #pragma omp parallel if(5) num_threads(s) proc_bind(close) reduction(^: e,f,arr[0:5][:argc]) reduction(default, &&: g) reduction(task, +: argc)
 // CHECK-NEXT: foo()
 // CHECK: template<> long tmain<long, 1>(long argc, long *argv) {
@@ -188,6 +204,8 @@
 // CHECK-NEXT: a = 2;
 // CHECK-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(1) copyin(S<long>::TS,thrp) proc_bind(master) reduction(+: c,arr1[argc]) reduction(max: e,arr[:1][0:10])
 // CHECK-NEXT: foo()
+// OMP51-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(1) copyin(S<long>::TS,thrp) proc_bind(primary) reduction(+: c,arr1[argc]) reduction(max: e,arr[:1][0:10])
+// OMP51-NEXT: foo()
 // CHECK-NEXT: #pragma omp parallel if(1) num_threads(s) proc_bind(close) reduction(^: e,f,arr[0:1][:argc]) reduction(default, &&: g) reduction(task, +: argc)
 // CHECK-NEXT: foo()
 
Index: clang/test/OpenMP/nvptx_target_parallel_proc_bind_codegen.cpp
===================================================================
--- clang/test/OpenMP/nvptx_target_parallel_proc_bind_codegen.cpp
+++ clang/test/OpenMP/nvptx_target_parallel_proc_bind_codegen.cpp
@@ -1,24 +1,24 @@
-// Test target codegen - host bc file has to be created first.
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
-
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
-
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
+#ifdef CK1
 // Check that the execution mode of all 3 target regions on the gpu is set to SPMD Mode.
-// CHECK-DAG: {{@__omp_offloading_.+l29}}_exec_mode = weak constant i8 0
-// CHECK-DAG: {{@__omp_offloading_.+l33}}_exec_mode = weak constant i8 0
-// CHECK-DAG: {{@__omp_offloading_.+l38}}_exec_mode = weak constant i8 0
+// CK1-DAG: {{@__omp_offloading_.+l29}}_exec_mode = weak constant i8 0
+// CK1-DAG: {{@__omp_offloading_.+l33}}_exec_mode = weak constant i8 0
+// CK1-DAG: {{@__omp_offloading_.+l38}}_exec_mode = weak constant i8 0
 
 template<typename tx>
 tx ftemplate(int n) {
@@ -53,57 +53,114 @@
   return a;
 }
 
-// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l29}}(
-// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
-// CHECK: call void @__kmpc_data_sharing_init_stack_spmd
-// CHECK: br label {{%?}}[[EXEC:.+]]
+// CK1-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l29}}(
+// CK1: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
+// CK1: call void @__kmpc_data_sharing_init_stack_spmd
+// CK1: br label {{%?}}[[EXEC:.+]]
 //
-// CHECK: [[EXEC]]
-// CHECK-NOT: call void @__kmpc_push_proc_bind
-// CHECK: {{call|invoke}} void [[OP1:@.+]](
-// CHECK: br label {{%?}}[[DONE:.+]]
+// CK1: [[EXEC]]
+// CK1-NOT: call void @__kmpc_push_proc_bind
+// CK1: {{call|invoke}} void [[OP1:@.+]](
+// CK1: br label {{%?}}[[DONE:.+]]
 //
-// CHECK: [[DONE]]
-// CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
-// CHECK: br label {{%?}}[[EXIT:.+]]
+// CK1: [[DONE]]
+// CK1: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
+// CK1: br label {{%?}}[[EXIT:.+]]
 //
-// CHECK: [[EXIT]]
-// CHECK: ret void
-// CHECK: }
-
-// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l33}}(
-// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
-// CHECK: call void @__kmpc_data_sharing_init_stack_spmd
-// CHECK: br label {{%?}}[[EXEC:.+]]
+// CK1: [[EXIT]]
+// CK1: ret void
+// CK1: }
+
+// CK1-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l33}}(
+// CK1: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
+// CK1: call void @__kmpc_data_sharing_init_stack_spmd
+// CK1: br label {{%?}}[[EXEC:.+]]
 //
-// CHECK: [[EXEC]]
-// CHECK-NOT: call void @__kmpc_push_proc_bind
-// CHECK: {{call|invoke}} void [[OP1:@.+]](
-// CHECK: br label {{%?}}[[DONE:.+]]
+// CK1: [[EXEC]]
+// CK1-NOT: call void @__kmpc_push_proc_bind
+// CK1: {{call|invoke}} void [[OP1:@.+]](
+// CK1: br label {{%?}}[[DONE:.+]]
 //
-// CHECK: [[DONE]]
-// CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
-// CHECK: br label {{%?}}[[EXIT:.+]]
+// CK1: [[DONE]]
+// CK1: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
+// CK1: br label {{%?}}[[EXIT:.+]]
 //
-// CHECK: [[EXIT]]
-// CHECK: ret void
-// CHECK: }
-
-// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l38}}(
-// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
-// CHECK: call void @__kmpc_data_sharing_init_stack_spmd
-// CHECK: br label {{%?}}[[EXEC:.+]]
+// CK1: [[EXIT]]
+// CK1: ret void
+// CK1: }
+
+// CK1-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l38}}(
+// CK1: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
+// CK1: call void @__kmpc_data_sharing_init_stack_spmd
+// CK1: br label {{%?}}[[EXEC:.+]]
 //
-// CHECK: [[EXEC]]
-// CHECK-NOT: call void @__kmpc_push_proc_bind
-// CHECK: {{call|invoke}} void [[OP1:@.+]](
-// CHECK: br label {{%?}}[[DONE:.+]]
+// CK1: [[EXEC]]
+// CK1-NOT: call void @__kmpc_push_proc_bind
+// CK1: {{call|invoke}} void [[OP1:@.+]](
+// CK1: br label {{%?}}[[DONE:.+]]
 //
-// CHECK: [[DONE]]
-// CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
-// CHECK: br label {{%?}}[[EXIT:.+]]
+// CK1: [[DONE]]
+// CK1: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
+// CK1: br label {{%?}}[[EXIT:.+]]
 //
-// CHECK: [[EXIT]]
-// CHECK: ret void
-// CHECK: }
+// CK1: [[EXIT]]
+// CK1: ret void
+// CK1: }
+#endif
+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
+#ifdef CK2
+// Check that the execution mode of all 3 target regions on the gpu is set to SPMD Mode.
+// CK2-DAG: {{@__omp_offloading_.+l132}}_exec_mode = weak constant i8 0
+
+template<typename tx>
+tx ftemplate(int n) {
+  tx a = 0;
+  short aa = 0;
+  tx b[10];
+
+  #pragma omp target parallel proc_bind(master)
+  {
+  }
+
+  return a;
+}
+
+int bar(int n){
+  int a = 0;
+
+  a += ftemplate<int>(n);
+
+  return a;
+}
+
+// CK2-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l132}}(
+// CK2: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
+// CK2: call void @__kmpc_data_sharing_init_stack_spmd
+// CK2: br label {{%?}}[[EXEC:.+]]
+//
+// CK2: [[EXEC]]
+// CK2-NOT: call void @__kmpc_push_proc_bind
+// CK2: {{call|invoke}} void [[OP1:@.+]](
+// CK2: br label {{%?}}[[DONE:.+]]
+//
+// CK2: [[DONE]]
+// CK2: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
+// CK2: br label {{%?}}[[EXIT:.+]]
+//
+// CK2: [[EXIT]]
+// CK2: ret void
+// CK2: }
+
+#endif
 #endif
Index: clang/test/OpenMP/distribute_parallel_for_proc_bind_codegen.cpp
===================================================================
--- clang/test/OpenMP/distribute_parallel_for_proc_bind_codegen.cpp
+++ clang/test/OpenMP/distribute_parallel_for_proc_bind_codegen.cpp
@@ -1,22 +1,23 @@
 // add -fopenmp-targets
-
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
-
-// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
+
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK1
+
 typedef __INTPTR_TYPE__ intptr_t;
 
-// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
-// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
-// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+// CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
 
 void foo();
 
@@ -37,7 +38,7 @@
 }
 
 int main() {
-  // CHECK-LABEL: @main
+  // CK1-LABEL: @main
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute parallel for proc_bind(spread)
@@ -49,50 +50,145 @@
   return tmain<int>();
 }
 
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL1:@.+]]()
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL2:@.+]]()
-// CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
-// CHECK: ret i32 [[CALL_RET]]
-
-// CHECK: define{{.+}} void [[OFFL1]](
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: define{{.+}} [[OFFL2]]()
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
-
-// CHECK: define{{.+}} [[TMAIN]]()
-// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
-// CHECK: call void [[OFFL3:@.+]]()
-
-// CHECK: define{{.+}} [[OFFL3]]()
-// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
-
-// CHECK: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
-// CHECK: [[GTID_ADDR:%.+]] = alloca i32*,
-// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
-// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
-// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
-// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
-// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
-// CHECK: ret void
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL1:@.+]]()
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL2:@.+]]()
+// CK1: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
+// CK1: ret i32 [[CALL_RET]]
+
+// CK1: define{{.+}} void [[OFFL1]](
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: define{{.+}} [[OFFL2]]()
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+
+// CK1: define{{.+}} [[TMAIN]]()
+// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK1: call void [[OFFL3:@.+]]()
+
+// CK1: define{{.+}} [[OFFL3]]()
+// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
+
+// CK1: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
+// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK1: ret void
+#endif
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK2
+
+typedef __INTPTR_TYPE__ intptr_t;
+
+// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+
+void foo();
+
+struct S {
+  intptr_t a, b, c;
+  S(intptr_t a) : a(a) {}
+  operator char() { return a; }
+  ~S() {}
+};
+
+template <typename T>
+T tmain() {
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for proc_bind(primary)
+  for(int i = 0; i < 1000; i++) {}
+  return T();
+}
+
+int main() {
+  // CK2-LABEL: @main
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for proc_bind(spread)
+  for(int i = 0; i < 1000; i++) {}
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for proc_bind(close)
+  for(int i = 0; i < 1000; i++) {}
+  return tmain<int>();
+}
+
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL1:@.+]]()
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL2:@.+]]()
+// CK2: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
+// CK2: ret i32 [[CALL_RET]]
+
+// CK2: define{{.+}} void [[OFFL1]](
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+
+// CK2: define{{.+}} [[OFFL2]]()
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+
+// CK2: define{{.+}} [[TMAIN]]()
+// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
+// CK2: call void [[OFFL3:@.+]]()
+
+// CK2: define{{.+}} [[OFFL3]]()
+// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
+
+// CK2: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
+// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
+// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
+// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
+// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
+// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
+// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
+// CK2: ret void
+#endif
 #endif
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -14069,10 +14069,21 @@
     Diag(KindKwLoc, diag::err_omp_unexpected_clause_value)
         << getListOfPossibleValues(OMPC_proc_bind,
                                    /*First=*/unsigned(OMP_PROC_BIND_master),
-                                   /*Last=*/5)
+                                   /*Last=*/
+                                   unsigned(LangOpts.OpenMP > 50
+                                                ? OMP_PROC_BIND_primary
+                                                : OMP_PROC_BIND_spread) +
+                                       1)
         << getOpenMPClauseName(OMPC_proc_bind);
     return nullptr;
   }
+  if (Kind == OMP_PROC_BIND_primary && LangOpts.OpenMP < 51)
+    Diag(KindKwLoc, diag::err_omp_unexpected_clause_value)
+        << getListOfPossibleValues(OMPC_proc_bind,
+                                   /*First=*/unsigned(OMP_PROC_BIND_master),
+                                   /*Last=*/
+                                   unsigned(OMP_PROC_BIND_spread) + 1)
+        << getOpenMPClauseName(OMPC_proc_bind);
   return new (Context)
       OMPProcBindClause(Kind, KindKwLoc, StartLoc, LParenLoc, EndLoc);
 }

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

[PATCH] D99622: [OpenMP51] Accept `primary` as proc bind affinity policy in Clang

Reply via email to