This is an archive of the discontinued LLVM Phabricator instance.

Differential D81478

[OPENMP50]Codegen for scan directives in parallel for regions.
ClosedPublic

Authored by ABataev on Jun 9 2020, 9:35 AM.

Details

Reviewers
jdoerfert
Commits
rG1ec469cf4c20: [OPENMP50]Codegen for scan directives in parallel for regions.
Summary

Added codegen for scan directives in parallel for regions.

Emits the code for the directive with inscan reductions.
Original code:

#pragma omp parallel for reduction(inscan, op : ...)
for() {
  <input phase>;
  #pragma omp scan (in)exclusive(...)
  <scan phase>
}

is transformed to something:

 #pragma omp parallel
{
size num_iters = <num_iters>;
<type> buffer[num_iters];
 #pragma omp for
for (i: 0..<num_iters>) {
  <input phase>;
  buffer[i] = red;
}
 #pragma omp barrier
for (int k = 0; k != ceil(log2(num_iters)); ++k)
for (size cnt = last_iter; cnt >= pow(2, k); --k)
  buffer[i] op= buffer[i-pow(2,k)];
 #pragma omp for
for (0..<num_iters>) {
  red = InclusiveScan ? buffer[i] : buffer[i-1];
  <scan phase>;
}
}

Diff Detail

Event Timeline

ABataev created this revision.Jun 9 2020, 9:35 AM
Herald added a project: Restricted Project. · View Herald TranscriptJun 9 2020, 9:35 AM
Herald added subscribers: sstefan1, guansong, yaxunl. · View Herald Transcript
Harbormaster failed remote builds in B59650: Diff 269571!Jun 9 2020, 10:58 AM
ABataev updated this revision to Diff 271115.Jun 16 2020, 9:05 AM

Rebase

Harbormaster failed remote builds in B60494: Diff 271115!Jun 16 2020, 10:27 AM
jdoerfert added inline comments.Jun 16 2020, 11:34 AM
clang/lib/CodeGen/CGStmtOpenMP.cpp
3634

This looks pretty much like D81658, right? Could we avoid the duplication and maybe use a templated function (or something else)?

ABataev marked an inline comment as done.Jun 16 2020, 11:38 AM
ABataev added inline comments.
clang/lib/CodeGen/CGStmtOpenMP.cpp
3634

The duplication is quite small. Here we don't need to check for lastprivates update, we need to check for the cancellation region. I don't think that the outlined functions are going to be much simpler and easier to read.

jdoerfert added inline comments.Jun 16 2020, 5:40 PM
clang/lib/CodeGen/CGStmtOpenMP.cpp
3634

I mean, these ~25 lines will exist 3 times at least. 2 times tacking lastprivate, which we can do for the third time at no extra cost, and 2 times with the cancelation RAII. Any change will need to be duplicated 3 times as well, etc.

If we do

template<typename RAII>
void emitWorksharingHelper(..., bool &HasLastprivate)
    if (llvm::any_of(S.getClausesOfKind<OMPReductionClause>(),
                     [](const OMPReductionClause *C) {
                       return C->getModifier() == OMPC_REDUCTION_inscan;
                     })) {
      const auto &&NumIteratorsGen = [&S](CodeGenFunction &CGF) {
        OMPLocalDeclMapRAII Scope(CGF);
        OMPLoopScope LoopScope(CGF, S);
        return CGF.EmitScalarExpr(S.getNumIterations());
      };
      const auto &&FirstGen = [&S](CodeGenFunction &CGF) {
        RAII CancelRegion(CGF, OMPD_for, S.hasCancel());
        (void)CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
                                         emitForLoopBounds,
                                         emitDispatchForLoopBounds);
        // Emit an implicit barrier at the end.
        CGF.CGM.getOpenMPRuntime().emitBarrierCall(CGF, S.getBeginLoc(),
                                                   OMPD_for);
      };
      const auto &&SecondGen = [&S, &HasLastprivate](CodeGenFunction &CGF) {
        RAII CancelRegion(CGF, OMPD_for, S.hasCancel());
        HasLastprivate = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
                                         emitForLoopBounds,
                                         emitDispatchForLoopBounds);
      };
      emitScanBasedDirective(CGF, S, NumIteratorsGen, FirstGen, SecondGen);
    } else {
      OMPCancelStackRAII CancelRegion(CGF, OMPD_parallel_for, S.hasCancel());
        CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), emitForLoopBounds,
                                   emitDispatchForLoopBounds);
    }

We can call it three times:

emitWorksharingHelper<OMPCancelStackRAII>(...);
emitWorksharingHelper<OMPCancelStackRAII>(...);
emitWorksharingHelper<OMPDummyRAII>(...);

Wouldn't that be cleaner?

ABataev marked an inline comment as done.Jun 17 2020, 4:36 AM
ABataev added inline comments.
clang/lib/CodeGen/CGStmtOpenMP.cpp
3634
  1. It requires the new OMPDummyRAII class.
  2. Thу functionality under control of the else branch also must be affected.
  3. We don't always need to check for HasLastprivate.
  4. The NumIteratorsGen might be different.
  5. We'll have a lot of trouble with adding new functionality to this function, especially if this functionality is unique for one construct and not required for others.
jdoerfert added inline comments.Jun 17 2020, 9:32 AM
clang/lib/CodeGen/CGStmtOpenMP.cpp
3634
  1. It requires the new OMPDummyRAII class.

Sure, is it more than this?

struct OMPDummyRAII {
  OMPDummyRAII(...) {}
};
  1. Thу functionality under control of the else branch also must be affected.

Right: sed -e/OMPCancelStackRAII/RAII/ :)

  1. We don't always need to check for HasLastprivate.

There is no cost in doing so though.

  1. The NumIteratorsGen might be different.

Might be or is? Eyeballing this it looks like 3 times exact the same code.

  1. We'll have a lot of trouble with adding new functionality to this function, especially if this functionality is unique for one construct and not required for others.

