mirror of
https://github.com/intel/llvm.git
synced 2026-02-05 13:21:04 +08:00
[OPENMP] Codegen for #pragma omp target parallel for.
llvm-svn: 317719
This commit is contained in:
Alexey Bataev
@@ -406,6 +406,9 @@ public:
|
||||
/// \brief Skip no-op (attributed, compound) container stmts and skip captured
|
||||
/// stmt at the top, if \a IgnoreCaptured is true.
|
||||
Stmt *IgnoreContainers(bool IgnoreCaptured = false);
|
||||
const Stmt *IgnoreContainers(bool IgnoreCaptured = false) const {
|
||||
return const_cast<Stmt *>(this)->IgnoreContainers(IgnoreCaptured);
|
||||
}
|
||||
|
||||
const Stmt *stripLabelLikeStatements() const;
|
||||
Stmt *stripLabelLikeStatements() {
|
||||
|
||||
@@ -899,7 +899,9 @@ public:
|
||||
}
|
||||
const Stmt *getBody() const {
|
||||
// This relies on the loop form is already checked by Sema.
|
||||
Stmt *Body = getAssociatedStmt()->IgnoreContainers(true);
|
||||
const Stmt *Body = getAssociatedStmt()->IgnoreContainers(true);
|
||||
while(const auto *CS = dyn_cast<CapturedStmt>(Body))
|
||||
Body = CS->getCapturedStmt();
|
||||
Body = cast<ForStmt>(Body)->getBody();
|
||||
for (unsigned Cnt = 1; Cnt < CollapsedNum; ++Cnt) {
|
||||
Body = Body->IgnoreContainers();
|
||||
|
||||
@@ -908,7 +908,6 @@ void clang::getOpenMPCaptureRegions(
|
||||
case OMPD_atomic:
|
||||
case OMPD_target_data:
|
||||
case OMPD_target:
|
||||
case OMPD_target_parallel_for:
|
||||
case OMPD_target_parallel_for_simd:
|
||||
case OMPD_target_simd:
|
||||
case OMPD_task:
|
||||
@@ -926,6 +925,7 @@ void clang::getOpenMPCaptureRegions(
|
||||
CaptureRegions.push_back(DKind);
|
||||
break;
|
||||
case OMPD_target_parallel:
|
||||
case OMPD_target_parallel_for:
|
||||
CaptureRegions.push_back(OMPD_target);
|
||||
CaptureRegions.push_back(OMPD_parallel);
|
||||
break;
|
||||
|
||||
@@ -7121,6 +7121,10 @@ void CGOpenMPRuntime::scanForTargetRegionsFunctions(const Stmt *S,
|
||||
CodeGenFunction::EmitOMPTargetTeamsDeviceFunction(
|
||||
CGM, ParentName, cast<OMPTargetTeamsDirective>(*S));
|
||||
break;
|
||||
case Stmt::OMPTargetParallelForDirectiveClass:
|
||||
CodeGenFunction::EmitOMPTargetParallelForDeviceFunction(
|
||||
CGM, ParentName, cast<OMPTargetParallelForDirective>(*S));
|
||||
break;
|
||||
default:
|
||||
llvm_unreachable("Unknown target directive for OpenMP device codegen.");
|
||||
}
|
||||
|
||||
@@ -276,6 +276,7 @@ getExecutionModeForDirective(CodeGenModule &CGM,
|
||||
case OMPD_target_teams:
|
||||
return CGOpenMPRuntimeNVPTX::ExecutionMode::Generic;
|
||||
case OMPD_target_parallel:
|
||||
case OMPD_target_parallel_for:
|
||||
return CGOpenMPRuntimeNVPTX::ExecutionMode::Spmd;
|
||||
default:
|
||||
llvm_unreachable("Unsupported directive on NVPTX device.");
|
||||
|
||||
@@ -3683,7 +3683,7 @@ static void emitCommonOMPTargetDirective(CodeGenFunction &CGF,
|
||||
const RegionCodeGenTy &CodeGen) {
|
||||
assert(isOpenMPTargetExecutionDirective(S.getDirectiveKind()));
|
||||
CodeGenModule &CGM = CGF.CGM;
|
||||
const CapturedStmt &CS = *cast<CapturedStmt>(S.getAssociatedStmt());
|
||||
const CapturedStmt &CS = *S.getCapturedStmt(OMPD_target);
|
||||
|
||||
llvm::Function *Fn = nullptr;
|
||||
llvm::Constant *FnID = nullptr;
|
||||
@@ -4133,16 +4133,41 @@ void CodeGenFunction::EmitOMPTargetParallelDirective(
|
||||
emitCommonOMPTargetDirective(*this, S, CodeGen);
|
||||
}
|
||||
|
||||
static void emitTargetParallelForRegion(CodeGenFunction &CGF,
|
||||
const OMPTargetParallelForDirective &S,
|
||||
PrePostActionTy &Action) {
|
||||
Action.Enter(CGF);
|
||||
// Emit directive as a combined directive that consists of two implicit
|
||||
// directives: 'parallel' with 'for' directive.
|
||||
auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &) {
|
||||
CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), emitForLoopBounds,
|
||||
emitDispatchForLoopBounds);
|
||||
};
|
||||
emitCommonOMPParallelDirective(CGF, S, OMPD_for, CodeGen,
|
||||
emitEmptyBoundParameters);
|
||||
}
|
||||
|
||||
void CodeGenFunction::EmitOMPTargetParallelForDeviceFunction(
|
||||
CodeGenModule &CGM, StringRef ParentName,
|
||||
const OMPTargetParallelForDirective &S) {
|
||||
// Emit SPMD target parallel for region as a standalone region.
|
||||
auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
|
||||
emitTargetParallelForRegion(CGF, S, Action);
|
||||
};
|
||||
llvm::Function *Fn;
|
||||
llvm::Constant *Addr;
|
||||
// Emit target region as a standalone region.
|
||||
CGM.getOpenMPRuntime().emitTargetOutlinedFunction(
|
||||
S, ParentName, Fn, Addr, /*IsOffloadEntry=*/true, CodeGen);
|
||||
assert(Fn && Addr && "Target device function emission failed.");
|
||||
}
|
||||
|
||||
void CodeGenFunction::EmitOMPTargetParallelForDirective(
|
||||
const OMPTargetParallelForDirective &S) {
|
||||
OMPLexicalScope Scope(*this, S, /*AsInlined=*/true);
|
||||
CGM.getOpenMPRuntime().emitInlinedDirective(
|
||||
*this, OMPD_target_parallel_for,
|
||||
[&S](CodeGenFunction &CGF, PrePostActionTy &) {
|
||||
OMPLoopScope PreInitScope(CGF, S);
|
||||
CGF.EmitStmt(
|
||||
cast<CapturedStmt>(S.getAssociatedStmt())->getCapturedStmt());
|
||||
});
|
||||
auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
|
||||
emitTargetParallelForRegion(CGF, S, Action);
|
||||
};
|
||||
emitCommonOMPTargetDirective(*this, S, CodeGen);
|
||||
}
|
||||
|
||||
/// Emit a helper variable and return corresponding lvalue.
|
||||
|
||||
@@ -2889,6 +2889,10 @@ public:
|
||||
static void
|
||||
EmitOMPTargetParallelDeviceFunction(CodeGenModule &CGM, StringRef ParentName,
|
||||
const OMPTargetParallelDirective &S);
|
||||
/// Emit device code for the target parallel for directive.
|
||||
static void EmitOMPTargetParallelForDeviceFunction(
|
||||
CodeGenModule &CGM, StringRef ParentName,
|
||||
const OMPTargetParallelForDirective &S);
|
||||
static void
|
||||
EmitOMPTargetTeamsDeviceFunction(CodeGenModule &CGM, StringRef ParentName,
|
||||
const OMPTargetTeamsDirective &S);
|
||||
|
||||
@@ -2084,7 +2084,8 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
break;
|
||||
}
|
||||
case OMPD_target_teams:
|
||||
case OMPD_target_parallel: {
|
||||
case OMPD_target_parallel:
|
||||
case OMPD_target_parallel_for: {
|
||||
Sema::CapturedParamNameType ParamsTarget[] = {
|
||||
std::make_pair(StringRef(), QualType()) // __context with shared vars
|
||||
};
|
||||
@@ -2119,7 +2120,6 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
case OMPD_atomic:
|
||||
case OMPD_target_data:
|
||||
case OMPD_target:
|
||||
case OMPD_target_parallel_for:
|
||||
case OMPD_target_parallel_for_simd:
|
||||
case OMPD_target_simd: {
|
||||
Sema::CapturedParamNameType Params[] = {
|
||||
@@ -6309,13 +6309,23 @@ StmtResult Sema::ActOnOpenMPTargetParallelForDirective(
|
||||
// The point of exit cannot be a branch out of the structured block.
|
||||
// longjmp() and throw() must not violate the entry/exit criteria.
|
||||
CS->getCapturedDecl()->setNothrow();
|
||||
for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_target_parallel_for);
|
||||
ThisCaptureLevel > 1; --ThisCaptureLevel) {
|
||||
CS = cast<CapturedStmt>(CS->getCapturedStmt());
|
||||
// 1.2.2 OpenMP Language Terminology
|
||||
// Structured block - An executable statement with a single entry at the
|
||||
// top and a single exit at the bottom.
|
||||
// The point of exit cannot be a branch out of the structured block.
|
||||
// longjmp() and throw() must not violate the entry/exit criteria.
|
||||
CS->getCapturedDecl()->setNothrow();
|
||||
}
|
||||
|
||||
OMPLoopDirective::HelperExprs B;
|
||||
// In presence of clause 'collapse' or 'ordered' with number of loops, it will
|
||||
// define the nested loops number.
|
||||
unsigned NestedLoopCount =
|
||||
CheckOpenMPLoop(OMPD_target_parallel_for, getCollapseNumberExpr(Clauses),
|
||||
getOrderedNumberExpr(Clauses), AStmt, *this, *DSAStack,
|
||||
getOrderedNumberExpr(Clauses), CS, *this, *DSAStack,
|
||||
VarsWithImplicitDSA, B);
|
||||
if (NestedLoopCount == 0)
|
||||
return StmtError();
|
||||
@@ -7285,6 +7295,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
|
||||
case OMPC_if:
|
||||
switch (DKind) {
|
||||
case OMPD_target_parallel:
|
||||
case OMPD_target_parallel_for:
|
||||
// If this clause applies to the nested 'parallel' region, capture within
|
||||
// the 'target' region, otherwise do not capture.
|
||||
if (NameModifier == OMPD_unknown || NameModifier == OMPD_parallel)
|
||||
@@ -7297,7 +7308,6 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
|
||||
case OMPD_parallel_for_simd:
|
||||
case OMPD_target:
|
||||
case OMPD_target_simd:
|
||||
case OMPD_target_parallel_for:
|
||||
case OMPD_target_parallel_for_simd:
|
||||
case OMPD_target_teams:
|
||||
case OMPD_target_teams_distribute:
|
||||
@@ -7351,6 +7361,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
|
||||
case OMPC_num_threads:
|
||||
switch (DKind) {
|
||||
case OMPD_target_parallel:
|
||||
case OMPD_target_parallel_for:
|
||||
CaptureRegion = OMPD_target;
|
||||
break;
|
||||
case OMPD_cancel:
|
||||
@@ -7360,7 +7371,6 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
|
||||
case OMPD_parallel_for_simd:
|
||||
case OMPD_target:
|
||||
case OMPD_target_simd:
|
||||
case OMPD_target_parallel_for:
|
||||
case OMPD_target_parallel_for_simd:
|
||||
case OMPD_target_teams:
|
||||
case OMPD_target_teams_distribute:
|
||||
@@ -7538,6 +7548,68 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
|
||||
}
|
||||
break;
|
||||
case OMPC_schedule:
|
||||
switch (DKind) {
|
||||
case OMPD_target_parallel_for:
|
||||
CaptureRegion = OMPD_target;
|
||||
break;
|
||||
case OMPD_parallel_for:
|
||||
case OMPD_parallel_for_simd:
|
||||
case OMPD_target_parallel_for_simd:
|
||||
case OMPD_target_teams_distribute_parallel_for:
|
||||
case OMPD_target_teams_distribute_parallel_for_simd:
|
||||
case OMPD_teams_distribute_parallel_for:
|
||||
case OMPD_teams_distribute_parallel_for_simd:
|
||||
case OMPD_distribute_parallel_for:
|
||||
case OMPD_distribute_parallel_for_simd:
|
||||
// Do not capture thread_limit-clause expressions.
|
||||
break;
|
||||
case OMPD_task:
|
||||
case OMPD_taskloop:
|
||||
case OMPD_taskloop_simd:
|
||||
case OMPD_target_data:
|
||||
case OMPD_target_enter_data:
|
||||
case OMPD_target_exit_data:
|
||||
case OMPD_target_update:
|
||||
case OMPD_teams:
|
||||
case OMPD_teams_distribute:
|
||||
case OMPD_teams_distribute_simd:
|
||||
case OMPD_target_teams_distribute:
|
||||
case OMPD_target_teams_distribute_simd:
|
||||
case OMPD_target:
|
||||
case OMPD_target_simd:
|
||||
case OMPD_target_parallel:
|
||||
case OMPD_cancel:
|
||||
case OMPD_parallel:
|
||||
case OMPD_parallel_sections:
|
||||
case OMPD_threadprivate:
|
||||
case OMPD_taskyield:
|
||||
case OMPD_barrier:
|
||||
case OMPD_taskwait:
|
||||
case OMPD_cancellation_point:
|
||||
case OMPD_flush:
|
||||
case OMPD_declare_reduction:
|
||||
case OMPD_declare_simd:
|
||||
case OMPD_declare_target:
|
||||
case OMPD_end_declare_target:
|
||||
case OMPD_simd:
|
||||
case OMPD_for:
|
||||
case OMPD_for_simd:
|
||||
case OMPD_sections:
|
||||
case OMPD_section:
|
||||
case OMPD_single:
|
||||
case OMPD_master:
|
||||
case OMPD_critical:
|
||||
case OMPD_taskgroup:
|
||||
case OMPD_distribute:
|
||||
case OMPD_ordered:
|
||||
case OMPD_atomic:
|
||||
case OMPD_distribute_simd:
|
||||
case OMPD_target_teams:
|
||||
llvm_unreachable("Unexpected OpenMP directive with schedule clause");
|
||||
case OMPD_unknown:
|
||||
llvm_unreachable("Unknown OpenMP directive");
|
||||
}
|
||||
break;
|
||||
case OMPC_dist_schedule:
|
||||
case OMPC_firstprivate:
|
||||
case OMPC_lastprivate:
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
451
clang/test/OpenMP/target_parallel_for_codegen_registration.cpp
Normal file
451
clang/test/OpenMP/target_parallel_for_codegen_registration.cpp
Normal file
@@ -0,0 +1,451 @@
|
||||
// Test host codegen.
|
||||
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
|
||||
|
||||
// Test target parallel for codegen - host bc file has to be created first.
|
||||
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
|
||||
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefix=TCHECK
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefix=TCHECK
|
||||
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
|
||||
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s -check-prefix=TCHECK
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefix=TCHECK
|
||||
|
||||
// Check that no target code is emmitted if no omptests flag was provided.
|
||||
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s -check-prefix=CHECK-NTARGET
|
||||
|
||||
// expected-no-diagnostics
|
||||
#ifndef HEADER
|
||||
#define HEADER
|
||||
|
||||
// CHECK-DAG: [[SA:%.+]] = type { [4 x i32] }
|
||||
// CHECK-DAG: [[SB:%.+]] = type { [8 x i32] }
|
||||
// CHECK-DAG: [[SC:%.+]] = type { [16 x i32] }
|
||||
// CHECK-DAG: [[SD:%.+]] = type { [32 x i32] }
|
||||
// CHECK-DAG: [[SE:%.+]] = type { [64 x i32] }
|
||||
// CHECK-DAG: [[ST1:%.+]] = type { [228 x i32] }
|
||||
// CHECK-DAG: [[ST2:%.+]] = type { [1128 x i32] }
|
||||
// CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
|
||||
// CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
|
||||
// CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* }
|
||||
|
||||
// TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
|
||||
|
||||
// CHECK-DAG: $[[REGFN:\.omp_offloading\..+]] = comdat
|
||||
|
||||
// CHECK-DAG: [[A1:@.+]] = internal global [[SA]]
|
||||
// CHECK-DAG: [[A2:@.+]] = global [[SA]]
|
||||
// CHECK-DAG: [[B1:@.+]] = global [[SB]]
|
||||
// CHECK-DAG: [[B2:@.+]] = global [[SB]]
|
||||
// CHECK-DAG: [[C1:@.+]] = internal global [[SC]]
|
||||
// CHECK-DAG: [[D1:@.+]] = global [[SD]]
|
||||
// CHECK-DAG: [[E1:@.+]] = global [[SE]]
|
||||
// CHECK-DAG: [[T1:@.+]] = global [[ST1]]
|
||||
// CHECK-DAG: [[T2:@.+]] = global [[ST2]]
|
||||
|
||||
// CHECK-NTARGET-DAG: [[SA:%.+]] = type { [4 x i32] }
|
||||
// CHECK-NTARGET-DAG: [[SB:%.+]] = type { [8 x i32] }
|
||||
// CHECK-NTARGET-DAG: [[SC:%.+]] = type { [16 x i32] }
|
||||
// CHECK-NTARGET-DAG: [[SD:%.+]] = type { [32 x i32] }
|
||||
// CHECK-NTARGET-DAG: [[SE:%.+]] = type { [64 x i32] }
|
||||
// CHECK-NTARGET-DAG: [[ST1:%.+]] = type { [228 x i32] }
|
||||
// CHECK-NTARGET-DAG: [[ST2:%.+]] = type { [1128 x i32] }
|
||||
// CHECK-NTARGET-NOT: type { i8*, i8*, %
|
||||
// CHECK-NTARGET-NOT: type { i32, %
|
||||
|
||||
// We have 7 target regions
|
||||
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// TCHECK-NOT: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
// CHECK-DAG: {{@.+}} = private constant i8 0
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
|
||||
// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
|
||||
|
||||
// CHECK-NTARGET-NOT: private constant i8 0
|
||||
// CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i
|
||||
|
||||
// CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
|
||||
// CHECK-DAG: [[ENTRY1:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR1]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR2:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME2:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY2:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR2]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR3:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME3:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY3:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR3]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR4:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME4:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY4:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR4]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR5:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME5:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY5:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR5]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR6:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME6:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY6:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR6]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR7:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME7:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY7:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR7]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR8:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME8:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY8:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR8]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR9:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME9:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY9:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR9]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR10:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME10:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY10:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR10]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR11:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME11:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY11:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR11]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// CHECK-DAG: [[NAMEPTR12:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME12:.+]]\00"
|
||||
// CHECK-DAG: [[ENTRY12:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR12]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
|
||||
// TCHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY1:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR1]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR2:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME2:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY2:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR2]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR3:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME3:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY3:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR3]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR4:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME4:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY4:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR4]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR5:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME5:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY5:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR5]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR6:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME6:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY6:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR6]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR7:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME7:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY7:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR7]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR8:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME8:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY8:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR8]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR9:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME9:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY9:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR9]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR10:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME10:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY10:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR10]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR11:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME11:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY11:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR11]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// TCHECK-DAG: [[NAMEPTR12:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME12:.+]]\00"
|
||||
// TCHECK-DAG: [[ENTRY12:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR12]], i32 0, i32 0), i[[SZ]] 0, i32 0, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
|
||||
// CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]]
|
||||
// CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]]
|
||||
// CHECK: [[DEVBEGIN:@.+]] = external constant i8
|
||||
// CHECK: [[DEVEND:@.+]] = external constant i8
|
||||
// CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]])
|
||||
// CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 1, [[DEVTY]]* getelementptr inbounds ([1 x [[DEVTY]]], [1 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, comdat($[[REGFN]])
|
||||
|
||||
// We have 4 initializers, one for the 500 priority, another one for 501, or more for the default priority, and the last one for the offloading registration function.
|
||||
// CHECK: @llvm.global_ctors = appending global [4 x { i32, void ()*, i8* }] [
|
||||
// CHECK-SAME: { i32, void ()*, i8* } { i32 500, void ()* [[P500:@[^,]+]], i8* null },
|
||||
// CHECK-SAME: { i32, void ()*, i8* } { i32 501, void ()* [[P501:@[^,]+]], i8* null },
|
||||
// CHECK-SAME: { i32, void ()*, i8* } { i32 65535, void ()* [[PMAX:@[^,]+]], i8* null },
|
||||
// CHECK-SAME: { i32, void ()*, i8* } { i32 0, void ()* bitcast (void (i8*)* @[[REGFN]] to void ()*), i8* bitcast (void (i8*)* @[[REGFN]] to i8*) }]
|
||||
|
||||
// CHECK-NTARGET: @llvm.global_ctors = appending global [3 x { i32, void ()*, i8* }] [
|
||||
|
||||
extern int *R;
|
||||
|
||||
struct SA {
|
||||
int arr[4];
|
||||
void foo() {
|
||||
int a = *R;
|
||||
a += 1;
|
||||
*R = a;
|
||||
}
|
||||
SA() {
|
||||
int a = *R;
|
||||
a += 2;
|
||||
*R = a;
|
||||
}
|
||||
~SA() {
|
||||
int a = *R;
|
||||
a += 3;
|
||||
*R = a;
|
||||
}
|
||||
};
|
||||
|
||||
struct SB {
|
||||
int arr[8];
|
||||
void foo() {
|
||||
int a = *R;
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
a += 4;
|
||||
*R = a;
|
||||
}
|
||||
SB() {
|
||||
int a = *R;
|
||||
a += 5;
|
||||
*R = a;
|
||||
}
|
||||
~SB() {
|
||||
int a = *R;
|
||||
a += 6;
|
||||
*R = a;
|
||||
}
|
||||
};
|
||||
|
||||
struct SC {
|
||||
int arr[16];
|
||||
void foo() {
|
||||
int a = *R;
|
||||
a += 7;
|
||||
*R = a;
|
||||
}
|
||||
SC() {
|
||||
int a = *R;
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
a += 8;
|
||||
*R = a;
|
||||
}
|
||||
~SC() {
|
||||
int a = *R;
|
||||
a += 9;
|
||||
*R = a;
|
||||
}
|
||||
};
|
||||
|
||||
struct SD {
|
||||
int arr[32];
|
||||
void foo() {
|
||||
int a = *R;
|
||||
a += 10;
|
||||
*R = a;
|
||||
}
|
||||
SD() {
|
||||
int a = *R;
|
||||
a += 11;
|
||||
*R = a;
|
||||
}
|
||||
~SD() {
|
||||
int a = *R;
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
a += 12;
|
||||
*R = a;
|
||||
}
|
||||
};
|
||||
|
||||
struct SE {
|
||||
int arr[64];
|
||||
void foo() {
|
||||
int a = *R;
|
||||
#pragma omp target parallel for if(target: 0)
|
||||
for (int i = 0; i < 10; ++i)
|
||||
a += 13;
|
||||
*R = a;
|
||||
}
|
||||
SE() {
|
||||
int a = *R;
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
a += 14;
|
||||
*R = a;
|
||||
}
|
||||
~SE() {
|
||||
int a = *R;
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
a += 15;
|
||||
*R = a;
|
||||
}
|
||||
};
|
||||
|
||||
template <int x>
|
||||
struct ST {
|
||||
int arr[128 + x];
|
||||
void foo() {
|
||||
int a = *R;
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
a += 16 + x;
|
||||
*R = a;
|
||||
}
|
||||
ST() {
|
||||
int a = *R;
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
a += 17 + x;
|
||||
*R = a;
|
||||
}
|
||||
~ST() {
|
||||
int a = *R;
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
a += 18 + x;
|
||||
*R = a;
|
||||
}
|
||||
};
|
||||
|
||||
// We have to make sure we us all the target regions:
|
||||
//CHECK-DAG: define internal void @[[NAME1]](
|
||||
//CHECK-DAG: call void @[[NAME1]](
|
||||
//CHECK-DAG: define internal void @[[NAME2]](
|
||||
//CHECK-DAG: call void @[[NAME2]](
|
||||
//CHECK-DAG: define internal void @[[NAME3]](
|
||||
//CHECK-DAG: call void @[[NAME3]](
|
||||
//CHECK-DAG: define internal void @[[NAME4]](
|
||||
//CHECK-DAG: call void @[[NAME4]](
|
||||
//CHECK-DAG: define internal void @[[NAME5]](
|
||||
//CHECK-DAG: call void @[[NAME5]](
|
||||
//CHECK-DAG: define internal void @[[NAME6]](
|
||||
//CHECK-DAG: call void @[[NAME6]](
|
||||
//CHECK-DAG: define internal void @[[NAME7]](
|
||||
//CHECK-DAG: call void @[[NAME7]](
|
||||
//CHECK-DAG: define internal void @[[NAME8]](
|
||||
//CHECK-DAG: call void @[[NAME8]](
|
||||
//CHECK-DAG: define internal void @[[NAME9]](
|
||||
//CHECK-DAG: call void @[[NAME9]](
|
||||
//CHECK-DAG: define internal void @[[NAME10]](
|
||||
//CHECK-DAG: call void @[[NAME10]](
|
||||
//CHECK-DAG: define internal void @[[NAME11]](
|
||||
//CHECK-DAG: call void @[[NAME11]](
|
||||
//CHECK-DAG: define internal void @[[NAME12]](
|
||||
//CHECK-DAG: call void @[[NAME12]](
|
||||
|
||||
//TCHECK-DAG: define void @[[NAME1]](
|
||||
//TCHECK-DAG: define void @[[NAME2]](
|
||||
//TCHECK-DAG: define void @[[NAME3]](
|
||||
//TCHECK-DAG: define void @[[NAME4]](
|
||||
//TCHECK-DAG: define void @[[NAME5]](
|
||||
//TCHECK-DAG: define void @[[NAME6]](
|
||||
//TCHECK-DAG: define void @[[NAME7]](
|
||||
//TCHECK-DAG: define void @[[NAME8]](
|
||||
//TCHECK-DAG: define void @[[NAME9]](
|
||||
//TCHECK-DAG: define void @[[NAME10]](
|
||||
//TCHECK-DAG: define void @[[NAME11]](
|
||||
//TCHECK-DAG: define void @[[NAME12]](
|
||||
|
||||
// CHECK-NTARGET-NOT: __tgt_target
|
||||
// CHECK-NTARGET-NOT: __tgt_register_lib
|
||||
// CHECK-NTARGET-NOT: __tgt_unregister_lib
|
||||
|
||||
// TCHECK-NOT: __tgt_target
|
||||
// TCHECK-NOT: __tgt_register_lib
|
||||
// TCHECK-NOT: __tgt_unregister_lib
|
||||
|
||||
// We have 2 initializers with priority 500
|
||||
//CHECK: define internal void [[P500]](
|
||||
//CHECK: call void @{{.+}}()
|
||||
//CHECK: call void @{{.+}}()
|
||||
//CHECK-NOT: call void @{{.+}}()
|
||||
//CHECK: ret void
|
||||
|
||||
// We have 1 initializers with priority 501
|
||||
//CHECK: define internal void [[P501]](
|
||||
//CHECK: call void @{{.+}}()
|
||||
//CHECK-NOT: call void @{{.+}}()
|
||||
//CHECK: ret void
|
||||
|
||||
// We have 6 initializers with default priority
|
||||
//CHECK: define internal void [[PMAX]](
|
||||
//CHECK: call void @{{.+}}()
|
||||
//CHECK: call void @{{.+}}()
|
||||
//CHECK: call void @{{.+}}()
|
||||
//CHECK: call void @{{.+}}()
|
||||
//CHECK: call void @{{.+}}()
|
||||
//CHECK: call void @{{.+}}()
|
||||
//CHECK-NOT: call void @{{.+}}()
|
||||
//CHECK: ret void
|
||||
|
||||
// Check registration and unregistration
|
||||
|
||||
//CHECK: define internal void @[[UNREGFN:.+]](i8*)
|
||||
//CHECK-SAME: comdat($[[REGFN]]) {
|
||||
//CHECK: call i32 @__tgt_unregister_lib([[DSCTY]]* [[DESC]])
|
||||
//CHECK: ret void
|
||||
//CHECK: declare i32 @__tgt_unregister_lib([[DSCTY]]*)
|
||||
|
||||
//CHECK: define linkonce hidden void @[[REGFN]](i8*)
|
||||
//CHECK-SAME: comdat {
|
||||
//CHECK: call i32 @__tgt_register_lib([[DSCTY]]* [[DESC]])
|
||||
//CHECK: call i32 @__cxa_atexit(void (i8*)* @[[UNREGFN]], i8* bitcast ([[DSCTY]]* [[DESC]] to i8*),
|
||||
//CHECK: ret void
|
||||
//CHECK: declare i32 @__tgt_register_lib([[DSCTY]]*)
|
||||
|
||||
static __attribute__((init_priority(500))) SA a1;
|
||||
SA a2;
|
||||
SB __attribute__((init_priority(500))) b1;
|
||||
SB __attribute__((init_priority(501))) b2;
|
||||
static SC c1;
|
||||
SD d1;
|
||||
SE e1;
|
||||
ST<100> t1;
|
||||
ST<1000> t2;
|
||||
|
||||
|
||||
int bar(int a){
|
||||
int r = a;
|
||||
|
||||
a1.foo();
|
||||
a2.foo();
|
||||
b1.foo();
|
||||
b2.foo();
|
||||
c1.foo();
|
||||
d1.foo();
|
||||
e1.foo();
|
||||
t1.foo();
|
||||
t2.foo();
|
||||
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
++r;
|
||||
|
||||
return r + *R;
|
||||
}
|
||||
|
||||
// Check metadata is properly generated:
|
||||
// CHECK: !omp_offload.info = !{!{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID:-?[0-9]+]], i32 [[FILEID:-?[0-9]+]], !"_ZN2SB3fooEv", i32 195, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SDD1Ev", i32 247, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SEC1Ev", i32 265, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SED1Ev", i32 272, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EE3fooEv", i32 284, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EEC1Ev", i32 291, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_Z3bari", i32 415, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EED1Ev", i32 298, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EEC1Ev", i32 291, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EED1Ev", i32 298, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EE3fooEv", i32 284, i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SCC1Ev", i32 221, i32 {{[0-9]+}}}
|
||||
|
||||
// TCHECK: !omp_offload.info = !{!{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID:-?[0-9]+]], i32 [[FILEID:-?[0-9]+]], !"_ZN2SB3fooEv", i32 195, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SDD1Ev", i32 247, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SEC1Ev", i32 265, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SED1Ev", i32 272, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EE3fooEv", i32 284, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EEC1Ev", i32 291, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_Z3bari", i32 415, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EED1Ev", i32 298, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EEC1Ev", i32 291, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EED1Ev", i32 298, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EE3fooEv", i32 284, i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SCC1Ev", i32 221, i32 {{[0-9]+}}}
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,68 @@
|
||||
// Test host codegen.
|
||||
// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
|
||||
|
||||
// Test target parallel for codegen - host bc file has to be created first.
|
||||
// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
|
||||
// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefix=TCHECK
|
||||
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefix=TCHECK
|
||||
// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
|
||||
// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s -check-prefix=TCHECK
|
||||
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefix=TCHECK
|
||||
|
||||
// expected-no-diagnostics
|
||||
#ifndef HEADER
|
||||
#define HEADER
|
||||
|
||||
// CHECK: [[CA:%.+]] = type { i32* }
|
||||
|
||||
// CHECK: define {{.*}}i32 @[[NNAME:.+]](i32 {{.*}}%{{.+}})
|
||||
int nested(int a){
|
||||
// CHECK: call void @__omp_offloading_[[FILEID:[0-9a-f]+_[0-9a-f]+]]_[[NNAME]]_l[[T1L:[0-9]+]](
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
++a;
|
||||
|
||||
// CHECK: call void @"[[LNAME:.+]]"([[CA]]*
|
||||
auto F = [&](){
|
||||
#pragma omp parallel
|
||||
{
|
||||
#pragma omp target parallel for
|
||||
for (int i = 0; i < 10; ++i)
|
||||
++a;
|
||||
}
|
||||
};
|
||||
|
||||
F();
|
||||
|
||||
return a;
|
||||
}
|
||||
|
||||
// CHECK: define {{.*}}void @__omp_offloading_[[FILEID]]_[[NNAME]]_l[[T1L]](
|
||||
// TCHECK: define {{.*}}void @__omp_offloading_[[FILEID:[0-9a-f]+_[0-9a-f]+]]_[[NNAME:.+]]_l[[T1L:[0-9]+]](
|
||||
|
||||
// CHECK: define {{.*}}void @"[[LNAME]]"(
|
||||
// CHECK: call void {{.*}}@__kmpc_fork_call{{.+}}[[PNAME:@.+]] to
|
||||
|
||||
// CHECK: define {{.*}}void [[PNAME]](
|
||||
// CHECK: call void @__omp_offloading_[[FILEID]]_[[NNAME]]_l[[T2L:[0-9]+]](
|
||||
|
||||
// CHECK: define {{.*}}void @__omp_offloading_[[FILEID]]_[[NNAME]]_l[[T2L]](
|
||||
// TCHECK: define {{.*}}void @__omp_offloading_[[FILEID]]_[[NNAME:.+]]_l[[T2L:[0-9]+]](
|
||||
|
||||
|
||||
// Check metadata is properly generated:
|
||||
// CHECK: !omp_offload.info = !{!{{[0-9]+}}, !{{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 {{-?[0-9]+}}, i32 {{-?[0-9]+}}, !"[[NNAME]]", i32 [[T1L]], i32 {{[0-9]+}}}
|
||||
// CHECK-DAG: = !{i32 0, i32 {{-?[0-9]+}}, i32 {{-?[0-9]+}}, !"[[NNAME]]", i32 [[T2L]], i32 {{[0-9]+}}}
|
||||
|
||||
// TCHECK: !omp_offload.info = !{!{{[0-9]+}}, !{{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 {{-?[0-9]+}}, i32 {{-?[0-9]+}}, !"[[NNAME]]", i32 [[T1L]], i32 {{[0-9]+}}}
|
||||
// TCHECK-DAG: = !{i32 0, i32 {{-?[0-9]+}}, i32 {{-?[0-9]+}}, !"[[NNAME]]", i32 [[T2L]], i32 {{[0-9]+}}}
|
||||
#endif
|
||||
113
clang/test/OpenMP/target_parallel_for_debug_codegen.cpp
Normal file
113
clang/test/OpenMP/target_parallel_for_debug_codegen.cpp
Normal file
@@ -0,0 +1,113 @@
|
||||
// RUN: %clang_cc1 -DCK1 -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 -DCK1 -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 - -debug-info-kind=limited | FileCheck %s
|
||||
// expected-no-diagnostics
|
||||
|
||||
int main() {
|
||||
/* int(*b)[a]; */
|
||||
/* int *(**c)[a]; */
|
||||
bool bb;
|
||||
int a;
|
||||
int b[10][10];
|
||||
int c[10][10][10];
|
||||
#pragma omp target parallel for firstprivate(a, b) map(tofrom \
|
||||
: c) map(tofrom \
|
||||
: bb) if (a)
|
||||
for (int i = 0; i < 10; ++i) {
|
||||
int &f = c[1][1][1];
|
||||
int &g = a;
|
||||
int &h = b[1][1];
|
||||
int d = 15;
|
||||
a = 5;
|
||||
b[0][a] = 10;
|
||||
c[0][0][a] = 11;
|
||||
b[0][a] = c[0][0][a];
|
||||
bb |= b[0][a];
|
||||
}
|
||||
#pragma omp target parallel for firstprivate(a) map(tofrom \
|
||||
: c, b) map(to \
|
||||
: bb)
|
||||
for (int i = 0; i < 10; ++i) {
|
||||
int &f = c[1][1][1];
|
||||
int &g = a;
|
||||
int &h = b[1][1];
|
||||
int d = 15;
|
||||
a = 5;
|
||||
b[0][a] = 10;
|
||||
c[0][0][a] = 11;
|
||||
b[0][a] = c[0][0][a];
|
||||
d = bb;
|
||||
}
|
||||
#pragma omp target parallel for map(tofrom \
|
||||
: a, c, b) map(from \
|
||||
: bb)
|
||||
for (int i = 0; i < 10; ++i) {
|
||||
int &f = c[1][1][1];
|
||||
int &g = a;
|
||||
int &h = b[1][1];
|
||||
int d = 15;
|
||||
a = 5;
|
||||
b[0][a] = 10;
|
||||
c[0][0][a] = 11;
|
||||
b[0][a] = c[0][0][a];
|
||||
bb = b[0][a];
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
|
||||
// CHECK: define internal void @__omp_offloading{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* noalias{{[^,]+}}, i1 {{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
|
||||
// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8* {{[^)]+}})
|
||||
|
||||
// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
|
||||
|
||||
// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
|
||||
// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
|
||||
|
||||
// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
|
||||
// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
|
||||
|
||||
// CHECK: define internal void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
|
||||
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
|
||||
// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8* {{[^)]+}})
|
||||
|
||||
// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
|
||||
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
|
||||
|
||||
// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
|
||||
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
|
||||
// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
|
||||
|
||||
// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
|
||||
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
|
||||
// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
|
||||
|
||||
// CHECK: define internal void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 addrspace(1)* noalias{{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
|
||||
// CHECK: addrspacecast i32 addrspace(1)* %{{.+}} to i32*
|
||||
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
|
||||
// CHECK: call void [[NONDEBUG_WRAPPER:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8* {{[^)]+}})
|
||||
|
||||
// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 addrspace(1)* noalias{{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
|
||||
// CHECK: addrspacecast i32 addrspace(1)* %{{.+}} to i32*
|
||||
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
|
||||
|
||||
// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
|
||||
// CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
|
||||
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
|
||||
// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
|
||||
|
||||
// CHECK: define void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
|
||||
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
|
||||
// CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
|
||||
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
|
||||
// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
|
||||
|
||||
Reference in New Issue
Block a user