maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OpenMP] Remove simplified device runtime handling 

The old device runtime had a "simplified" version that prevented many of
the runtime features from being initialized. The old device runtime was
deleted in LLVM 14 and is no longer in use. Selectively deactivating
features is now done using specific flags rather than the old technique.
This patch simply removes the extra logic required for handling the old
simple runtime scheme.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D133802
This commit is contained in:
Joseph Huber 2022-09-13 14:38:00 -05:00
parent b1cd393f9e
commit 2d26ecb1fb
17 changed files with 475 additions and 898 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
clang/include/clang/Basic/LangOptions.def
View File

@ -246,7 +246,6 @@ LANGOPT(OpenMPUseTLS , 1, 0, "Use TLS for threadprivates or runtime calls")
LANGOPT(OpenMPIsDevice , 1, 0, "Generate code only for OpenMP target device") 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(OpenMPCUDAMode , 1, 0, "Generate code for OpenMP pragmas in SIMT/SPMD mode")
LANGOPT(OpenMPIRBuilder , 1, 0, "Use the experimental OpenMP-IR-Builder codegen path.") LANGOPT(OpenMPIRBuilder , 1, 0, "Use the experimental OpenMP-IR-Builder codegen path.")
LANGOPT(OpenMPCUDAForceFullRuntime , 1, 0, "Force to use full runtime in all constructs when offloading to CUDA devices")
LANGOPT(OpenMPCUDANumSMs , 32, 0, "Number of SMs for CUDA devices.") LANGOPT(OpenMPCUDANumSMs , 32, 0, "Number of SMs for CUDA devices.")
LANGOPT(OpenMPCUDABlocksPerSM , 32, 0, "Number of blocks per SM for CUDA devices.") LANGOPT(OpenMPCUDABlocksPerSM , 32, 0, "Number of blocks per SM for CUDA devices.")
LANGOPT(OpenMPCUDAReductionBufNum , 32, 1024, "Number of the reduction records in the intermediate reduction buffer used for the teams reductions.") LANGOPT(OpenMPCUDAReductionBufNum , 32, 1024, "Number of the reduction records in the intermediate reduction buffer used for the teams reductions.")

4
clang/include/clang/Driver/Options.td
View File

@ -2562,10 +2562,6 @@ def fopenmp_cuda_mode : Flag<["-"], "fopenmp-cuda-mode">, Group<f_Group>,
Flags<[CC1Option, NoArgumentUnused, HelpHidden]>; Flags<[CC1Option, NoArgumentUnused, HelpHidden]>;
def fno_openmp_cuda_mode : Flag<["-"], "fno-openmp-cuda-mode">, Group<f_Group>, def fno_openmp_cuda_mode : Flag<["-"], "fno-openmp-cuda-mode">, Group<f_Group>,
Flags<[NoArgumentUnused, HelpHidden]>; 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 fopenmp_cuda_number_of_sm_EQ : Joined<["-"], "fopenmp-cuda-number-of-sm=">, Group<f_Group>, def fopenmp_cuda_number_of_sm_EQ : Joined<["-"], "fopenmp-cuda-number-of-sm=">, Group<f_Group>,
Flags<[CC1Option, NoArgumentUnused, HelpHidden]>; Flags<[CC1Option, NoArgumentUnused, HelpHidden]>;
def fopenmp_cuda_blocks_per_sm_EQ : Joined<["-"], "fopenmp-cuda-blocks-per-sm=">, Group<f_Group>, def fopenmp_cuda_blocks_per_sm_EQ : Joined<["-"], "fopenmp-cuda-blocks-per-sm=">, Group<f_Group>,

39
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
View File