What new functionality do you expect to be unique between them? If there is stuff planned, we might not want to unify the implementation, if it is just a hunch, we want to go ahead and have a single implemenation so it is easier to read and do global modifications. Finding out that something needs special handling is at the end of the day easier than finding all other (unrelated) locations that need to be changed.

ABataev marked an inline comment as done.Jun 17 2020, 10:03 AM
ABataev added inline comments.
clang/lib/CodeGen/CGStmtOpenMP.cpp
3634

What new functionality do you expect to be unique between them? If there is stuff planned, we might not want to unify the implementation, if it is just a hunch, we want to go ahead and have a single implemenation so it is easier to read and do global modifications. Finding out that something needs special handling is at the end of the day easier than finding all other (unrelated) locations that need to be changed.

I have a different opinion here. When you will some specific implementation for one construct that is not required for other, you can easily miss it, and it might lead to unexpected results, like a compiler crash or something else. Adding step by step is much better, because you not only modify the code for the one particular construct, but also definitely will modify the tests for this construct. Having a common solution, you may easily miss something.

ABataev updated this revision to Diff 271437.Jun 17 2020, 12:03 PM

Rebase + avoid copy paste.

Harbormaster failed remote builds in B60679: Diff 271437!Jun 17 2020, 1:30 PM
jdoerfert accepted this revision.Jun 18 2020, 7:09 AM

LGTM and cleaner than my sketch, thanks a lot :)

This revision is now accepted and ready to land.Jun 18 2020, 7:09 AM
Closed by commit rG1ec469cf4c20: [OPENMP50]Codegen for scan directives in parallel for regions. (authored by ABataev). · Explain WhyJun 18 2020, 9:13 AM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
clang/
lib/
CodeGen/
107 lines
test/
OpenMP/
313 lines

Diff 271742

