r341073 - [OPENMP][NVPTX] Add options -f[no-]openmp-cuda-force-full-runtime.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Thu Aug 30 07:45:24 PDT 2018
Author: abataev
Date: Thu Aug 30 07:45:24 2018
New Revision: 341073
URL: http://llvm.org/viewvc/llvm-project?rev=341073&view=rev
Log:
[OPENMP][NVPTX] Add options -f[no-]openmp-cuda-force-full-runtime.
Added options -f[no-]openmp-cuda-force-full-runtime to [not] force use
of the full runtime for OpenMP offloading to CUDA devices.
Added:
cfe/trunk/test/OpenMP/nvptx_force_full_runtime_SPMD_codegen.cpp
Modified:
cfe/trunk/include/clang/Basic/LangOptions.def
cfe/trunk/include/clang/Driver/Options.td
cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
cfe/trunk/lib/Driver/ToolChains/Clang.cpp
cfe/trunk/lib/Frontend/CompilerInvocation.cpp
cfe/trunk/test/Driver/openmp-offload-gpu.c
Modified: cfe/trunk/include/clang/Basic/LangOptions.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.def?rev=341073&r1=341072&r2=341073&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.def (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.def Thu Aug 30 07:45:24 2018
@@ -203,6 +203,7 @@ LANGOPT(OpenMPSimd , 1, 0, "Use S
LANGOPT(OpenMPUseTLS , 1, 0, "Use TLS for threadprivates or runtime calls")
LANGOPT(OpenMPIsDevice , 1, 0, "Generate code only for OpenMP target device")
LANGOPT(OpenMPCUDAMode , 1, 0, "Generate code for OpenMP pragmas in SIMT/SPMD mode")
+LANGOPT(OpenMPCUDAForceFullRuntime , 1, 0, "Force to use full runtime in all constructs when offloading to CUDA devices")
LANGOPT(OpenMPHostCXXExceptions , 1, 0, "C++ exceptions handling in the host code.")
LANGOPT(RenderScript , 1, 0, "RenderScript")
Modified: cfe/trunk/include/clang/Driver/Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/Options.td?rev=341073&r1=341072&r2=341073&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/Options.td (original)
+++ cfe/trunk/include/clang/Driver/Options.td Thu Aug 30 07:45:24 2018
@@ -1531,6 +1531,10 @@ def fopenmp_cuda_mode : Flag<["-"], "fop
Flags<[CC1Option, NoArgumentUnused, HelpHidden]>;
def fno_openmp_cuda_mode : Flag<["-"], "fno-openmp-cuda-mode">, Group<f_Group>,
Flags<[NoArgumentUnused, HelpHidden]>;
+def fopenmp_cuda_force_full_runtime : Flag<["-"], "fopenmp-cuda-force-full-runtime">, Group<f_Group>,
+ Flags<[CC1Option, NoArgumentUnused, HelpHidden]>;
+def fno_openmp_cuda_force_full_runtime : Flag<["-"], "fno-openmp-cuda-force-full-runtime">, Group<f_Group>,
+ Flags<[NoArgumentUnused, HelpHidden]>;
def fno_optimize_sibling_calls : Flag<["-"], "fno-optimize-sibling-calls">, Group<f_Group>;
def foptimize_sibling_calls : Flag<["-"], "foptimize-sibling-calls">, Group<f_Group>;
def fno_escaping_block_tail_calls : Flag<["-"], "fno-escaping-block-tail-calls">, Group<f_Group>, Flags<[CC1Option]>;
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp?rev=341073&r1=341072&r2=341073&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp Thu Aug 30 07:45:24 2018
@@ -1218,7 +1218,8 @@ void CGOpenMPRuntimeNVPTX::emitSPMDEntry
EST.ExitBB = CGF.createBasicBlock(".exit");
// Initialize the OMP state in the runtime; called by all active threads.
- bool RequiresFullRuntime = !supportsLightweightRuntime(CGF.getContext(), D);
+ bool RequiresFullRuntime = CGM.getLangOpts().OpenMPCUDAForceFullRuntime ||
+ !supportsLightweightRuntime(CGF.getContext(), D);
llvm::Value *Args[] = {
getThreadLimit(CGF, /*IsInSPMDExecutionMode=*/true),
/*RequiresOMPRuntime=*/
Modified: cfe/trunk/lib/Driver/ToolChains/Clang.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Driver/ToolChains/Clang.cpp?rev=341073&r1=341072&r2=341073&view=diff
==============================================================================
--- cfe/trunk/lib/Driver/ToolChains/Clang.cpp (original)
+++ cfe/trunk/lib/Driver/ToolChains/Clang.cpp Thu Aug 30 07:45:24 2018
@@ -4039,8 +4039,16 @@ void Clang::ConstructJob(Compilation &C,
// When in OpenMP offloading mode with NVPTX target, forward
// cuda-mode flag
- Args.AddLastArg(CmdArgs, options::OPT_fopenmp_cuda_mode,
- options::OPT_fno_openmp_cuda_mode);
+ if (Args.hasFlag(options::OPT_fopenmp_cuda_mode,
+ options::OPT_fno_openmp_cuda_mode, /*Default=*/false))
+ CmdArgs.push_back("-fopenmp-cuda-mode");
+
+ // When in OpenMP offloading mode with NVPTX target, check if full runtime
+ // is required.
+ if (Args.hasFlag(options::OPT_fopenmp_cuda_force_full_runtime,
+ options::OPT_fno_openmp_cuda_force_full_runtime,
+ /*Default=*/false))
+ CmdArgs.push_back("-fopenmp-cuda-force-full-runtime");
break;
default:
// By default, if Clang doesn't know how to generate useful OpenMP code
Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=341073&r1=341072&r2=341073&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Thu Aug 30 07:45:24 2018
@@ -2677,10 +2677,15 @@ static void ParseLangArgs(LangOptions &O
<< Opts.OMPHostIRFile;
}
- // set CUDA mode for OpenMP target NVPTX if specified in options
+ // Set CUDA mode for OpenMP target NVPTX if specified in options
Opts.OpenMPCUDAMode = Opts.OpenMPIsDevice && T.isNVPTX() &&
Args.hasArg(options::OPT_fopenmp_cuda_mode);
+ // Set CUDA mode for OpenMP target NVPTX if specified in options
+ Opts.OpenMPCUDAForceFullRuntime =
+ Opts.OpenMPIsDevice && T.isNVPTX() &&
+ Args.hasArg(options::OPT_fopenmp_cuda_force_full_runtime);
+
// Record whether the __DEPRECATED define was requested.
Opts.Deprecated = Args.hasFlag(OPT_fdeprecated_macro,
OPT_fno_deprecated_macro,
Modified: cfe/trunk/test/Driver/openmp-offload-gpu.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Driver/openmp-offload-gpu.c?rev=341073&r1=341072&r2=341073&view=diff
==============================================================================
--- cfe/trunk/test/Driver/openmp-offload-gpu.c (original)
+++ cfe/trunk/test/Driver/openmp-offload-gpu.c Thu Aug 30 07:45:24 2018
@@ -216,3 +216,26 @@
// HAS_DEBUG: nvlink
// HAS_DEBUG-SAME: "-g"
+// RUN: %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-mode 2>&1 \
+// RUN: | FileCheck -check-prefix=CUDA_MODE %s
+// RUN: %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-mode -fopenmp-cuda-mode 2>&1 \
+// RUN: | FileCheck -check-prefix=CUDA_MODE %s
+// CUDA_MODE: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"
+// CUDA_MODE-SAME: "-fopenmp-cuda-mode"
+// RUN: %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-mode 2>&1 \
+// RUN: | FileCheck -check-prefix=NO_CUDA_MODE %s
+// RUN: %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-mode -fno-openmp-cuda-mode 2>&1 \
+// RUN: | FileCheck -check-prefix=NO_CUDA_MODE %s
+// NO_CUDA_MODE-NOT: "-{{fno-|f}}openmp-cuda-mode"
+
+// RUN: %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-force-full-runtime 2>&1 \
+// RUN: | FileCheck -check-prefix=FULL_RUNTIME %s
+// RUN: %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-force-full-runtime -fopenmp-cuda-force-full-runtime 2>&1 \
+// RUN: | FileCheck -check-prefix=FULL_RUNTIME %s
+// FULL_RUNTIME: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"
+// FULL_RUNTIME-SAME: "-fopenmp-cuda-force-full-runtime"
+// RUN: %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fno-openmp-cuda-force-full-runtime 2>&1 \
+// RUN: | FileCheck -check-prefix=NO_FULL_RUNTIME %s
+// RUN: %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-force-full-runtime -fno-openmp-cuda-force-full-runtime 2>&1 \
+// RUN: | FileCheck -check-prefix=NO_FULL_RUNTIME %s
+// NO_FULL_RUNTIME-NOT: "-{{fno-|f}}openmp-cuda-force-full-runtime"
Added: cfe/trunk/test/OpenMP/nvptx_force_full_runtime_SPMD_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_force_full_runtime_SPMD_codegen.cpp?rev=341073&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_force_full_runtime_SPMD_codegen.cpp (added)
+++ cfe/trunk/test/OpenMP/nvptx_force_full_runtime_SPMD_codegen.cpp Thu Aug 30 07:45:24 2018
@@ -0,0 +1,328 @@
+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -verify -fopenmp -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 -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 - -fopenmp-cuda-force-full-runtime | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -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 -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 - -fopenmp-cuda-force-full-runtime | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -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 - -fopenmp-cuda-force-full-runtime | FileCheck %s
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+// CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1
+// CHECK: @__omp_offloading_{{.+}}_l52_exec_mode = weak constant i8 1
+// CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1
+
+void foo() {
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+#pragma omp target teams distribute parallel for simd
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for simd schedule(static)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for simd schedule(static, 1)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for simd schedule(auto)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for simd schedule(runtime)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for simd schedule(dynamic)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for simd schedule(guided)
+ for (int i = 0; i < 10; ++i)
+ ;
+int a;
+// CHECK: call void @__kmpc_kernel_init(
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+#pragma omp target teams distribute parallel for lastprivate(a)
+ for (int i = 0; i < 10; ++i)
+ a = i;
+#pragma omp target teams distribute parallel for schedule(static)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for schedule(static, 1)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for schedule(auto)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for schedule(runtime)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for schedule(dynamic)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams distribute parallel for schedule(guided)
+ for (int i = 0; i < 10; ++i)
+ ;
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+#pragma omp target teams
+#pragma omp distribute parallel for simd
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for simd schedule(static)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for simd schedule(static, 1)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for simd schedule(auto)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for simd schedule(runtime)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for simd schedule(dynamic)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for simd schedule(guided)
+ for (int i = 0; i < 10; ++i)
+ ;
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+#pragma omp target teams
+#pragma omp distribute parallel for
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for schedule(static)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for schedule(static, 1)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for schedule(auto)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for schedule(runtime)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for schedule(dynamic)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target teams
+#pragma omp distribute parallel for schedule(guided)
+ for (int i = 0; i < 10; ++i)
+ ;
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for schedule(static)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for schedule(static, 1)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for schedule(auto)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for schedule(runtime)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for schedule(dynamic)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for schedule(guided)
+ for (int i = 0; i < 10; ++i)
+ ;
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+#pragma omp target parallel for
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel for schedule(static)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel for schedule(static, 1)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel for schedule(auto)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel for schedule(runtime)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel for schedule(dynamic)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel for schedule(guided)
+ for (int i = 0; i < 10; ++i)
+ ;
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+#pragma omp target parallel
+#pragma omp for simd
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel
+#pragma omp for simd schedule(static)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel
+#pragma omp for simd schedule(static, 1)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel
+#pragma omp for simd schedule(auto)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel
+#pragma omp for simd schedule(runtime)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel
+#pragma omp for simd schedule(dynamic)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target parallel
+#pragma omp for simd schedule(guided)
+ for (int i = 0; i < 10; ++i)
+ ;
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+#pragma omp target
+#pragma omp parallel
+#pragma omp for simd ordered
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel
+#pragma omp for simd schedule(static)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel
+#pragma omp for simd schedule(static, 1)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel
+#pragma omp for simd schedule(auto)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel
+#pragma omp for simd schedule(runtime)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel
+#pragma omp for simd schedule(dynamic)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel
+#pragma omp for simd schedule(guided)
+ for (int i = 0; i < 10; ++i)
+ ;
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}})
+#pragma omp target
+#pragma omp parallel for
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel for schedule(static)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel for schedule(static, 1)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel for schedule(auto)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel for schedule(runtime)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel for schedule(dynamic)
+ for (int i = 0; i < 10; ++i)
+ ;
+#pragma omp target
+#pragma omp parallel for schedule(guided)
+ for (int i = 0; i < 10; ++i)
+ ;
+}
+
+#endif
+
More information about the cfe-commits
mailing list