@ -73,30 +73,15 @@ private:
CGOpenMPRuntimeGPU::ExecutionMode SavedExecMode = CGOpenMPRuntimeGPU::ExecutionMode SavedExecMode =
CGOpenMPRuntimeGPU::EM_Unknown; CGOpenMPRuntimeGPU::EM_Unknown;
CGOpenMPRuntimeGPU::ExecutionMode &ExecMode; CGOpenMPRuntimeGPU::ExecutionMode &ExecMode;
bool SavedRuntimeMode = false;
bool *RuntimeMode = nullptr;
public: public:
/// Constructor for Non-SPMD mode. ExecutionRuntimeModesRAII(CGOpenMPRuntimeGPU::ExecutionMode &ExecMode,
ExecutionRuntimeModesRAII(CGOpenMPRuntimeGPU::ExecutionMode &ExecMode) CGOpenMPRuntimeGPU::ExecutionMode EntryMode)
: ExecMode(ExecMode) { : ExecMode(ExecMode) {
SavedExecMode = ExecMode; SavedExecMode = ExecMode;
ExecMode = CGOpenMPRuntimeGPU::EM_NonSPMD; ExecMode = EntryMode;
}
/// Constructor for SPMD mode.
ExecutionRuntimeModesRAII(CGOpenMPRuntimeGPU::ExecutionMode &ExecMode,
bool &RuntimeMode, bool FullRuntimeMode)
: ExecMode(ExecMode), RuntimeMode(&RuntimeMode) {
SavedExecMode = ExecMode;
SavedRuntimeMode = RuntimeMode;
ExecMode = CGOpenMPRuntimeGPU::EM_SPMD;
RuntimeMode = FullRuntimeMode;
}
~ExecutionRuntimeModesRAII() {
ExecMode = SavedExecMode;
if (RuntimeMode)
*RuntimeMode = SavedRuntimeMode;
} }
~ExecutionRuntimeModesRAII() { ExecMode = SavedExecMode; }
}; };
/// GPU Configuration: This information can be derived from cuda registers, /// GPU Configuration: This information can be derived from cuda registers,
@ -1012,7 +997,7 @@ void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
llvm::Constant *&OutlinedFnID, llvm::Constant *&OutlinedFnID,
bool IsOffloadEntry, bool IsOffloadEntry,
const RegionCodeGenTy &CodeGen) { const RegionCodeGenTy &CodeGen) {
ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode); ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode, EM_NonSPMD);
EntryFunctionState EST; EntryFunctionState EST;
WrapperFunctionsMap.clear(); WrapperFunctionsMap.clear();
@ -1047,7 +1032,7 @@ void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
void CGOpenMPRuntimeGPU::emitKernelInit(CodeGenFunction &CGF, void CGOpenMPRuntimeGPU::emitKernelInit(CodeGenFunction &CGF,
EntryFunctionState &EST, bool IsSPMD) { EntryFunctionState &EST, bool IsSPMD) {
CGBuilderTy &Bld = CGF.Builder; CGBuilderTy &Bld = CGF.Builder;
Bld.restoreIP(OMPBuilder.createTargetInit(Bld, IsSPMD, requiresFullRuntime())); Bld.restoreIP(OMPBuilder.createTargetInit(Bld, IsSPMD, true));
IsInTargetMasterThreadRegion = IsSPMD; IsInTargetMasterThreadRegion = IsSPMD;
if (!IsSPMD) if (!IsSPMD)
emitGenericVarsProlog(CGF, EST.Loc); emitGenericVarsProlog(CGF, EST.Loc);
@ -1060,7 +1045,7 @@ void CGOpenMPRuntimeGPU::emitKernelDeinit(CodeGenFunction &CGF,
emitGenericVarsEpilog(CGF); emitGenericVarsEpilog(CGF);
CGBuilderTy &Bld = CGF.Builder; CGBuilderTy &Bld = CGF.Builder;
OMPBuilder.createTargetDeinit(Bld, IsSPMD, requiresFullRuntime()); OMPBuilder.createTargetDeinit(Bld, IsSPMD, true);
} }
void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D, void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
@ -1069,10 +1054,7 @@ void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
llvm::Constant *&OutlinedFnID, llvm::Constant *&OutlinedFnID,
bool IsOffloadEntry, bool IsOffloadEntry,
const RegionCodeGenTy &CodeGen) { const RegionCodeGenTy &CodeGen) {
ExecutionRuntimeModesRAII ModeRAII( ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode, EM_SPMD);
CurrentExecutionMode, RequiresFullRuntime,
CGM.getLangOpts().OpenMPCUDAForceFullRuntime ||
!supportsLightweightRuntime(CGM.getContext(), D));
EntryFunctionState EST; EntryFunctionState EST;
// Emit target region as a standalone region. // Emit target region as a standalone region.
@ -1184,11 +1166,8 @@ static const ModeFlagsTy UndefinedMode =
unsigned CGOpenMPRuntimeGPU::getDefaultLocationReserved2Flags() const { unsigned CGOpenMPRuntimeGPU::getDefaultLocationReserved2Flags() const {
switch (getExecutionMode()) { switch (getExecutionMode()) {
case EM_SPMD: case EM_SPMD:
if (requiresFullRuntime()) return KMP_IDENT_SPMD_MODE & (~KMP_IDENT_SIMPLE_RT_MODE);
return KMP_IDENT_SPMD_MODE & (~KMP_IDENT_SIMPLE_RT_MODE);
return KMP_IDENT_SPMD_MODE | KMP_IDENT_SIMPLE_RT_MODE;
case EM_NonSPMD: case EM_NonSPMD:
assert(requiresFullRuntime() && "Expected full runtime.");
return (~KMP_IDENT_SPMD_MODE) & (~KMP_IDENT_SIMPLE_RT_MODE); return (~KMP_IDENT_SPMD_MODE) & (~KMP_IDENT_SIMPLE_RT_MODE);
case EM_Unknown: case EM_Unknown:
return UndefinedMode; return UndefinedMode;

5
clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
View File

@ -42,8 +42,6 @@ private:
ExecutionMode getExecutionMode() const; ExecutionMode getExecutionMode() const;
bool requiresFullRuntime() const { return RequiresFullRuntime; }
/// Get barrier to synchronize all threads in a block. /// Get barrier to synchronize all threads in a block.
void syncCTAThreads(CodeGenFunction &CGF); void syncCTAThreads(CodeGenFunction &CGF);
@ -386,9 +384,6 @@ private:
/// to emit optimized code. /// to emit optimized code.
ExecutionMode CurrentExecutionMode = EM_Unknown; ExecutionMode CurrentExecutionMode = EM_Unknown;
/// Check if the full runtime is required (default - yes).
bool RequiresFullRuntime = true;
/// true if we're emitting the code for the target region and next parallel /// true if we're emitting the code for the target region and next parallel
/// region is L0 for sure. /// region is L0 for sure.
bool IsInTargetMasterThreadRegion = false; bool IsInTargetMasterThreadRegion = false;

7
clang/lib/Driver/ToolChains/Clang.cpp
View File

@ -6116,13 +6116,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
options::OPT_fno_openmp_target_debug, /*Default=*/false)) options::OPT_fno_openmp_target_debug, /*Default=*/false))
CmdArgs.push_back("-fopenmp-target-debug"); CmdArgs.push_back("-fopenmp-target-debug");
// 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");
// When in OpenMP offloading mode, forward assumptions information about // When in OpenMP offloading mode, forward assumptions information about
// thread and team counts in the device. // thread and team counts in the device.
if (Args.hasFlag(options::OPT_fopenmp_assume_teams_oversubscription, if (Args.hasFlag(options::OPT_fopenmp_assume_teams_oversubscription,

8
clang/lib/Frontend/CompilerInvocation.cpp
View File

@ -3479,9 +3479,6 @@ void CompilerInvocation::GenerateLangArgs(const LangOptions &Opts,
if (Opts.OpenMPCUDAMode) if (Opts.OpenMPCUDAMode)
GenerateArg(Args, OPT_fopenmp_cuda_mode, SA); GenerateArg(Args, OPT_fopenmp_cuda_mode, SA);
if (Opts.OpenMPCUDAForceFullRuntime)
GenerateArg(Args, OPT_fopenmp_cuda_force_full_runtime, SA);
// The arguments used to set Optimize, OptimizeSize and NoInlineDefine are // The arguments used to set Optimize, OptimizeSize and NoInlineDefine are
// generated from CodeGenOptions. // generated from CodeGenOptions.
@ -3933,11 +3930,6 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
Opts.OpenMPCUDAMode = Opts.OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN()) && Opts.OpenMPCUDAMode = Opts.OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN()) &&
Args.hasArg(options::OPT_fopenmp_cuda_mode); Args.hasArg(options::OPT_fopenmp_cuda_mode);
// Set CUDA mode for OpenMP target NVPTX/AMDGCN if specified in options
Opts.OpenMPCUDAForceFullRuntime =
Opts.OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN()) &&
Args.hasArg(options::OPT_fopenmp_cuda_force_full_runtime);
// FIXME: Eliminate this dependency. // FIXME: Eliminate this dependency.
unsigned Opt = getOptimizationLevel(Args, IK, Diags), unsigned Opt = getOptimizationLevel(Args, IK, Diags),
OptSize = getOptimizationLevelSize(Args); OptSize = getOptimizationLevelSize(Args);

20
clang/test/Driver/openmp-offload-gpu.c
View File

@ -207,26 +207,6 @@
// RUN: | FileCheck -check-prefix=NO_CUDA_MODE %s // RUN: | FileCheck -check-prefix=NO_CUDA_MODE %s
// NO_CUDA_MODE-NOT: "-{{fno-|f}}openmp-cuda-mode" // NO_CUDA_MODE-NOT: "-{{fno-|f}}openmp-cuda-mode"
// RUN: %clang -### -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 -### -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
// RUN: %clang -### -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fopenmp-cuda-force-full-runtime 2>&1 \
// RUN: | FileCheck -check-prefix=FULL_RUNTIME %s
// RUN: %clang -### -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fno-openmp-cuda-force-full-runtime -fopenmp-cuda-force-full-runtime 2>&1 \
// RUN: | FileCheck -check-prefix=FULL_RUNTIME %s
// FULL_RUNTIME: "-cc1"{{.*}}"-triple" "{{nvptx64-nvidia-cuda|amdgcn-amd-amdhsa}}"
// FULL_RUNTIME-SAME: "-fopenmp-cuda-force-full-runtime"
// RUN: %clang -### -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 -### -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
// RUN: %clang -### -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %s -fno-openmp-cuda-force-full-runtime 2>&1 \
// RUN: | FileCheck -check-prefix=NO_FULL_RUNTIME %s
// RUN: %clang -### -fopenmp=libomp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfx906 %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"
// RUN: %clang -### -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-teams-reduction-recs-num=2048 2>&1 \ // RUN: %clang -### -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_60 %s -fopenmp-cuda-teams-reduction-recs-num=2048 2>&1 \
// RUN: | FileCheck -check-prefix=CUDA_RED_RECS %s // RUN: | FileCheck -check-prefix=CUDA_RED_RECS %s
// CUDA_RED_RECS: "-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda" // CUDA_RED_RECS: "-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"

2
clang/test/OpenMP/amdgcn_target_codegen.cpp
View File

@ -27,7 +27,7 @@ int test_amdgcn_target_tid_threads_simd() {
int arr[N]; int arr[N];
// CHECK: call i32 @__kmpc_target_init(%struct.ident_t* addrspacecast (%struct.ident_t addrspace(1)* @1 to %struct.ident_t*), i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init(%struct.ident_t* addrspacecast (%struct.ident_t addrspace(1)* @1 to %struct.ident_t*), i8 2, i1 false, i1 true)
#pragma omp target simd #pragma omp target simd
for (int i = 0; i < N; i++) { for (int i = 0; i < N; i++) {
arr[i] = 1; arr[i] = 1;

165
clang/test/OpenMP/nvptx_SPMD_codegen.cpp
View File

@ -11,28 +11,21 @@
int a; int a;
// CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1 // CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1
// CHECK-DAG: [[DISTR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2050, i32 3, i32 {{[0-9]+}}, i8* getelementptr inbounds
// CHECK-DAG: [[FOR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 514, i32 3, i32 {{[0-9]+}}, i8* getelementptr inbounds
// CHECK-DAG: [[LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 3, i32 {{[0-9]+}}, i8* getelementptr inbounds
// CHECK-DAG: [[DISTR_FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2050, i32 1, i32 {{[0-9]+}}, i8* getelementptr inbounds // CHECK-DAG: [[DISTR_FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2050, i32 1, i32 {{[0-9]+}}, i8* getelementptr inbounds
// CHECK-DAG: [[FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 1, i32 {{[0-9]+}}, i8* getelementptr inbounds // CHECK-DAG: [[FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 1, i32 {{[0-9]+}}, i8* getelementptr inbounds
// CHECK-DAG: [[BAR_LIGHT:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 66, i32 3, i32 {{[0-9]+}}, i8* getelementptr inbounds
// CHECK-DAG: [[BAR_FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 66, i32 1, i32 {{[0-9]+}}, i8* getelementptr inbounds // CHECK-DAG: [[BAR_FULL:@.+]] = private unnamed_addr constant %struct.ident_t { i32 0, i32 66, i32 1, i32 {{[0-9]+}}, i8* getelementptr inbounds
// CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1 // CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1
void foo() { void foo() {
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_LIGHT]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FOR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[DISTR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[FOR_LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[LIGHT]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
@ -67,18 +60,15 @@ void foo() {
for (int i = 0; i < 10; ++i) for (int i = 0; i < 10; ++i)
; ;
int a; int a;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_LIGHT]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FOR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[DISTR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[FOR_LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[LIGHT]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
@ -112,15 +102,15 @@ int a;
#pragma omp target teams distribute parallel for schedule(guided) #pragma omp target teams distribute parallel for schedule(guided)
for (int i = 0; i < 10; ++i) for (int i = 0; i < 10; ++i)
; ;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_LIGHT]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FOR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[DISTR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init( // CHECK: call i32 @__kmpc_target_init(
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
@ -172,18 +162,15 @@ int a;
#pragma omp distribute parallel for simd schedule(guided) #pragma omp distribute parallel for simd schedule(guided)
for (int i = 0; i < 10; ++i) for (int i = 0; i < 10; ++i)
; ;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_LIGHT]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FOR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[DISTR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[FOR_LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[LIGHT]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
@ -224,18 +211,6 @@ int a;
#pragma omp distribute parallel for schedule(guided) #pragma omp distribute parallel for schedule(guided)
for (int i = 0; i < 10; ++i) for (int i = 0; i < 10; ++i)
; ;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[DISTR_LIGHT]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
@ -248,6 +223,14 @@ int a;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]] // CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[DISTR_FULL]]
// CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
#pragma omp target #pragma omp target
#pragma omp teams #pragma omp teams
#pragma omp distribute parallel for #pragma omp distribute parallel for
@ -283,15 +266,12 @@ int a;
#pragma omp distribute parallel for schedule(guided) #pragma omp distribute parallel for schedule(guided)
for (int i = 0; i < 10; ++i) for (int i = 0; i < 10; ++i)
; ;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FOR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[FOR_LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
@ -321,18 +301,15 @@ int a;
#pragma omp target parallel for schedule(guided) #pragma omp target parallel for schedule(guided)
for (int i = 0; i < 10; ++i) for (int i = 0; i < 10; ++i)
; ;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FOR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[LIGHT]] // CHECK-DAG: [[BAR_FULL]]
// CHECK-DAG: [[BAR_LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[FOR_LIGHT]] // CHECK-DAG: [[BAR_FULL]]
// CHECK-DAG: [[LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[BAR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[BAR_FULL]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]] // CHECK-DAG: [[BAR_FULL]]
@ -376,14 +353,6 @@ int a;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]] // CHECK-DAG: [[BAR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK-DAG: [[BAR_LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]] // CHECK-DAG: [[BAR_FULL]]
@ -396,6 +365,11 @@ int a;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]] // CHECK-DAG: [[BAR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
// CHECK-DAG: [[BAR_FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]]
#pragma omp target #pragma omp target
#pragma omp parallel #pragma omp parallel
#pragma omp for simd ordered #pragma omp for simd ordered
@ -431,15 +405,10 @@ int a;
#pragma omp for simd schedule(guided) #pragma omp for simd schedule(guided)
for (int i = 0; i < 10; ++i) for (int i = 0; i < 10; ++i)
; ;
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FOR_LIGHT]] // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[LIGHT]] // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK-DAG: [[FULL]]
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false)
// CHECK-DAG: [[FOR_LIGHT]]
// CHECK-DAG: [[LIGHT]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-DAG: [[FULL]] // CHECK-DAG: [[FULL]]
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)

8
clang/test/OpenMP/nvptx_distribute_parallel_generic_mode_codegen.cpp
View File

@ -45,7 +45,7 @@ int main(int argc, char **argv) {
// CHECK4-NEXT: [[TMP2:%.*]] = load i32*, i32** [[A_ADDR]], align 8 // CHECK4-NEXT: [[TMP2:%.*]] = load i32*, i32** [[A_ADDR]], align 8
// CHECK4-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32* // CHECK4-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32*
// CHECK4-NEXT: [[TMP3:%.*]] = load [10 x i32]*, [10 x i32]** [[D_ADDR]], align 8 // CHECK4-NEXT: [[TMP3:%.*]] = load [10 x i32]*, [10 x i32]** [[D_ADDR]], align 8
// CHECK4-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 false) // CHECK4-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 true)
// CHECK4-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP4]], -1 // CHECK4-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP4]], -1
// CHECK4-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK4-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK4: user_code.entry: // CHECK4: user_code.entry:
@ -57,7 +57,7 @@ int main(int argc, char **argv) {
// CHECK4-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK4-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK4-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK4-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK4-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x i32]* [[TMP0]], [10 x i32]* [[TMP1]], i32* [[TMP2]], i64 [[TMP7]], [10 x i32]* [[TMP3]]) #[[ATTR5:[0-9]+]] // CHECK4-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x i32]* [[TMP0]], [10 x i32]* [[TMP1]], i32* [[TMP2]], i64 [[TMP7]], [10 x i32]* [[TMP3]]) #[[ATTR5:[0-9]+]]
// CHECK4-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK4-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK4-NEXT: ret void // CHECK4-NEXT: ret void
// CHECK4: worker.exit: // CHECK4: worker.exit:
// CHECK4-NEXT: ret void // CHECK4-NEXT: ret void
@ -378,7 +378,7 @@ int main(int argc, char **argv) {
// CHECK5-NEXT: [[TMP1:%.*]] = load [10 x i32]*, [10 x i32]** [[C_ADDR]], align 4 // CHECK5-NEXT: [[TMP1:%.*]] = load [10 x i32]*, [10 x i32]** [[C_ADDR]], align 4
// CHECK5-NEXT: [[TMP2:%.*]] = load i32*, i32** [[A_ADDR]], align 4 // CHECK5-NEXT: [[TMP2:%.*]] = load i32*, i32** [[A_ADDR]], align 4
// CHECK5-NEXT: [[TMP3:%.*]] = load [10 x i32]*, [10 x i32]** [[D_ADDR]], align 4 // CHECK5-NEXT: [[TMP3:%.*]] = load [10 x i32]*, [10 x i32]** [[D_ADDR]], align 4
// CHECK5-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 false) // CHECK5-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 true)
// CHECK5-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP4]], -1 // CHECK5-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP4]], -1
// CHECK5-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK5-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK5: user_code.entry: // CHECK5: user_code.entry:
@ -389,7 +389,7 @@ int main(int argc, char **argv) {
// CHECK5-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK5-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK5-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK5-NEXT: store i32 [[TMP5]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK5-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x i32]* [[TMP0]], [10 x i32]* [[TMP1]], i32* [[TMP2]], i32 [[TMP7]], [10 x i32]* [[TMP3]]) #[[ATTR5:[0-9]+]] // CHECK5-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x i32]* [[TMP0]], [10 x i32]* [[TMP1]], i32* [[TMP2]], i32 [[TMP7]], [10 x i32]* [[TMP3]]) #[[ATTR5:[0-9]+]]
// CHECK5-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK5-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK5-NEXT: ret void // CHECK5-NEXT: ret void
// CHECK5: worker.exit: // CHECK5: worker.exit:
// CHECK5-NEXT: ret void // CHECK5-NEXT: ret void

326
clang/test/OpenMP/nvptx_force_full_runtime_SPMD_codegen.cpp
View File

@ -1,326 +0,0 @@
// 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
void foo() {
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
#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 i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
#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 i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
#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 i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
#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 i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
#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 i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
#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 i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
#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 i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
#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 i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
#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

16
clang/test/OpenMP/nvptx_target_simd_codegen.cpp
View File

@ -61,32 +61,32 @@ int bar(int n){
} }
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l32}}( // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l32}}(
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-NOT: call void @__kmpc_for_static_init // CHECK-NOT: call void @__kmpc_for_static_init
// CHECK-NOT: call void @__kmpc_for_static_fini // CHECK-NOT: call void @__kmpc_for_static_fini
// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 false) // CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 true)
// CHECK: ret void // CHECK: ret void
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l37}}( // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l37}}(
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-NOT: call void @__kmpc_for_static_init // CHECK-NOT: call void @__kmpc_for_static_init
// CHECK-NOT: call void @__kmpc_for_static_fini // CHECK-NOT: call void @__kmpc_for_static_fini
// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 false) // CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 true)
// CHECK: ret void // CHECK: ret void
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l42}}( // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l42}}(
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-NOT: call void @__kmpc_for_static_init // CHECK-NOT: call void @__kmpc_for_static_init
// CHECK-NOT: call void @__kmpc_for_static_fini // CHECK-NOT: call void @__kmpc_for_static_fini
// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 false) // CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 true)
// CHECK: ret void // CHECK: ret void
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l47}}( // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+l47}}(
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK-NOT: call void @__kmpc_for_static_init // CHECK-NOT: call void @__kmpc_for_static_init
// CHECK-NOT: call void @__kmpc_for_static_fini // CHECK-NOT: call void @__kmpc_for_static_fini
// CHECK-NOT: call void @__kmpc_nvptx_end_reduce_nowait( // CHECK-NOT: call void @__kmpc_nvptx_end_reduce_nowait(
// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 false) // CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 true)
// CHECK: ret void // CHECK: ret void
#endif #endif

72
clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_codegen.cpp
View File

@ -87,7 +87,7 @@ int bar(int n){
// CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32* // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32*
// CHECK1-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 8 // CHECK1-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 8
// CHECK1-NEXT: [[CONV1:%.*]] = bitcast i64* [[L_ADDR]] to i32* // CHECK1-NEXT: [[CONV1:%.*]] = bitcast i64* [[L_ADDR]] to i32*
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 false) // CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 true)
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK1: user_code.entry: // CHECK1: user_code.entry:
@ -103,7 +103,7 @@ int bar(int n){
// CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK1-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i32]* [[TMP0]], i64 [[TMP6]]) #[[ATTR3:[0-9]+]] // CHECK1-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i32]* [[TMP0]], i64 [[TMP6]]) #[[ATTR3:[0-9]+]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
// CHECK1: worker.exit: // CHECK1: worker.exit:
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
@ -410,7 +410,7 @@ int bar(int n){
// CHECK1-NEXT: store [1000 x i16]* [[AA]], [1000 x i16]** [[AA_ADDR]], align 8 // CHECK1-NEXT: store [1000 x i16]* [[AA]], [1000 x i16]** [[AA_ADDR]], align 8
// CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32* // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32*
// CHECK1-NEXT: [[TMP0:%.*]] = load [1000 x i16]*, [1000 x i16]** [[AA_ADDR]], align 8 // CHECK1-NEXT: [[TMP0:%.*]] = load [1000 x i16]*, [1000 x i16]** [[AA_ADDR]], align 8
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK1: user_code.entry: // CHECK1: user_code.entry:
@ -422,7 +422,7 @@ int bar(int n){
// CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK1-NEXT: call void @__omp_outlined__2(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i16]* [[TMP0]]) #[[ATTR3]] // CHECK1-NEXT: call void @__omp_outlined__2(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i16]* [[TMP0]]) #[[ATTR3]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
// CHECK1: worker.exit: // CHECK1: worker.exit:
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
@ -665,7 +665,7 @@ int bar(int n){
// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 // CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
// CHECK1-NEXT: store [10 x i32]* [[B]], [10 x i32]** [[B_ADDR]], align 8 // CHECK1-NEXT: store [10 x i32]* [[B]], [10 x i32]** [[B_ADDR]], align 8
// CHECK1-NEXT: [[TMP0:%.*]] = load [10 x i32]*, [10 x i32]** [[B_ADDR]], align 8 // CHECK1-NEXT: [[TMP0:%.*]] = load [10 x i32]*, [10 x i32]** [[B_ADDR]], align 8
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK1: user_code.entry: // CHECK1: user_code.entry:
@ -673,7 +673,7 @@ int bar(int n){
// CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK1-NEXT: call void @__omp_outlined__4(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x i32]* [[TMP0]]) #[[ATTR3]] // CHECK1-NEXT: call void @__omp_outlined__4(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x i32]* [[TMP0]]) #[[ATTR3]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
// CHECK1: worker.exit: // CHECK1: worker.exit:
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
@ -856,7 +856,7 @@ int bar(int n){
// CHECK1-NEXT: store i64 [[F]], i64* [[F_ADDR]], align 8 // CHECK1-NEXT: store i64 [[F]], i64* [[F_ADDR]], align 8
// CHECK1-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 8 // CHECK1-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 8
// CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[F_ADDR]] to i32* // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[F_ADDR]] to i32*
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK1: user_code.entry: // CHECK1: user_code.entry:
@ -868,7 +868,7 @@ int bar(int n){
// CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK1-NEXT: call void @__omp_outlined__6(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x [10 x i32]]* [[TMP0]], i64 [[TMP4]]) #[[ATTR3]] // CHECK1-NEXT: call void @__omp_outlined__6(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x [10 x i32]]* [[TMP0]], i64 [[TMP4]]) #[[ATTR3]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
// CHECK1: worker.exit: // CHECK1: worker.exit:
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
@ -1089,7 +1089,7 @@ int bar(int n){
// CHECK1-NEXT: store [10 x [10 x i32]]* [[C]], [10 x [10 x i32]]** [[C_ADDR]], align 8 // CHECK1-NEXT: store [10 x [10 x i32]]* [[C]], [10 x [10 x i32]]** [[C_ADDR]], align 8
// CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32* // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32*
// CHECK1-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 8 // CHECK1-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 8
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK1: user_code.entry: // CHECK1: user_code.entry:
@ -1101,7 +1101,7 @@ int bar(int n){
// CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK1-NEXT: call void @__omp_outlined__8(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [10 x [10 x i32]]* [[TMP0]]) #[[ATTR3]] // CHECK1-NEXT: call void @__omp_outlined__8(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [10 x [10 x i32]]* [[TMP0]]) #[[ATTR3]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
// CHECK1: worker.exit: // CHECK1: worker.exit:
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
@ -1408,7 +1408,7 @@ int bar(int n){
// CHECK1-NEXT: store i32* [[V]], i32** [[V_ADDR]], align 8 // CHECK1-NEXT: store i32* [[V]], i32** [[V_ADDR]], align 8
// CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32* // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32*
// CHECK1-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 8 // CHECK1-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 8
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK1: user_code.entry: // CHECK1: user_code.entry:
@ -1421,7 +1421,7 @@ int bar(int n){
// CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK1-NEXT: call void @__omp_outlined__10(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i32]* [[TMP0]], i32* [[TMP5]]) #[[ATTR3]] // CHECK1-NEXT: call void @__omp_outlined__10(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i32]* [[TMP0]], i32* [[TMP5]]) #[[ATTR3]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
// CHECK1: worker.exit: // CHECK1: worker.exit:
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
@ -1681,7 +1681,7 @@ int bar(int n){
// CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32* // CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32*
// CHECK2-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 8 // CHECK2-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 8
// CHECK2-NEXT: [[CONV1:%.*]] = bitcast i64* [[L_ADDR]] to i32* // CHECK2-NEXT: [[CONV1:%.*]] = bitcast i64* [[L_ADDR]] to i32*
// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 false) // CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 true)
// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK2: user_code.entry: // CHECK2: user_code.entry:
@ -1697,7 +1697,7 @@ int bar(int n){
// CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK2-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i32]* [[TMP0]], i64 [[TMP6]]) #[[ATTR3:[0-9]+]] // CHECK2-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i32]* [[TMP0]], i64 [[TMP6]]) #[[ATTR3:[0-9]+]]
// CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
// CHECK2: worker.exit: // CHECK2: worker.exit:
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
@ -2004,7 +2004,7 @@ int bar(int n){
// CHECK2-NEXT: store [1000 x i16]* [[AA]], [1000 x i16]** [[AA_ADDR]], align 8 // CHECK2-NEXT: store [1000 x i16]* [[AA]], [1000 x i16]** [[AA_ADDR]], align 8
// CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32* // CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32*
// CHECK2-NEXT: [[TMP0:%.*]] = load [1000 x i16]*, [1000 x i16]** [[AA_ADDR]], align 8 // CHECK2-NEXT: [[TMP0:%.*]] = load [1000 x i16]*, [1000 x i16]** [[AA_ADDR]], align 8
// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK2: user_code.entry: // CHECK2: user_code.entry:
@ -2016,7 +2016,7 @@ int bar(int n){
// CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK2-NEXT: call void @__omp_outlined__2(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i16]* [[TMP0]]) #[[ATTR3]] // CHECK2-NEXT: call void @__omp_outlined__2(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i16]* [[TMP0]]) #[[ATTR3]]
// CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
// CHECK2: worker.exit: // CHECK2: worker.exit:
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
@ -2259,7 +2259,7 @@ int bar(int n){
// CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 // CHECK2-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
// CHECK2-NEXT: store [10 x i32]* [[B]], [10 x i32]** [[B_ADDR]], align 8 // CHECK2-NEXT: store [10 x i32]* [[B]], [10 x i32]** [[B_ADDR]], align 8
// CHECK2-NEXT: [[TMP0:%.*]] = load [10 x i32]*, [10 x i32]** [[B_ADDR]], align 8 // CHECK2-NEXT: [[TMP0:%.*]] = load [10 x i32]*, [10 x i32]** [[B_ADDR]], align 8
// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK2: user_code.entry: // CHECK2: user_code.entry:
@ -2267,7 +2267,7 @@ int bar(int n){
// CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK2-NEXT: call void @__omp_outlined__4(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x i32]* [[TMP0]]) #[[ATTR3]] // CHECK2-NEXT: call void @__omp_outlined__4(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x i32]* [[TMP0]]) #[[ATTR3]]
// CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
// CHECK2: worker.exit: // CHECK2: worker.exit:
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
@ -2450,7 +2450,7 @@ int bar(int n){
// CHECK2-NEXT: store i64 [[F]], i64* [[F_ADDR]], align 8 // CHECK2-NEXT: store i64 [[F]], i64* [[F_ADDR]], align 8
// CHECK2-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 8 // CHECK2-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 8
// CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[F_ADDR]] to i32* // CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[F_ADDR]] to i32*
// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK2: user_code.entry: // CHECK2: user_code.entry:
@ -2462,7 +2462,7 @@ int bar(int n){
// CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK2-NEXT: call void @__omp_outlined__6(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x [10 x i32]]* [[TMP0]], i64 [[TMP4]]) #[[ATTR3]] // CHECK2-NEXT: call void @__omp_outlined__6(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x [10 x i32]]* [[TMP0]], i64 [[TMP4]]) #[[ATTR3]]
// CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
// CHECK2: worker.exit: // CHECK2: worker.exit:
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
@ -2683,7 +2683,7 @@ int bar(int n){
// CHECK2-NEXT: store [10 x [10 x i32]]* [[C]], [10 x [10 x i32]]** [[C_ADDR]], align 8 // CHECK2-NEXT: store [10 x [10 x i32]]* [[C]], [10 x [10 x i32]]** [[C_ADDR]], align 8
// CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32* // CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32*
// CHECK2-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 8 // CHECK2-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 8
// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK2: user_code.entry: // CHECK2: user_code.entry:
@ -2695,7 +2695,7 @@ int bar(int n){
// CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK2-NEXT: call void @__omp_outlined__8(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [10 x [10 x i32]]* [[TMP0]]) #[[ATTR3]] // CHECK2-NEXT: call void @__omp_outlined__8(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [10 x [10 x i32]]* [[TMP0]]) #[[ATTR3]]
// CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
// CHECK2: worker.exit: // CHECK2: worker.exit:
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
@ -2997,7 +2997,7 @@ int bar(int n){
// CHECK2-NEXT: store i32* [[V]], i32** [[V_ADDR]], align 8 // CHECK2-NEXT: store i32* [[V]], i32** [[V_ADDR]], align 8
// CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32* // CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[N_ADDR]] to i32*
// CHECK2-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 8 // CHECK2-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 8
// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK2: user_code.entry: // CHECK2: user_code.entry:
@ -3010,7 +3010,7 @@ int bar(int n){
// CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK2-NEXT: call void @__omp_outlined__10(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i32]* [[TMP0]], i32* [[TMP5]]) #[[ATTR3]] // CHECK2-NEXT: call void @__omp_outlined__10(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], [1000 x i32]* [[TMP0]], i32* [[TMP5]]) #[[ATTR3]]
// CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
// CHECK2: worker.exit: // CHECK2: worker.exit:
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
@ -3268,7 +3268,7 @@ int bar(int n){
// CHECK3-NEXT: store [1000 x i32]* [[A]], [1000 x i32]** [[A_ADDR]], align 4 // CHECK3-NEXT: store [1000 x i32]* [[A]], [1000 x i32]** [[A_ADDR]], align 4
// CHECK3-NEXT: store i32 [[L]], i32* [[L_ADDR]], align 4 // CHECK3-NEXT: store i32 [[L]], i32* [[L_ADDR]], align 4
// CHECK3-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 4 // CHECK3-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 4
// CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 false) // CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 true)
// CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK3: user_code.entry: // CHECK3: user_code.entry:
@ -3282,7 +3282,7 @@ int bar(int n){
// CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK3-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP4]], [1000 x i32]* [[TMP0]], i32 [[TMP6]]) #[[ATTR3:[0-9]+]] // CHECK3-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP4]], [1000 x i32]* [[TMP0]], i32 [[TMP6]]) #[[ATTR3:[0-9]+]]
// CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
// CHECK3: worker.exit: // CHECK3: worker.exit:
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
@ -3575,7 +3575,7 @@ int bar(int n){
// CHECK3-NEXT: store i32 [[N]], i32* [[N_ADDR]], align 4 // CHECK3-NEXT: store i32 [[N]], i32* [[N_ADDR]], align 4
// CHECK3-NEXT: store [1000 x i16]* [[AA]], [1000 x i16]** [[AA_ADDR]], align 4 // CHECK3-NEXT: store [1000 x i16]* [[AA]], [1000 x i16]** [[AA_ADDR]], align 4
// CHECK3-NEXT: [[TMP0:%.*]] = load [1000 x i16]*, [1000 x i16]** [[AA_ADDR]], align 4 // CHECK3-NEXT: [[TMP0:%.*]] = load [1000 x i16]*, [1000 x i16]** [[AA_ADDR]], align 4
// CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK3: user_code.entry: // CHECK3: user_code.entry:
@ -3586,7 +3586,7 @@ int bar(int n){
// CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK3-NEXT: call void @__omp_outlined__2(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP4]], [1000 x i16]* [[TMP0]]) #[[ATTR3]] // CHECK3-NEXT: call void @__omp_outlined__2(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP4]], [1000 x i16]* [[TMP0]]) #[[ATTR3]]
// CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
// CHECK3: worker.exit: // CHECK3: worker.exit:
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
@ -3820,7 +3820,7 @@ int bar(int n){
// CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 // CHECK3-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4
// CHECK3-NEXT: store [10 x i32]* [[B]], [10 x i32]** [[B_ADDR]], align 4 // CHECK3-NEXT: store [10 x i32]* [[B]], [10 x i32]** [[B_ADDR]], align 4
// CHECK3-NEXT: [[TMP0:%.*]] = load [10 x i32]*, [10 x i32]** [[B_ADDR]], align 4 // CHECK3-NEXT: [[TMP0:%.*]] = load [10 x i32]*, [10 x i32]** [[B_ADDR]], align 4
// CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK3: user_code.entry: // CHECK3: user_code.entry:
@ -3828,7 +3828,7 @@ int bar(int n){
// CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK3-NEXT: call void @__omp_outlined__4(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x i32]* [[TMP0]]) #[[ATTR3]] // CHECK3-NEXT: call void @__omp_outlined__4(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x i32]* [[TMP0]]) #[[ATTR3]]
// CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
// CHECK3: worker.exit: // CHECK3: worker.exit:
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
@ -4004,7 +4004,7 @@ int bar(int n){
// CHECK3-NEXT: store [10 x [10 x i32]]* [[C]], [10 x [10 x i32]]** [[C_ADDR]], align 4 // CHECK3-NEXT: store [10 x [10 x i32]]* [[C]], [10 x [10 x i32]]** [[C_ADDR]], align 4
// CHECK3-NEXT: store i32 [[F]], i32* [[F_ADDR]], align 4 // CHECK3-NEXT: store i32 [[F]], i32* [[F_ADDR]], align 4
// CHECK3-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 4 // CHECK3-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 4
// CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK3: user_code.entry: // CHECK3: user_code.entry:
@ -4015,7 +4015,7 @@ int bar(int n){
// CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK3-NEXT: call void @__omp_outlined__6(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x [10 x i32]]* [[TMP0]], i32 [[TMP4]]) #[[ATTR3]] // CHECK3-NEXT: call void @__omp_outlined__6(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], [10 x [10 x i32]]* [[TMP0]], i32 [[TMP4]]) #[[ATTR3]]
// CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
// CHECK3: worker.exit: // CHECK3: worker.exit:
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
@ -4225,7 +4225,7 @@ int bar(int n){
// CHECK3-NEXT: store i32 [[N]], i32* [[N_ADDR]], align 4 // CHECK3-NEXT: store i32 [[N]], i32* [[N_ADDR]], align 4
// CHECK3-NEXT: store [10 x [10 x i32]]* [[C]], [10 x [10 x i32]]** [[C_ADDR]], align 4 // CHECK3-NEXT: store [10 x [10 x i32]]* [[C]], [10 x [10 x i32]]** [[C_ADDR]], align 4
// CHECK3-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 4 // CHECK3-NEXT: [[TMP0:%.*]] = load [10 x [10 x i32]]*, [10 x [10 x i32]]** [[C_ADDR]], align 4
// CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK3: user_code.entry: // CHECK3: user_code.entry:
@ -4236,7 +4236,7 @@ int bar(int n){
// CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK3-NEXT: call void @__omp_outlined__8(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP4]], [10 x [10 x i32]]* [[TMP0]]) #[[ATTR3]] // CHECK3-NEXT: call void @__omp_outlined__8(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP4]], [10 x [10 x i32]]* [[TMP0]]) #[[ATTR3]]
// CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
// CHECK3: worker.exit: // CHECK3: worker.exit:
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
@ -4542,7 +4542,7 @@ int bar(int n){
// CHECK3-NEXT: store [1000 x i32]* [[A]], [1000 x i32]** [[A_ADDR]], align 4 // CHECK3-NEXT: store [1000 x i32]* [[A]], [1000 x i32]** [[A_ADDR]], align 4
// CHECK3-NEXT: store i32* [[V]], i32** [[V_ADDR]], align 4 // CHECK3-NEXT: store i32* [[V]], i32** [[V_ADDR]], align 4
// CHECK3-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 4 // CHECK3-NEXT: [[TMP0:%.*]] = load [1000 x i32]*, [1000 x i32]** [[A_ADDR]], align 4
// CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 false) // CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i8 2, i1 false, i1 true)
// CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK3: user_code.entry: // CHECK3: user_code.entry:
@ -4554,7 +4554,7 @@ int bar(int n){
// CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK3-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK3-NEXT: call void @__omp_outlined__10(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP4]], [1000 x i32]* [[TMP0]], i32* [[TMP5]]) #[[ATTR3]] // CHECK3-NEXT: call void @__omp_outlined__10(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP4]], [1000 x i32]* [[TMP0]], i32* [[TMP5]]) #[[ATTR3]]
// CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void
// CHECK3: worker.exit: // CHECK3: worker.exit:
// CHECK3-NEXT: ret void // CHECK3-NEXT: ret void

8
clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_generic_mode_codegen.cpp
View File

@ -45,7 +45,7 @@ int main(int argc, char **argv) {
// CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32* // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[ARGC_ADDR]] to i32*
// CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[A_ADDR]], align 8 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[A_ADDR]], align 8
// CHECK1-NEXT: [[CONV1:%.*]] = bitcast i64* [[DOTCAPTURE_EXPR__ADDR]] to i32* // CHECK1-NEXT: [[CONV1:%.*]] = bitcast i64* [[DOTCAPTURE_EXPR__ADDR]] to i32*
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 false) // CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 true)
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK1: user_code.entry: // CHECK1: user_code.entry:
@ -61,7 +61,7 @@ int main(int argc, char **argv) {
// CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK1-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK1-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], i32* [[TMP0]], i64 [[TMP6]]) #[[ATTR3:[0-9]+]] // CHECK1-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i64 [[TMP4]], i32* [[TMP0]], i64 [[TMP6]]) #[[ATTR3:[0-9]+]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
// CHECK1: worker.exit: // CHECK1: worker.exit:
// CHECK1-NEXT: ret void // CHECK1-NEXT: ret void
@ -353,7 +353,7 @@ int main(int argc, char **argv) {
// CHECK2-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 4 // CHECK2-NEXT: store i32* [[A]], i32** [[A_ADDR]], align 4
// CHECK2-NEXT: store i32 [[DOTCAPTURE_EXPR_]], i32* [[DOTCAPTURE_EXPR__ADDR]], align 4 // CHECK2-NEXT: store i32 [[DOTCAPTURE_EXPR_]], i32* [[DOTCAPTURE_EXPR__ADDR]], align 4
// CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[A_ADDR]], align 4 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[A_ADDR]], align 4
// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 false) // CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 true)
// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK2: user_code.entry: // CHECK2: user_code.entry:
@ -367,7 +367,7 @@ int main(int argc, char **argv) {
// CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4 // CHECK2-NEXT: store i32 [[TMP2]], i32* [[DOTTHREADID_TEMP_]], align 4
// CHECK2-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP4]], i32* [[TMP0]], i32 [[TMP6]]) #[[ATTR3:[0-9]+]] // CHECK2-NEXT: call void @__omp_outlined__(i32* [[DOTTHREADID_TEMP_]], i32* [[DOTZERO_ADDR]], i32 [[TMP4]], i32* [[TMP0]], i32 [[TMP6]]) #[[ATTR3:[0-9]+]]
// CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 false) // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i8 2, i1 true)
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void
// CHECK2: worker.exit: // CHECK2: worker.exit:
// CHECK2-NEXT: ret void // CHECK2-NEXT: ret void

664
clang/test/OpenMP/nvptx_target_teams_distribute_parallel_for_simd_codegen.cpp
View File

File diff suppressed because it is too large Load Diff

16
clang/test/OpenMP/nvptx_target_teams_distribute_simd_codegen.cpp
View File

@ -70,24 +70,24 @@ int bar(int n){
} }
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+}}_l37( // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+}}_l37(
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 false) // CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 true)
// CHECK: call void @__kmpc_distribute_static_init_4({{.+}}, {{.+}}, {{.+}} 91, // CHECK: call void @__kmpc_distribute_static_init_4({{.+}}, {{.+}}, {{.+}} 91,
// CHECK: call void @__kmpc_distribute_static_fini( // CHECK: call void @__kmpc_distribute_static_fini(
// CHECK: ret void // CHECK: ret void
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+}}_l43( // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+}}_l43(
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 false) // CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 true)
// CHECK: call void @__kmpc_distribute_static_init_4({{.+}}, {{.+}}, {{.+}} 91, // CHECK: call void @__kmpc_distribute_static_init_4({{.+}}, {{.+}}, {{.+}} 91,
// CHECK: call void @__kmpc_distribute_static_fini( // CHECK: call void @__kmpc_distribute_static_fini(
// CHECK: ret void // CHECK: ret void
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+}}_l48( // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+}}_l48(
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 false) // CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 true)
// CHECK: call void @__kmpc_distribute_static_init_4({{.+}}, {{.+}}, {{.+}} 91, // CHECK: call void @__kmpc_distribute_static_init_4({{.+}}, {{.+}}, {{.+}} 91,
// CHECK: call void @__kmpc_distribute_static_fini( // CHECK: call void @__kmpc_distribute_static_fini(
@ -95,8 +95,8 @@ int bar(int n){
// CHECK: define {{.*}}void {{@__omp_offloading_.+}}_l53({{.+}}, i{{32|64}} [[F_IN:%.+]]) // CHECK: define {{.*}}void {{@__omp_offloading_.+}}_l53({{.+}}, i{{32|64}} [[F_IN:%.+]])
// CHECK: store {{.+}} [[F_IN]], {{.+}}* {{.+}}, // CHECK: store {{.+}} [[F_IN]], {{.+}}* {{.+}},
// CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 false) // CHECK: call i32 @__kmpc_target_init({{.*}}, i8 2, i1 false, i1 true)
// CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 false) // CHECK: call void @__kmpc_target_deinit({{.*}}, i8 2, i1 true)
// CHECK: store {{.+}} 99, {{.+}}* [[COMB_UB:%.+]], align // CHECK: store {{.+}} 99, {{.+}}* [[COMB_UB:%.+]], align
// CHECK: call void @__kmpc_distribute_static_init_4({{.+}}, {{.+}}, {{.+}} 91, {{.+}}, {{.+}}, {{.+}}* [[COMB_UB]], // CHECK: call void @__kmpc_distribute_static_init_4({{.+}}, {{.+}}, {{.+}} 91, {{.+}}, {{.+}}, {{.+}}* [[COMB_UB]],

12
clang/test/OpenMP/target_parallel_for_debug_codegen.cpp
View File

@ -89,7 +89,7 @@ int main() {
// CHECK1-NEXT: [[TMP6:%.*]] = addrspacecast i8 addrspace(1)* [[TMP5]] to i8*, !dbg [[DBG41]] // CHECK1-NEXT: [[TMP6:%.*]] = addrspacecast i8 addrspace(1)* [[TMP5]] to i8*, !dbg [[DBG41]]
// CHECK1-NEXT: store i8* [[TMP6]], i8** [[_TMP2]], align 8, !dbg [[DBG41]] // CHECK1-NEXT: store i8* [[TMP6]], i8** [[_TMP2]], align 8, !dbg [[DBG41]]
// CHECK1-NEXT: [[TMP7:%.*]] = load i8*, i8** [[_TMP2]], align 8, !dbg [[DBG41]] // CHECK1-NEXT: [[TMP7:%.*]] = load i8*, i8** [[_TMP2]], align 8, !dbg [[DBG41]]
// CHECK1-NEXT: [[TMP8:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 false), !dbg [[DBG41]] // CHECK1-NEXT: [[TMP8:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i8 2, i1 false, i1 true), !dbg [[DBG41]]
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP8]], -1, !dbg [[DBG41]] // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP8]], -1, !dbg [[DBG41]]
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]], !dbg [[DBG41]] // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]], !dbg [[DBG41]]
// CHECK1: user_code.entry: // CHECK1: user_code.entry:
@ -114,7 +114,7 @@ int main() {
// CHECK1-NEXT: [[TMP20:%.*]] = zext i1 [[TOBOOL]] to i32, !dbg [[DBG42]] // CHECK1-NEXT: [[TMP20:%.*]] = zext i1 [[TOBOOL]] to i32, !dbg [[DBG42]]
// CHECK1-NEXT: [[TMP21:%.*]] = bitcast [4 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**, !dbg [[DBG42]] // CHECK1-NEXT: [[TMP21:%.*]] = bitcast [4 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**, !dbg [[DBG42]]
// CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB6]], i32 [[TMP9]], i32 [[TMP20]], i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, [10 x [10 x [10 x i32]]]*, i64, [10 x [10 x i32]]*, i8*)* @__omp_outlined__ to i8*), i8* null, i8** [[TMP21]], i64 4), !dbg [[DBG42]] // CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB6]], i32 [[TMP9]], i32 [[TMP20]], i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, [10 x [10 x [10 x i32]]]*, i64, [10 x [10 x i32]]*, i8*)* @__omp_outlined__ to i8*), i8* null, i8** [[TMP21]], i64 4), !dbg [[DBG42]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB8:[0-9]+]], i8 2, i1 false), !dbg [[DBG45:![0-9]+]] // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB8:[0-9]+]], i8 2, i1 true), !dbg [[DBG45:![0-9]+]]
// CHECK1-NEXT: ret void, !dbg [[DBG46:![0-9]+]] // CHECK1-NEXT: ret void, !dbg [[DBG46:![0-9]+]]
// CHECK1: worker.exit: // CHECK1: worker.exit:
// CHECK1-NEXT: ret void, !dbg [[DBG41]] // CHECK1-NEXT: ret void, !dbg [[DBG41]]
@ -392,7 +392,7 @@ int main() {
// CHECK1-NEXT: [[TMP7:%.*]] = addrspacecast i8 addrspace(1)* [[TMP6]] to i8*, !dbg [[DBG146]] // CHECK1-NEXT: [[TMP7:%.*]] = addrspacecast i8 addrspace(1)* [[TMP6]] to i8*, !dbg [[DBG146]]
// CHECK1-NEXT: store i8* [[TMP7]], i8** [[_TMP2]], align 8, !dbg [[DBG146]] // CHECK1-NEXT: store i8* [[TMP7]], i8** [[_TMP2]], align 8, !dbg [[DBG146]]
// CHECK1-NEXT: [[TMP8:%.*]] = load i8*, i8** [[_TMP2]], align 8, !dbg [[DBG146]] // CHECK1-NEXT: [[TMP8:%.*]] = load i8*, i8** [[_TMP2]], align 8, !dbg [[DBG146]]
// CHECK1-NEXT: [[TMP9:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB10:[0-9]+]], i8 2, i1 false, i1 false), !dbg [[DBG146]] // CHECK1-NEXT: [[TMP9:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB10:[0-9]+]], i8 2, i1 false, i1 true), !dbg [[DBG146]]
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP9]], -1, !dbg [[DBG146]] // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP9]], -1, !dbg [[DBG146]]
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]], !dbg [[DBG146]] // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]], !dbg [[DBG146]]
// CHECK1: user_code.entry: // CHECK1: user_code.entry:
@ -414,7 +414,7 @@ int main() {
// CHECK1-NEXT: store i8* [[TMP8]], i8** [[TMP19]], align 8, !dbg [[DBG147]] // CHECK1-NEXT: store i8* [[TMP8]], i8** [[TMP19]], align 8, !dbg [[DBG147]]
// CHECK1-NEXT: [[TMP20:%.*]] = bitcast [4 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**, !dbg [[DBG147]] // CHECK1-NEXT: [[TMP20:%.*]] = bitcast [4 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**, !dbg [[DBG147]]
// CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB15]], i32 [[TMP10]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, [10 x [10 x [10 x i32]]]*, i64, [10 x [10 x i32]]*, i8*)* @__omp_outlined__2 to i8*), i8* null, i8** [[TMP20]], i64 4), !dbg [[DBG147]] // CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB15]], i32 [[TMP10]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, [10 x [10 x [10 x i32]]]*, i64, [10 x [10 x i32]]*, i8*)* @__omp_outlined__2 to i8*), i8* null, i8** [[TMP20]], i64 4), !dbg [[DBG147]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB17:[0-9]+]], i8 2, i1 false), !dbg [[DBG148:![0-9]+]] // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB17:[0-9]+]], i8 2, i1 true), !dbg [[DBG148:![0-9]+]]
// CHECK1-NEXT: ret void, !dbg [[DBG150:![0-9]+]] // CHECK1-NEXT: ret void, !dbg [[DBG150:![0-9]+]]
// CHECK1: worker.exit: // CHECK1: worker.exit:
// CHECK1-NEXT: ret void, !dbg [[DBG146]] // CHECK1-NEXT: ret void, !dbg [[DBG146]]
@ -680,7 +680,7 @@ int main() {
// CHECK1-NEXT: [[TMP10:%.*]] = addrspacecast i8 addrspace(1)* [[TMP9]] to i8*, !dbg [[DBG236]] // CHECK1-NEXT: [[TMP10:%.*]] = addrspacecast i8 addrspace(1)* [[TMP9]] to i8*, !dbg [[DBG236]]
// CHECK1-NEXT: store i8* [[TMP10]], i8** [[_TMP3]], align 8, !dbg [[DBG236]] // CHECK1-NEXT: store i8* [[TMP10]], i8** [[_TMP3]], align 8, !dbg [[DBG236]]
// CHECK1-NEXT: [[TMP11:%.*]] = load i8*, i8** [[_TMP3]], align 8, !dbg [[DBG236]] // CHECK1-NEXT: [[TMP11:%.*]] = load i8*, i8** [[_TMP3]], align 8, !dbg [[DBG236]]
// CHECK1-NEXT: [[TMP12:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB19:[0-9]+]], i8 2, i1 false, i1 false), !dbg [[DBG236]] // CHECK1-NEXT: [[TMP12:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB19:[0-9]+]], i8 2, i1 false, i1 true), !dbg [[DBG236]]
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP12]], -1, !dbg [[DBG236]] // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP12]], -1, !dbg [[DBG236]]
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]], !dbg [[DBG236]] // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]], !dbg [[DBG236]]
// CHECK1: user_code.entry: // CHECK1: user_code.entry:
@ -698,7 +698,7 @@ int main() {
// CHECK1-NEXT: store i8* [[TMP11]], i8** [[TMP20]], align 8, !dbg [[DBG237]] // CHECK1-NEXT: store i8* [[TMP11]], i8** [[TMP20]], align 8, !dbg [[DBG237]]
// CHECK1-NEXT: [[TMP21:%.*]] = bitcast [4 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**, !dbg [[DBG237]] // CHECK1-NEXT: [[TMP21:%.*]] = bitcast [4 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**, !dbg [[DBG237]]
// CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB24]], i32 [[TMP13]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, [10 x [10 x [10 x i32]]]*, i32*, [10 x [10 x i32]]*, i8*)* @__omp_outlined__4 to i8*), i8* null, i8** [[TMP21]], i64 4), !dbg [[DBG237]] // CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB24]], i32 [[TMP13]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*, [10 x [10 x [10 x i32]]]*, i32*, [10 x [10 x i32]]*, i8*)* @__omp_outlined__4 to i8*), i8* null, i8** [[TMP21]], i64 4), !dbg [[DBG237]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB26:[0-9]+]], i8 2, i1 false), !dbg [[DBG238:![0-9]+]] // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB26:[0-9]+]], i8 2, i1 true), !dbg [[DBG238:![0-9]+]]
// CHECK1-NEXT: ret void, !dbg [[DBG240:![0-9]+]] // CHECK1-NEXT: ret void, !dbg [[DBG240:![0-9]+]]
// CHECK1: worker.exit: // CHECK1: worker.exit:
// CHECK1-NEXT: ret void, !dbg [[DBG236]] // CHECK1-NEXT: ret void, !dbg [[DBG236]]