clang/lib/CodeGen/CGStmtOpenMP.cpp

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 3,197 Lines • ▼ Show 20 Linesstatic void emitScanBasedDirective(
CGF.Builder.CreateCondBr(Cmp, LoopBB, ExitBB); CGF.Builder.CreateCondBr(Cmp, LoopBB, ExitBB);
auto DL1 = ApplyDebugLocation::CreateDefaultArtificial(CGF, S.getEndLoc()); auto DL1 = ApplyDebugLocation::CreateDefaultArtificial(CGF, S.getEndLoc());
CGF.EmitBlock(ExitBB); CGF.EmitBlock(ExitBB);
CGF.OMPFirstScanLoop = false; CGF.OMPFirstScanLoop = false;
SecondGen(CGF); SecondGen(CGF);
} }
void CodeGenFunction::EmitOMPForDirective(const OMPForDirective &S) { static bool emitWorksharingDirective(CodeGenFunction &CGF,
bool HasLastprivates = false; const OMPLoopDirective &S,
auto &&CodeGen = [&S, &HasLastprivates](CodeGenFunction &CGF, bool HasCancel) {
PrePostActionTy &) { bool HasLastprivates;
if (llvm::any_of(S.getClausesOfKind<OMPReductionClause>(), if (llvm::any_of(S.getClausesOfKind<OMPReductionClause>(),
[](const OMPReductionClause *C) { [](const OMPReductionClause *C) {
return C->getModifier() == OMPC_REDUCTION_inscan; return C->getModifier() == OMPC_REDUCTION_inscan;
})) { })) {
const auto &&NumIteratorsGen = [&S](CodeGenFunction &CGF) { const auto &&NumIteratorsGen = [&S](CodeGenFunction &CGF) {
OMPLocalDeclMapRAII Scope(CGF); CodeGenFunction::OMPLocalDeclMapRAII Scope(CGF);
OMPLoopScope LoopScope(CGF, S); OMPLoopScope LoopScope(CGF, S);
return CGF.EmitScalarExpr(S.getNumIterations()); return CGF.EmitScalarExpr(S.getNumIterations());
}; };
const auto &&FirstGen = [&S](CodeGenFunction &CGF) { const auto &&FirstGen = [&S, HasCancel](CodeGenFunction &CGF) {
OMPCancelStackRAII CancelRegion(CGF, OMPD_for, S.hasCancel()); CodeGenFunction::OMPCancelStackRAII CancelRegion(
CGF, S.getDirectiveKind(), HasCancel);
(void)CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), (void)CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
emitForLoopBounds, emitForLoopBounds,
emitDispatchForLoopBounds); emitDispatchForLoopBounds);
// Emit an implicit barrier at the end. // Emit an implicit barrier at the end.
CGF.CGM.getOpenMPRuntime().emitBarrierCall(CGF, S.getBeginLoc(), CGF.CGM.getOpenMPRuntime().emitBarrierCall(CGF, S.getBeginLoc(),
OMPD_for); OMPD_for);
}; };
const auto &&SecondGen = [&S, &HasLastprivates](CodeGenFunction &CGF) { const auto &&SecondGen = [&S, HasCancel,
OMPCancelStackRAII CancelRegion(CGF, OMPD_for, S.hasCancel()); &HasLastprivates](CodeGenFunction &CGF) {
CodeGenFunction::OMPCancelStackRAII CancelRegion(
CGF, S.getDirectiveKind(), HasCancel);
HasLastprivates = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), HasLastprivates = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
emitForLoopBounds, emitForLoopBounds,
emitDispatchForLoopBounds); emitDispatchForLoopBounds);
}; };
emitScanBasedDirective(CGF, S, NumIteratorsGen, FirstGen, SecondGen); emitScanBasedDirective(CGF, S, NumIteratorsGen, FirstGen, SecondGen);
} else { } else {
OMPCancelStackRAII CancelRegion(CGF, OMPD_for, S.hasCancel()); CodeGenFunction::OMPCancelStackRAII CancelRegion(CGF, S.getDirectiveKind(),
HasCancel);
HasLastprivates = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), HasLastprivates = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
emitForLoopBounds, emitForLoopBounds,
emitDispatchForLoopBounds); emitDispatchForLoopBounds);
} }
return HasLastprivates;
}
void CodeGenFunction::EmitOMPForDirective(const OMPForDirective &S) {
bool HasLastprivates = false;
auto &&CodeGen = [&S, &HasLastprivates](CodeGenFunction &CGF,
PrePostActionTy &) {
HasLastprivates = emitWorksharingDirective(CGF, S, S.hasCancel());
}; };
{ {
auto LPCRegion = auto LPCRegion =
CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S); CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
OMPLexicalScope Scope(*this, S, OMPD_unknown); OMPLexicalScope Scope(*this, S, OMPD_unknown);
CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_for, CodeGen, CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_for, CodeGen,
S.hasCancel()); S.hasCancel());
} }
// Emit an implicit barrier at the end. // Emit an implicit barrier at the end.
if (!S.getSingleClause<OMPNowaitClause>() || HasLastprivates) if (!S.getSingleClause<OMPNowaitClause>() || HasLastprivates)
CGM.getOpenMPRuntime().emitBarrierCall(*this, S.getBeginLoc(), OMPD_for); CGM.getOpenMPRuntime().emitBarrierCall(*this, S.getBeginLoc(), OMPD_for);
// Check for outer lastprivate conditional update. // Check for outer lastprivate conditional update.
checkForLastprivateConditionalUpdate(*this, S); checkForLastprivateConditionalUpdate(*this, S);
} }
void CodeGenFunction::EmitOMPForSimdDirective(const OMPForSimdDirective &S) { void CodeGenFunction::EmitOMPForSimdDirective(const OMPForSimdDirective &S) {
bool HasLastprivates = false; bool HasLastprivates = false;
auto &&CodeGen = [&S, &HasLastprivates](CodeGenFunction &CGF, auto &&CodeGen = [&S, &HasLastprivates](CodeGenFunction &CGF,
PrePostActionTy &) { PrePostActionTy &) {
if (llvm::any_of(S.getClausesOfKind<OMPReductionClause>(), HasLastprivates = emitWorksharingDirective(CGF, S, /*HasCancel=*/false);
[](const OMPReductionClause *C) {
return C->getModifier() == OMPC_REDUCTION_inscan;
})) {
const auto &&NumIteratorsGen = [&S](CodeGenFunction &CGF) {
OMPLocalDeclMapRAII Scope(CGF);
OMPLoopScope LoopScope(CGF, S);
return CGF.EmitScalarExpr(S.getNumIterations());
};
const auto &&FirstGen = [&S](CodeGenFunction &CGF) {
(void)CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
emitForLoopBounds,
emitDispatchForLoopBounds);
// Emit an implicit barrier at the end.
CGF.CGM.getOpenMPRuntime().emitBarrierCall(CGF, S.getBeginLoc(),
OMPD_for);
};
const auto &&SecondGen = [&S, &HasLastprivates](CodeGenFunction &CGF) {
HasLastprivates = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
emitForLoopBounds,
emitDispatchForLoopBounds);
};
emitScanBasedDirective(CGF, S, NumIteratorsGen, FirstGen, SecondGen);
} else {
HasLastprivates = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
emitForLoopBounds,
emitDispatchForLoopBounds);
}
}; };
{ {
auto LPCRegion = auto LPCRegion =
CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S); CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
OMPLexicalScope Scope(*this, S, OMPD_unknown); OMPLexicalScope Scope(*this, S, OMPD_unknown);
CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_simd, CodeGen); CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_simd, CodeGen);
} }
▲ Show 20 Lines • Show All 319 Lines • ▼ Show 20 Lines
} }
void CodeGenFunction::EmitOMPParallelForDirective( void CodeGenFunction::EmitOMPParallelForDirective(
const OMPParallelForDirective &S) { const OMPParallelForDirective &S) {
// Emit directive as a combined directive that consists of two implicit // Emit directive as a combined directive that consists of two implicit
// directives: 'parallel' with 'for' directive. // directives: 'parallel' with 'for' directive.
auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) { auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
Action.Enter(CGF); Action.Enter(CGF);
OMPCancelStackRAII CancelRegion(CGF, OMPD_parallel_for, S.hasCancel()); (void)emitWorksharingDirective(CGF, S, S.hasCancel());
CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), emitForLoopBounds,
emitDispatchForLoopBounds);
}; };
{ {
auto LPCRegion = auto LPCRegion =
CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S); CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
emitCommonOMPParallelDirective(*this, S, OMPD_for, CodeGen, emitCommonOMPParallelDirective(*this, S, OMPD_for, CodeGen,
emitEmptyBoundParameters); emitEmptyBoundParameters);
} }
// Check for outer lastprivate conditional update. // Check for outer lastprivate conditional update.
checkForLastprivateConditionalUpdate(*this, S); checkForLastprivateConditionalUpdate(*this, S);
} }
void CodeGenFunction::EmitOMPParallelForSimdDirective( void CodeGenFunction::EmitOMPParallelForSimdDirective(
const OMPParallelForSimdDirective &S) { const OMPParallelForSimdDirective &S) {
// Emit directive as a combined directive that consists of two implicit // Emit directive as a combined directive that consists of two implicit
// directives: 'parallel' with 'for' directive. // directives: 'parallel' with 'for' directive.
auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) { auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
Action.Enter(CGF); Action.Enter(CGF);
CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), emitForLoopBounds, CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), emitForLoopBounds,
emitDispatchForLoopBounds); emitDispatchForLoopBounds);
}; };
{ {
auto LPCRegion = auto LPCRegion =
CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S); CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
emitCommonOMPParallelDirective(*this, S, OMPD_simd, CodeGen, emitCommonOMPParallelDirective(*this, S, OMPD_simd, CodeGen,
emitEmptyBoundParameters); emitEmptyBoundParameters);
jdoerfertUnsubmitted
Not Done
ReplyInline Actions

This looks pretty much like D81658, right? Could we avoid the duplication and maybe use a templated function (or something else)?

jdoerfert: This looks pretty much like D81658, right? Could we avoid the duplication and maybe use a…
ABataevAuthorUnsubmitted
Done
ReplyInline Actions

The duplication is quite small. Here we don't need to check for lastprivates update, we need to check for the cancellation region. I don't think that the outlined functions are going to be much simpler and easier to read.

ABataev: The duplication is quite small. Here we don't need to check for lastprivates update, we need to…
jdoerfertUnsubmitted
Not Done
ReplyInline Actions

I mean, these ~25 lines will exist 3 times at least. 2 times tacking lastprivate, which we can do for the third time at no extra cost, and 2 times with the cancelation RAII. Any change will need to be duplicated 3 times as well, etc.

If we do

template<typename RAII>
void emitWorksharingHelper(..., bool &HasLastprivate)
    if (llvm::any_of(S.getClausesOfKind<OMPReductionClause>(),
                     [](const OMPReductionClause *C) {
                       return C->getModifier() == OMPC_REDUCTION_inscan;
                     })) {
      const auto &&NumIteratorsGen = [&S](CodeGenFunction &CGF) {
        OMPLocalDeclMapRAII Scope(CGF);
        OMPLoopScope LoopScope(CGF, S);
        return CGF.EmitScalarExpr(S.getNumIterations());
      };
      const auto &&FirstGen = [&S](CodeGenFunction &CGF) {
        RAII CancelRegion(CGF, OMPD_for, S.hasCancel());
        (void)CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
                                         emitForLoopBounds,
                                         emitDispatchForLoopBounds);
        // Emit an implicit barrier at the end.
        CGF.CGM.getOpenMPRuntime().emitBarrierCall(CGF, S.getBeginLoc(),
                                                   OMPD_for);
      };
      const auto &&SecondGen = [&S, &HasLastprivate](CodeGenFunction &CGF) {
        RAII CancelRegion(CGF, OMPD_for, S.hasCancel());
        HasLastprivate = CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(),
                                         emitForLoopBounds,
                                         emitDispatchForLoopBounds);
      };
      emitScanBasedDirective(CGF, S, NumIteratorsGen, FirstGen, SecondGen);
    } else {
      OMPCancelStackRAII CancelRegion(CGF, OMPD_parallel_for, S.hasCancel());
        CGF.EmitOMPWorksharingLoop(S, S.getEnsureUpperBound(), emitForLoopBounds,
                                   emitDispatchForLoopBounds);
    }

We can call it three times:

emitWorksharingHelper<OMPCancelStackRAII>(...);
emitWorksharingHelper<OMPCancelStackRAII>(...);
emitWorksharingHelper<OMPDummyRAII>(...);

Wouldn't that be cleaner?

jdoerfert: I mean, these ~25 lines will exist 3 times at least. 2 times tacking lastprivate, which we can…
ABataevAuthorUnsubmitted
Done
ReplyInline Actions
  1. It requires the new OMPDummyRAII class.
  2. Thу functionality under control of the else branch also must be affected.
  3. We don't always need to check for HasLastprivate.
  4. The NumIteratorsGen might be different.
  5. We'll have a lot of trouble with adding new functionality to this function, especially if this functionality is unique for one construct and not required for others.
ABataev: 1. It requires the new `OMPDummyRAII` class. 2. Thу functionality under control of the `else`…
jdoerfertUnsubmitted
Not Done
ReplyInline Actions
  1. It requires the new OMPDummyRAII class.

Sure, is it more than this?

struct OMPDummyRAII {
  OMPDummyRAII(...) {}
};
  1. Thу functionality under control of the else branch also must be affected.

Right: sed -e/OMPCancelStackRAII/RAII/ :)

  1. We don't always need to check for HasLastprivate.

There is no cost in doing so though.

  1. The NumIteratorsGen might be different.

Might be or is? Eyeballing this it looks like 3 times exact the same code.

  1. We'll have a lot of trouble with adding new functionality to this function, especially if this functionality is unique for one construct and not required for others.

What new functionality do you expect to be unique between them? If there is stuff planned, we might not want to unify the implementation, if it is just a hunch, we want to go ahead and have a single implemenation so it is easier to read and do global modifications. Finding out that something needs special handling is at the end of the day easier than finding all other (unrelated) locations that need to be changed.

jdoerfert: > 1. It requires the new OMPDummyRAII class. Sure, is it more than this? ``` struct…
ABataevAuthorUnsubmitted
Done
ReplyInline Actions

What new functionality do you expect to be unique between them? If there is stuff planned, we might not want to unify the implementation, if it is just a hunch, we want to go ahead and have a single implemenation so it is easier to read and do global modifications. Finding out that something needs special handling is at the end of the day easier than finding all other (unrelated) locations that need to be changed.

I have a different opinion here. When you will some specific implementation for one construct that is not required for other, you can easily miss it, and it might lead to unexpected results, like a compiler crash or something else. Adding step by step is much better, because you not only modify the code for the one particular construct, but also definitely will modify the tests for this construct. Having a common solution, you may easily miss something.

ABataev: > What new functionality do you expect to be unique between them? If there is stuff planned, we…
} }
// Check for outer lastprivate conditional update. // Check for outer lastprivate conditional update.
checkForLastprivateConditionalUpdate(*this, S); checkForLastprivateConditionalUpdate(*this, S);
} }
void CodeGenFunction::EmitOMPParallelMasterDirective( void CodeGenFunction::EmitOMPParallelMasterDirective(
const OMPParallelMasterDirective &S) { const OMPParallelMasterDirective &S) {
// Emit directive as a combined directive that consists of two implicit // Emit directive as a combined directive that consists of two implicit
▲ Show 20 Lines • Show All 2,896 Lines • Show Last 20 Lines

clang/test/OpenMP/parallel_for_scan_codegen.cpp

  • This file was added.
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -triple x86_64-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -triple x86_64-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
void foo();
void bar();
// CHECK: define void @{{.*}}baz{{.*}}(i32 %n)
void baz(int n) {
static float a[10];
static double b;
// CHECK: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: call i8* @llvm.stacksave()
// CHECK: [[A_BUF_SIZE:%.+]] = mul nuw i64 10, [[NUM_ELEMS:%[^,]+]]
// float a_buffer[10][n];
// CHECK: [[A_BUF:%.+]] = alloca float, i64 [[A_BUF_SIZE]],
// double b_buffer[10];
// CHECK: [[B_BUF:%.+]] = alloca double, i64 10,
#pragma omp parallel for reduction(inscan, +:a[:n], b)
for (int i = 0; i < 10; ++i) {
// CHECK: call void @__kmpc_for_static_init_4(
// CHECK: call i8* @llvm.stacksave()
// CHECK: store float 0.000000e+00, float* %
// CHECK: store double 0.000000e+00, double* [[B_PRIV_ADDR:%.+]],
// CHECK: br label %[[DISPATCH:[^,]+]]
// CHECK: [[INPUT_PHASE:.+]]:
// CHECK: call void @{{.+}}foo{{.+}}()
// a_buffer[i][0..n] = a_priv[[0..n];
// CHECK: [[BASE_IDX_I:%.+]] = load i32, i32* [[IV_ADDR:%.+]],
// CHECK: [[BASE_IDX:%.+]] = zext i32 [[BASE_IDX_I]] to i64
// CHECK: [[IDX:%.+]] = mul nsw i64 [[BASE_IDX]], [[NUM_ELEMS]]
// CHECK: [[A_BUF_IDX:%.+]] = getelementptr inbounds float, float* [[A_BUF]], i64 [[IDX]]
// CHECK: [[A_PRIV:%.+]] = getelementptr inbounds [10 x float], [10 x float]* [[A_PRIV_ADDR:%.+]], i64 0, i64 0
// CHECK: [[BYTES:%.+]] = mul nuw i64 [[NUM_ELEMS:%.+]], 4
// CHECK: [[DEST:%.+]] = bitcast float* [[A_BUF_IDX]] to i8*
// CHECK: [[SRC:%.+]] = bitcast float* [[A_PRIV]] to i8*
// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* {{.*}}[[DEST]], i8* {{.*}}[[SRC]], i64 [[BYTES]], i1 false)
// b_buffer[i] = b_priv;
// CHECK: [[B_BUF_IDX:%.+]] = getelementptr inbounds double, double* [[B_BUF]], i64 [[BASE_IDX]]
// CHECK: [[B_PRIV:%.+]] = load double, double* [[B_PRIV_ADDR]],
// CHECK: store double [[B_PRIV]], double* [[B_BUF_IDX]],
// CHECK: br label %[[LOOP_CONTINUE:.+]]
// CHECK: [[DISPATCH]]:
// CHECK: br label %[[INPUT_PHASE]]
// CHECK: [[LOOP_CONTINUE]]:
// CHECK: call void @llvm.stackrestore(i8* %
// CHECK: call void @__kmpc_for_static_fini(
// CHECK: call void @__kmpc_barrier(
foo();
#pragma omp scan inclusive(a[:n], b)
// CHECK: [[LOG2_10:%.+]] = call double @llvm.log2.f64(double 1.000000e+01)
// CHECK: [[CEIL_LOG2_10:%.+]] = call double @llvm.ceil.f64(double [[LOG2_10]])
// CHECK: [[CEIL_LOG2_10_INT:%.+]] = fptoui double [[CEIL_LOG2_10]] to i32
// CHECK: br label %[[OUTER_BODY:[^,]+]]
// CHECK: [[OUTER_BODY]]:
// CHECK: [[K:%.+]] = phi i32 [ 0, %{{.+}} ], [ [[K_NEXT:%.+]], %{{.+}} ]
// CHECK: [[K2POW:%.+]] = phi i64 [ 1, %{{.+}} ], [ [[K2POW_NEXT:%.+]], %{{.+}} ]
// CHECK: [[CMP:%.+]] = icmp uge i64 9, [[K2POW]]
// CHECK: br i1 [[CMP]], label %[[INNER_BODY:[^,]+]], label %[[INNER_EXIT:[^,]+]]
// CHECK: [[INNER_BODY]]:
// CHECK: [[I:%.+]] = phi i64 [ 9, %[[OUTER_BODY]] ], [ [[I_PREV:%.+]], %{{.+}} ]
// a_buffer[i] += a_buffer[i-pow(2, k)];
// CHECK: [[IDX:%.+]] = mul nsw i64 [[I]], [[NUM_ELEMS]]
// CHECK: [[A_BUF_IDX:%.+]] = getelementptr inbounds float, float* [[A_BUF]], i64 [[IDX]]
// CHECK: [[IDX_SUB_K2POW:%.+]] = sub nuw i64 [[I]], [[K2POW]]
// CHECK: [[IDX:%.+]] = mul nsw i64 [[IDX_SUB_K2POW]], [[NUM_ELEMS]]
// CHECK: [[A_BUF_IDX_SUB_K2POW:%.+]] = getelementptr inbounds float, float* [[A_BUF]], i64 [[IDX]]
// CHECK: [[B_BUF_IDX:%.+]] = getelementptr inbounds double, double* [[B_BUF]], i64 [[I]]
// CHECK: [[IDX_SUB_K2POW:%.+]] = sub nuw i64 [[I]], [[K2POW]]
// CHECK: [[B_BUF_IDX_SUB_K2POW:%.+]] = getelementptr inbounds double, double* [[B_BUF]], i64 [[IDX_SUB_K2POW]]
// CHECK: [[A_BUF_END:%.+]] = getelementptr float, float* [[A_BUF_IDX]], i64 [[NUM_ELEMS]]
// CHECK: [[ISEMPTY:%.+]] = icmp eq float* [[A_BUF_IDX]], [[A_BUF_END]]
// CHECK: br i1 [[ISEMPTY]], label %[[RED_DONE:[^,]+]], label %[[RED_BODY:[^,]+]]
// CHECK: [[RED_BODY]]:
// CHECK: [[A_BUF_IDX_SUB_K2POW_ELEM:%.+]] = phi float* [ [[A_BUF_IDX_SUB_K2POW]], %[[INNER_BODY]] ], [ [[A_BUF_IDX_SUB_K2POW_NEXT:%.+]], %[[RED_BODY]] ]
// CHECK: [[A_BUF_IDX_ELEM:%.+]] = phi float* [ [[A_BUF_IDX]], %[[INNER_BODY]] ], [ [[A_BUF_IDX_NEXT:%.+]], %[[RED_BODY]] ]
// CHECK: [[A_BUF_IDX_VAL:%.+]] = load float, float* [[A_BUF_IDX_ELEM]],
// CHECK: [[A_BUF_IDX_SUB_K2POW_VAL:%.+]] = load float, float* [[A_BUF_IDX_SUB_K2POW_ELEM]],
// CHECK: [[RED:%.+]] = fadd float [[A_BUF_IDX_VAL]], [[A_BUF_IDX_SUB_K2POW_VAL]]
// CHECK: store float [[RED]], float* [[A_BUF_IDX_ELEM]],
// CHECK: [[A_BUF_IDX_NEXT]] = getelementptr float, float* [[A_BUF_IDX_ELEM]], i32 1
// CHECK: [[A_BUF_IDX_SUB_K2POW_NEXT]] = getelementptr float, float* [[A_BUF_IDX_SUB_K2POW_ELEM]], i32 1
// CHECK: [[DONE:%.+]] = icmp eq float* [[A_BUF_IDX_NEXT]], [[A_BUF_END]]
// CHECK: br i1 [[DONE]], label %[[RED_DONE]], label %[[RED_BODY]]
// CHECK: [[RED_DONE]]:
// b_buffer[i] += b_buffer[i-pow(2, k)];
// CHECK: [[B_BUF_IDX_VAL:%.+]] = load double, double* [[B_BUF_IDX]],
// CHECK: [[B_BUF_IDX_SUB_K2POW_VAL:%.+]] = load double, double* [[B_BUF_IDX_SUB_K2POW]],
// CHECK: [[RED:%.+]] = fadd double [[B_BUF_IDX_VAL]], [[B_BUF_IDX_SUB_K2POW_VAL]]
// CHECK: store double [[RED]], double* [[B_BUF_IDX]],
// --i;
// CHECK: [[I_PREV:%.+]] = sub nuw i64 [[I]], 1
// CHECK: [[CMP:%.+]] = icmp uge i64 [[I_PREV]], [[K2POW]]
// CHECK: br i1 [[CMP]], label %[[INNER_BODY]], label %[[INNER_EXIT]]
// CHECK: [[INNER_EXIT]]:
// ++k;
// CHECK: [[K_NEXT]] = add nuw i32 [[K]], 1
// k2pow <<= 1;
// CHECK: [[K2POW_NEXT]] = shl nuw i64 [[K2POW]], 1
// CHECK: [[CMP:%.+]] = icmp ne i32 [[K_NEXT]], [[CEIL_LOG2_10_INT]]
// CHECK: br i1 [[CMP]], label %[[OUTER_BODY]], label %[[OUTER_EXIT:[^,]+]]
// CHECK: [[OUTER_EXIT]]:
bar();
// CHECK: call void @__kmpc_for_static_init_4(
// CHECK: call i8* @llvm.stacksave()
// CHECK: store float 0.000000e+00, float* %
// CHECK: store double 0.000000e+00, double* [[B_PRIV_ADDR:%.+]],
// CHECK: br label %[[DISPATCH:[^,]+]]
// Skip the before scan body.
// CHECK: call void @{{.+}}foo{{.+}}()
// CHECK: [[EXIT_INSCAN:[^,]+]]:
// CHECK: br label %[[LOOP_CONTINUE:[^,]+]]
// CHECK: [[DISPATCH]]:
// a_priv[[0..n] = a_buffer[i][0..n];
// CHECK: [[BASE_IDX_I:%.+]] = load i32, i32* [[IV_ADDR:%.+]],
// CHECK: [[BASE_IDX:%.+]] = zext i32 [[BASE_IDX_I]] to i64
// CHECK: [[IDX:%.+]] = mul nsw i64 [[BASE_IDX]], [[NUM_ELEMS]]
// CHECK: [[A_BUF_IDX:%.+]] = getelementptr inbounds float, float* [[A_BUF]], i64 [[IDX]]
// CHECK: [[A_PRIV:%.+]] = getelementptr inbounds [10 x float], [10 x float]* [[A_PRIV_ADDR:%.+]], i64 0, i64 0
// CHECK: [[BYTES:%.+]] = mul nuw i64 [[NUM_ELEMS:%.+]], 4
// CHECK: [[DEST:%.+]] = bitcast float* [[A_PRIV]] to i8*
// CHECK: [[SRC:%.+]] = bitcast float* [[A_BUF_IDX]] to i8*
// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* {{.*}}[[DEST]], i8* {{.*}}[[SRC]], i64 [[BYTES]], i1 false)
// b_priv = b_buffer[i];
// CHECK: [[B_BUF_IDX:%.+]] = getelementptr inbounds double, double* [[B_BUF]], i64 [[BASE_IDX]]
// CHECK: [[B_BUF_IDX_VAL:%.+]] = load double, double* [[B_BUF_IDX]],
// CHECK: store double [[B_BUF_IDX_VAL]], double* [[B_PRIV_ADDR]],
// CHECK: br label %[[SCAN_PHASE:[^,]+]]
// CHECK: [[SCAN_PHASE]]:
// CHECK: call void @{{.+}}bar{{.+}}()
// CHECK: br label %[[EXIT_INSCAN]]
// CHECK: [[LOOP_CONTINUE]]:
// CHECK: call void @llvm.stackrestore(i8* %
// CHECK: call void @__kmpc_for_static_fini(
// CHECK: call void @llvm.stackrestore(i8*
}
// CHECK: call i8* @llvm.stacksave()
// CHECK: [[A_BUF_SIZE:%.+]] = mul nuw i64 10, [[NUM_ELEMS:%[^,]+]]
// float a_buffer[10][n];
// CHECK: [[A_BUF:%.+]] = alloca float, i64 [[A_BUF_SIZE]],
// double b_buffer[10];
// CHECK: [[B_BUF:%.+]] = alloca double, i64 10,
#pragma omp parallel for reduction(inscan, +:a[:n], b)
for (int i = 0; i < 10; ++i) {
// CHECK: call void @__kmpc_for_static_init_4(
// CHECK: call i8* @llvm.stacksave()
// CHECK: store float 0.000000e+00, float* %
// CHECK: store double 0.000000e+00, double* [[B_PRIV_ADDR:%.+]],
// CHECK: br label %[[DISPATCH:[^,]+]]
// Skip the before scan body.
// CHECK: call void @{{.+}}foo{{.+}}()
// CHECK: [[EXIT_INSCAN:[^,]+]]:
// a_buffer[i][0..n] = a_priv[[0..n];
// CHECK: [[BASE_IDX_I:%.+]] = load i32, i32* [[IV_ADDR:%.+]],
// CHECK: [[BASE_IDX:%.+]] = zext i32 [[BASE_IDX_I]] to i64
// CHECK: [[IDX:%.+]] = mul nsw i64 [[BASE_IDX]], [[NUM_ELEMS]]
// CHECK: [[A_BUF_IDX:%.+]] = getelementptr inbounds float, float* [[A_BUF]], i64 [[IDX]]
// CHECK: [[A_PRIV:%.+]] = getelementptr inbounds [10 x float], [10 x float]* [[A_PRIV_ADDR:%.+]], i64 0, i64 0
// CHECK: [[BYTES:%.+]] = mul nuw i64 [[NUM_ELEMS:%.+]], 4
// CHECK: [[DEST:%.+]] = bitcast float* [[A_BUF_IDX]] to i8*
// CHECK: [[SRC:%.+]] = bitcast float* [[A_PRIV]] to i8*
// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* {{.*}}[[DEST]], i8* {{.*}}[[SRC]], i64 [[BYTES]], i1 false)
// b_buffer[i] = b_priv;
// CHECK: [[B_BUF_IDX:%.+]] = getelementptr inbounds double, double* [[B_BUF]], i64 [[BASE_IDX]]
// CHECK: [[B_PRIV:%.+]] = load double, double* [[B_PRIV_ADDR]],
// CHECK: store double [[B_PRIV]], double* [[B_BUF_IDX]],
// CHECK: br label %[[LOOP_CONTINUE:[^,]+]]
// CHECK: [[DISPATCH]]:
// CHECK: br label %[[INPUT_PHASE:[^,]+]]
// CHECK: [[INPUT_PHASE]]:
// CHECK: call void @{{.+}}bar{{.+}}()
// CHECK: br label %[[EXIT_INSCAN]]
// CHECK: [[LOOP_CONTINUE]]:
// CHECK: call void @llvm.stackrestore(i8* %
// CHECK: call void @__kmpc_for_static_fini(
// CHECK: call void @__kmpc_barrier(
foo();
#pragma omp scan exclusive(a[:n], b)
// CHECK: [[LOG2_10:%.+]] = call double @llvm.log2.f64(double 1.000000e+01)
// CHECK: [[CEIL_LOG2_10:%.+]] = call double @llvm.ceil.f64(double [[LOG2_10]])
// CHECK: [[CEIL_LOG2_10_INT:%.+]] = fptoui double [[CEIL_LOG2_10]] to i32
// CHECK: br label %[[OUTER_BODY:[^,]+]]
// CHECK: [[OUTER_BODY]]:
// CHECK: [[K:%.+]] = phi i32 [ 0, %{{.+}} ], [ [[K_NEXT:%.+]], %{{.+}} ]
// CHECK: [[K2POW:%.+]] = phi i64 [ 1, %{{.+}} ], [ [[K2POW_NEXT:%.+]], %{{.+}} ]
// CHECK: [[CMP:%.+]] = icmp uge i64 9, [[K2POW]]
// CHECK: br i1 [[CMP]], label %[[INNER_BODY:[^,]+]], label %[[INNER_EXIT:[^,]+]]
// CHECK: [[INNER_BODY]]:
// CHECK: [[I:%.+]] = phi i64 [ 9, %[[OUTER_BODY]] ], [ [[I_PREV:%.+]], %{{.+}} ]
// a_buffer[i] += a_buffer[i-pow(2, k)];
// CHECK: [[IDX:%.+]] = mul nsw i64 [[I]], [[NUM_ELEMS]]
// CHECK: [[A_BUF_IDX:%.+]] = getelementptr inbounds float, float* [[A_BUF]], i64 [[IDX]]
// CHECK: [[IDX_SUB_K2POW:%.+]] = sub nuw i64 [[I]], [[K2POW]]
// CHECK: [[IDX:%.+]] = mul nsw i64 [[IDX_SUB_K2POW]], [[NUM_ELEMS]]
// CHECK: [[A_BUF_IDX_SUB_K2POW:%.+]] = getelementptr inbounds float, float* [[A_BUF]], i64 [[IDX]]
// CHECK: [[B_BUF_IDX:%.+]] = getelementptr inbounds double, double* [[B_BUF]], i64 [[I]]
// CHECK: [[IDX_SUB_K2POW:%.+]] = sub nuw i64 [[I]], [[K2POW]]
// CHECK: [[B_BUF_IDX_SUB_K2POW:%.+]] = getelementptr inbounds double, double* [[B_BUF]], i64 [[IDX_SUB_K2POW]]
// CHECK: [[A_BUF_END:%.+]] = getelementptr float, float* [[A_BUF_IDX]], i64 [[NUM_ELEMS]]
// CHECK: [[ISEMPTY:%.+]] = icmp eq float* [[A_BUF_IDX]], [[A_BUF_END]]
// CHECK: br i1 [[ISEMPTY]], label %[[RED_DONE:[^,]+]], label %[[RED_BODY:[^,]+]]
// CHECK: [[RED_BODY]]:
// CHECK: [[A_BUF_IDX_SUB_K2POW_ELEM:%.+]] = phi float* [ [[A_BUF_IDX_SUB_K2POW]], %[[INNER_BODY]] ], [ [[A_BUF_IDX_SUB_K2POW_NEXT:%.+]], %[[RED_BODY]] ]
// CHECK: [[A_BUF_IDX_ELEM:%.+]] = phi float* [ [[A_BUF_IDX]], %[[INNER_BODY]] ], [ [[A_BUF_IDX_NEXT:%.+]], %[[RED_BODY]] ]
// CHECK: [[A_BUF_IDX_VAL:%.+]] = load float, float* [[A_BUF_IDX_ELEM]],
// CHECK: [[A_BUF_IDX_SUB_K2POW_VAL:%.+]] = load float, float* [[A_BUF_IDX_SUB_K2POW_ELEM]],
// CHECK: [[RED:%.+]] = fadd float [[A_BUF_IDX_VAL]], [[A_BUF_IDX_SUB_K2POW_VAL]]
// CHECK: store float [[RED]], float* [[A_BUF_IDX_ELEM]],
// CHECK: [[A_BUF_IDX_NEXT]] = getelementptr float, float* [[A_BUF_IDX_ELEM]], i32 1
// CHECK: [[A_BUF_IDX_SUB_K2POW_NEXT]] = getelementptr float, float* [[A_BUF_IDX_SUB_K2POW_ELEM]], i32 1
// CHECK: [[DONE:%.+]] = icmp eq float* [[A_BUF_IDX_NEXT]], [[A_BUF_END]]
// CHECK: br i1 [[DONE]], label %[[RED_DONE]], label %[[RED_BODY]]
// CHECK: [[RED_DONE]]:
// b_buffer[i] += b_buffer[i-pow(2, k)];
// CHECK: [[B_BUF_IDX_VAL:%.+]] = load double, double* [[B_BUF_IDX]],
// CHECK: [[B_BUF_IDX_SUB_K2POW_VAL:%.+]] = load double, double* [[B_BUF_IDX_SUB_K2POW]],
// CHECK: [[RED:%.+]] = fadd double [[B_BUF_IDX_VAL]], [[B_BUF_IDX_SUB_K2POW_VAL]]
// CHECK: store double [[RED]], double* [[B_BUF_IDX]],
// --i;
// CHECK: [[I_PREV:%.+]] = sub nuw i64 [[I]], 1
// CHECK: [[CMP:%.+]] = icmp uge i64 [[I_PREV]], [[K2POW]]
// CHECK: br i1 [[CMP]], label %[[INNER_BODY]], label %[[INNER_EXIT]]
// CHECK: [[INNER_EXIT]]:
// ++k;
// CHECK: [[K_NEXT]] = add nuw i32 [[K]], 1
// k2pow <<= 1;
// CHECK: [[K2POW_NEXT]] = shl nuw i64 [[K2POW]], 1
// CHECK: [[CMP:%.+]] = icmp ne i32 [[K_NEXT]], [[CEIL_LOG2_10_INT]]
// CHECK: br i1 [[CMP]], label %[[OUTER_BODY]], label %[[OUTER_EXIT:[^,]+]]
// CHECK: [[OUTER_EXIT]]:
bar();
// CHECK: call void @__kmpc_for_static_init_4(
// CHECK: call i8* @llvm.stacksave()
// CHECK: store float 0.000000e+00, float* %
// CHECK: store double 0.000000e+00, double* [[B_PRIV_ADDR:%.+]],
// CHECK: br label %[[DISPATCH:[^,]+]]
// CHECK: [[SCAN_PHASE:.+]]:
// CHECK: call void @{{.+}}foo{{.+}}()
// CHECK: br label %[[LOOP_CONTINUE:.+]]
// CHECK: [[DISPATCH]]:
// if (i >0)
// a_priv[[0..n] = a_buffer[i-1][0..n];
// CHECK: [[BASE_IDX_I:%.+]] = load i32, i32* [[IV_ADDR:%.+]],
// CHECK: [[BASE_IDX:%.+]] = zext i32 [[BASE_IDX_I]] to i64
// CHECK: [[CMP:%.+]] = icmp eq i64 [[BASE_IDX]], 0
// CHECK: br i1 [[CMP]], label %[[IF_DONE:[^,]+]], label %[[IF_THEN:[^,]+]]
// CHECK: [[IF_THEN]]:
// CHECK: [[BASE_IDX_SUB_1:%.+]] = sub nuw i64 [[BASE_IDX]], 1
// CHECK: [[IDX:%.+]] = mul nsw i64 [[BASE_IDX_SUB_1]], [[NUM_ELEMS]]
// CHECK: [[A_BUF_IDX:%.+]] = getelementptr inbounds float, float* [[A_BUF]], i64 [[IDX]]
// CHECK: [[A_PRIV:%.+]] = getelementptr inbounds [10 x float], [10 x float]* [[A_PRIV_ADDR:%.+]], i64 0, i64 0
// CHECK: [[BYTES:%.+]] = mul nuw i64 [[NUM_ELEMS:%.+]], 4
// CHECK: [[DEST:%.+]] = bitcast float* [[A_PRIV]] to i8*
// CHECK: [[SRC:%.+]] = bitcast float* [[A_BUF_IDX]] to i8*
// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* {{.*}}[[DEST]], i8* {{.*}}[[SRC]], i64 [[BYTES]], i1 false)
// b_priv = b_buffer[i];
// CHECK: [[B_BUF_IDX:%.+]] = getelementptr inbounds double, double* [[B_BUF]], i64 [[BASE_IDX_SUB_1]]
// CHECK: [[B_BUF_IDX_VAL:%.+]] = load double, double* [[B_BUF_IDX]],
// CHECK: store double [[B_BUF_IDX_VAL]], double* [[B_PRIV_ADDR]],
// CHECK: br label %[[SCAN_PHASE]]
// CHECK: [[LOOP_CONTINUE]]:
// CHECK: call void @llvm.stackrestore(i8* %
// CHECK: call void @__kmpc_for_static_fini(
// CHECK: call void @llvm.stackrestore(i8*
}
}
#endif