Skip to content

Commit b1bd36a

Browse files
wsmoseswsmoses
authored
Correct CacheLICM behavior when not the outermost loop (rust-lang#236)
* Allow cachelicm within subregion * Correct CacheLICM
1 parent d4d3ae9 commit b1bd36a

File tree

5 files changed

+371
-19
lines changed

5 files changed

+371
-19
lines changed

‎enzyme/Enzyme/CacheUtility.cpp

+20-9
Original file line numberDiff line numberDiff line change
@@ -659,7 +659,8 @@ AllocaInst *CacheUtility::createCacheForScope(LimitContext ctx, Type *T,
659659
assert(ctx.Block);
660660
assert(T);
661661

662-
auto sublimits = getSubLimits(/*inForwardPass*/ true, nullptr, ctx);
662+
auto sublimits =
663+
getSubLimits(/*inForwardPass*/ true, nullptr, ctx, extraSize);
663664

664665
// List of types stored in the cache for each Loop-Chunk
665666
// This is stored from innner-most chunk to outermost
@@ -944,7 +945,11 @@ Value *CacheUtility::computeIndexOfChunk(
944945
/// innermost loop to outermost loop.
945946
CacheUtility::SubLimitType CacheUtility::getSubLimits(bool inForwardPass,
946947
IRBuilder<> *RB,
947-
LimitContext ctx) {
948+
LimitContext ctx,
949+
Value *extraSize) {
950+
// Store the LoopContext's in InnerMost => Outermost order
951+
std::vector<LoopContext> contexts;
952+
948953
// Given a ``SingleIteration'' Limit Context, return a chunking of
949954
// one loop with size 1, and header/preheader of the BasicBlock
950955
// This is done to create a context for a block outside a loop
@@ -954,7 +959,6 @@ CacheUtility::SubLimitType CacheUtility::getSubLimits(bool inForwardPass,
954959
LoopContext idx;
955960
auto subctx = ctx.Block;
956961
auto zero = ConstantInt::get(Type::getInt64Ty(newFunc->getContext()), 0);
957-
auto one = ConstantInt::get(Type::getInt64Ty(newFunc->getContext()), 1);
958962
// The iteration count is always zero so we can set it as such
959963
idx.var = nullptr; // = zero;
960964
idx.incvar = nullptr;
@@ -966,13 +970,9 @@ CacheUtility::SubLimitType CacheUtility::getSubLimits(bool inForwardPass,
966970
idx.dynamic = false;
967971
idx.parent = nullptr;
968972
idx.exitBlocks = {};
969-
SubLimitType sublimits;
970-
sublimits.push_back({one, {{idx, one}}});
971-
return sublimits;
973+
contexts.push_back(idx);
972974
}
973975

974-
// Store the LoopContext's in InnerMost => Outermost order
975-
std::vector<LoopContext> contexts;
976976
for (BasicBlock *blk = ctx.Block; blk != nullptr;) {
977977
LoopContext idx;
978978
if (!getContext(blk, idx, ctx.ReverseLimit)) {
@@ -1057,6 +1057,17 @@ CacheUtility::SubLimitType CacheUtility::getSubLimits(bool inForwardPass,
10571057
allocationBuilder.SetInsertPoint(&allocationPreheaders[i]->back());
10581058
limitMinus1 = unwrapM(contexts[i].maxLimit, allocationBuilder, prevMap,
10591059
UnwrapMode::AttemptFullUnwrap);
1060+
} else if (i == 0 && extraSize &&
1061+
unwrapM(extraSize, allocationBuilder, prevMap,
1062+
UnwrapMode::AttemptFullUnwrap) == nullptr) {
1063+
EmitWarning(
1064+
"NoOuterLimit", cast<Instruction>(extraSize)->getDebugLoc(),
1065+
newFunc, cast<Instruction>(extraSize)->getParent(),
1066+
"Could not compute outermost loop limit by moving extraSize value ",
1067+
*extraSize, " computed at block", contexts[i].header->getName(),
1068+
" function ", contexts[i].header->getParent()->getName());
1069+
allocationPreheaders[i] = contexts[i].preheader;
1070+
allocationBuilder.SetInsertPoint(&allocationPreheaders[i]->back());
10601071
}
10611072
assert(limitMinus1 != nullptr);
10621073

@@ -1295,7 +1306,7 @@ Value *CacheUtility::getCachePointer(bool inForwardPass, IRBuilder<> &BuilderM,
12951306
assert(ctx.Block);
12961307
assert(cache);
12971308

1298-
auto sublimits = getSubLimits(inForwardPass, &BuilderM, ctx);
1309+
auto sublimits = getSubLimits(inForwardPass, &BuilderM, ctx, extraSize);
12991310

13001311
ValueToValueMapTy available;
13011312

‎enzyme/Enzyme/CacheUtility.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -223,7 +223,7 @@ class CacheUtility {
223223
/*loop limits*/ std::vector<std::pair<LoopContext, llvm::Value *>>>>
224224
SubLimitType;
225225
SubLimitType getSubLimits(bool inForwardPass, llvm::IRBuilder<> *RB,
226-
LimitContext ctx);
226+
LimitContext ctx, llvm::Value *extraSize = nullptr);
227227

228228
private:
229229
/// Internal data structure used by getSubLimit to avoid computing the same

‎enzyme/Enzyme/GradientUtils.cpp

+27-9
Original file line numberDiff line numberDiff line change
@@ -199,12 +199,19 @@ Value *GradientUtils::unwrapM(Value *const val, IRBuilder<> &BuilderM,
199199
___res = \
200200
unwrapM(v, Builder, available, mode, origParent, permitCache); \
201201
if (!___res && mode == UnwrapMode::AttemptFullUnwrapWithLookup) { \
202+
bool noLookup = false; \
203+
if (auto opinst = dyn_cast<Instruction>(v)) \
204+
if (isOriginalBlock(*Builder.GetInsertBlock())) { \
205+
if (!DT.dominates(opinst, &*Builder.GetInsertPoint())) \
206+
noLookup = true; \
207+
} \
202208
if (origParent) \
203209
if (auto opinst = dyn_cast<Instruction>(v)) { \
204210
v = fixLCSSA(opinst, origParent, /*mergeIfTrue*/ false, \
205211
/*guaranteedVisible*/ false); \
206212
} \
207-
___res = lookupM(v, Builder, available, v != val); \
213+
if (!noLookup) \
214+
___res = lookupM(v, Builder, available, v != val); \
208215
} \
209216
if (___res) \
210217
assert(___res->getType() == v->getType() && "uw"); \
@@ -450,11 +457,6 @@ Value *GradientUtils::unwrapM(Value *const val, IRBuilder<> &BuilderM,
450457
auto toreturn = BuilderM.CreateGEP(ptr, ind, inst->getName() + "_unwrap");
451458
if (isa<GetElementPtrInst>(toreturn))
452459
cast<GetElementPtrInst>(toreturn)->setIsInBounds(inst->isInBounds());
453-
else {
454-
// llvm::errs() << "gep tr: " << *toreturn << " inst: " << *inst << "
455-
// ptr: " << *ptr << "\n"; llvm::errs() << "safe: " << *SAFE(inst,
456-
// getPointerOperand()) << "\n"; assert(0 && "illegal");
457-
}
458460
if (auto newi = dyn_cast<Instruction>(toreturn))
459461
newi->copyIRFlags(inst);
460462
if (permitCache)
@@ -629,7 +631,9 @@ Value *GradientUtils::unwrapM(Value *const val, IRBuilder<> &BuilderM,
629631
if (!isOriginalBlock(*ivctx)) {
630632
ivctx = originalForReverseBlock(*ivctx);
631633
}
632-
if (ivctx == phi->getParent() || DT.dominates(phi, ivctx)) {
634+
if ((ivctx == phi->getParent() || DT.dominates(phi, ivctx)) &&
635+
(!isOriginalBlock(*BuilderM.GetInsertBlock()) ||
636+
DT.dominates(phi, &*BuilderM.GetInsertPoint()))) {
633637
LoopContext lc;
634638
bool loopVar = false;
635639
if (getContext(phi->getParent(), lc) && lc.var == phi) {
@@ -1165,6 +1169,13 @@ Value *GradientUtils::unwrapM(Value *const val, IRBuilder<> &BuilderM,
11651169
mode == UnwrapMode::AttemptFullUnwrapWithLookup) {
11661170
assert(val->getName() != "<badref>");
11671171
Value *nval = val;
1172+
if (auto opinst = dyn_cast<Instruction>(nval))
1173+
if (isOriginalBlock(*BuilderM.GetInsertBlock())) {
1174+
if (!DT.dominates(opinst, &*BuilderM.GetInsertPoint())) {
1175+
assert(mode == UnwrapMode::AttemptFullUnwrapWithLookup);
1176+
return nullptr;
1177+
}
1178+
}
11681179
if (scope)
11691180
if (auto opinst = dyn_cast<Instruction>(nval)) {
11701181
nval = fixLCSSA(opinst, scope, /*mergeIfTrue*/ false,
@@ -3653,6 +3664,8 @@ Value *GradientUtils::lookupM(Value *val, IRBuilder<> &BuilderM,
36533664
}
36543665

36553666
if (ctx && lim && start && offset) {
3667+
Value *firstLim = lim;
3668+
Value *firstStart = start;
36563669
while (Loop *L = LI.getLoopFor(ctx)) {
36573670
BasicBlock *nctx = L->getLoopPreheader();
36583671
assert(nctx);
@@ -3676,12 +3689,12 @@ Value *GradientUtils::lookupM(Value *val, IRBuilder<> &BuilderM,
36763689
if (failed)
36773690
break;
36783691
IRBuilder<> nv(nctx->getTerminator());
3679-
Value *nlim = unwrapM(lim, nv,
3692+
Value *nlim = unwrapM(firstLim, nv,
36803693
/*available*/ ValueToValueMapTy(),
36813694
UnwrapMode::AttemptFullUnwrapWithLookup);
36823695
if (!nlim)
36833696
break;
3684-
Value *nstart = unwrapM(start, nv,
3697+
Value *nstart = unwrapM(firstStart, nv,
36853698
/*available*/ ValueToValueMapTy(),
36863699
UnwrapMode::AttemptFullUnwrapWithLookup);
36873700
if (!nstart)
@@ -3700,6 +3713,11 @@ Value *GradientUtils::lookupM(Value *val, IRBuilder<> &BuilderM,
37003713
bool forceSingleIter = false;
37013714
if (!getContext(ctx, tmp)) {
37023715
forceSingleIter = true;
3716+
} else if (auto inst = dyn_cast<Instruction>(lim)) {
3717+
if (inst->getParent() == ctx ||
3718+
!DT.dominates(inst->getParent(), ctx)) {
3719+
forceSingleIter = true;
3720+
}
37033721
}
37043722
LimitContext lctx(/*ReverseLimit*/ reverseBlocks.size() > 0, ctx,
37053723
forceSingleIter);

‎enzyme/test/Enzyme/ReverseMode/metacachelicm.ll

+163
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,163 @@
1+
; RUN: %opt < %s %loadEnzyme -enzyme -enzyme-preopt=false -mem2reg -instsimplify -simplifycfg -S | FileCheck %s
2+
3+
; Function Attrs: nounwind uwtable
4+
define dso_local void @compute(double* noalias nocapture %data, i64* noalias nocapture readnone %array, double* noalias nocapture %out) #0 {
5+
entry:
6+
br label %for.body5.preheader
7+
8+
for.cond.cleanup: ; preds = %for.cond.cleanup4
9+
store double 0.000000e+00, double* %data, align 8, !tbaa !2
10+
ret void
11+
12+
for.body5.preheader: ; preds = %entry, %for.cond.cleanup4
13+
%indvars.iv = phi i64 [ 0, %entry ], [ %indvars.iv.next, %for.cond.cleanup4 ]
14+
%call = tail call i64 @getSize() #2
15+
br label %for.body5
16+
17+
for.cond.cleanup4: ; preds = %for.body5
18+
%arrayidx7 = getelementptr inbounds double, double* %out, i64 %indvars.iv
19+
store double %add, double* %arrayidx7, align 8, !tbaa !2
20+
%indvars.iv.next = add nuw nsw i64 %indvars.iv, 1
21+
%exitcond29 = icmp eq i64 %indvars.iv.next, 10
22+
br i1 %exitcond29, label %for.cond.cleanup, label %for.body5.preheader
23+
24+
for.body5: ; preds = %for.body5, %for.body5.preheader
25+
%j.027 = phi i64 [ %inc, %for.body5 ], [ 0, %for.body5.preheader ]
26+
%res.026 = phi double [ %add, %for.body5 ], [ 0.000000e+00, %for.body5.preheader ]
27+
%arrayidx = getelementptr inbounds double, double* %data, i64 %j.027
28+
%i0 = load double, double* %arrayidx, align 8, !tbaa !2
29+
%mul = fmul double %i0, %i0
30+
%add = fadd double %res.026, %mul
31+
%inc = add nuw i64 %j.027, 1
32+
%exitcond = icmp eq i64 %inc, %call
33+
br i1 %exitcond, label %for.cond.cleanup4, label %for.body5
34+
}
35+
36+
declare dso_local i64 @getSize() local_unnamed_addr #1
37+
38+
; Function Attrs: nounwind
39+
declare void @llvm.assume(i1) #2
40+
41+
; Function Attrs: nounwind uwtable
42+
define dso_local void @call(double* %data, double* %d_data, i64* %array, double* %out, double* %d_out) local_unnamed_addr #0 {
43+
entry:
44+
tail call void (i8*, ...) @__enzyme_autodiff(i8* bitcast (void (double*, i64*, double*)* @compute to i8*), double* %data, double* %d_data, i64* %array, double* %out, double* %d_out) #2
45+
ret void
46+
}
47+
48+
declare dso_local void @__enzyme_autodiff(i8*, ...) local_unnamed_addr #1
49+
50+
attributes #0 = { nounwind uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" }
51+
attributes #1 = { "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" }
52+
attributes #2 = { nounwind }
53+
54+
!llvm.module.flags = !{!0}
55+
!llvm.ident = !{!1}
56+
57+
!0 = !{i32 1, !"wchar_size", i32 4}
58+
!1 = !{!"clang version 7.1.0 "}
59+
!2 = !{!3, !3, i64 0}
60+
!3 = !{!"double", !4, i64 0}
61+
!4 = !{!"omnipotent char", !5, i64 0}
62+
!5 = !{!"Simple C/C++ TBAA"}
63+
64+
; CHECK: define internal void @diffecompute(double* noalias nocapture %data, double* nocapture %"data'", i64* noalias nocapture readnone %array, double* noalias nocapture %out, double* nocapture %"out'")
65+
; CHECK-NEXT: entry:
66+
; CHECK-NEXT: %malloccall = tail call noalias nonnull dereferenceable(80) dereferenceable_or_null(80) i8* @malloc(i64 80)
67+
; CHECK-NEXT: %_malloccache = bitcast i8* %malloccall to i64*
68+
; CHECK-NEXT: %malloccall3 = tail call noalias nonnull dereferenceable(80) dereferenceable_or_null(80) i8* @malloc(i64 80)
69+
; CHECK-NEXT: %i0_malloccache = bitcast i8* %malloccall3 to double**
70+
; CHECK-NEXT: br label %for.body5.preheader
71+
72+
; CHECK: for.cond.cleanup: ; preds = %for.cond.cleanup4
73+
; CHECK-NEXT: store double 0.000000e+00, double* %data, align 8, !tbaa !2
74+
; CHECK-NEXT: store double 0.000000e+00, double* %"data'", align 8
75+
; CHECK-NEXT: br label %invertfor.cond.cleanup4
76+
77+
; CHECK: for.body5.preheader: ; preds = %for.cond.cleanup4, %entry
78+
; CHECK-NEXT: %iv = phi i64 [ %iv.next, %for.cond.cleanup4 ], [ 0, %entry ]
79+
; CHECK-NEXT: %iv.next = add nuw nsw i64 %iv, 1
80+
; CHECK-NEXT: %call = tail call i64 @getSize()
81+
; CHECK-NEXT: %0 = add i64 %call, -1
82+
; CHECK-NEXT: %1 = getelementptr inbounds i64, i64* %_malloccache, i64 %iv
83+
; CHECK-NEXT: store i64 %0, i64* %1, align 8, !invariant.group !6
84+
; CHECK-NEXT: %2 = getelementptr inbounds double*, double** %i0_malloccache, i64 %iv
85+
; CHECK-NEXT: %mallocsize = mul nuw nsw i64 %call, 8
86+
; CHECK-NEXT: %malloccall5 = tail call noalias nonnull i8* @malloc(i64 %mallocsize)
87+
; CHECK-NEXT: %i0_malloccache6 = bitcast i8* %malloccall5 to double*
88+
; CHECK-NEXT: store double* %i0_malloccache6, double** %2, align 8, !invariant.group !7
89+
; CHECK-NEXT: %3 = getelementptr inbounds double*, double** %i0_malloccache, i64 %iv
90+
; CHECK-NEXT: %4 = load double*, double** %3, align 8, !dereferenceable !8, !invariant.group !7
91+
; CHECK-NEXT: %5 = bitcast double* %4 to i8*
92+
; CHECK-NEXT: %6 = bitcast double* %data to i8*
93+
; CHECK-NEXT: %7 = mul nuw nsw i64 8, %call
94+
; CHECK-NEXT: call void @llvm.memcpy.p0i8.p0i8.i64(i8* nonnull align 8 %5, i8* nonnull align 8 %6, i64 %7, i1 false)
95+
; CHECK-NEXT: br label %for.body5
96+
97+
; CHECK: for.cond.cleanup4: ; preds = %for.body5
98+
; CHECK-NEXT: %arrayidx7 = getelementptr inbounds double, double* %out, i64 %iv
99+
; CHECK-NEXT: store double %add, double* %arrayidx7, align 8, !tbaa !2
100+
; CHECK-NEXT: %exitcond29 = icmp eq i64 %iv.next, 10
101+
; CHECK-NEXT: br i1 %exitcond29, label %for.cond.cleanup, label %for.body5.preheader
102+
103+
; CHECK: for.body5: ; preds = %for.body5, %for.body5.preheader
104+
; CHECK-NEXT: %iv1 = phi i64 [ %iv.next2, %for.body5 ], [ 0, %for.body5.preheader ]
105+
; CHECK-NEXT: %res.026 = phi double [ %add, %for.body5 ], [ 0.000000e+00, %for.body5.preheader ]
106+
; CHECK-NEXT: %iv.next2 = add nuw nsw i64 %iv1, 1
107+
; CHECK-NEXT: %arrayidx = getelementptr inbounds double, double* %data, i64 %iv1
108+
; CHECK-NEXT: %i0 = load double, double* %arrayidx, align 8, !tbaa !2
109+
; CHECK-NEXT: %mul = fmul double %i0, %i0
110+
; CHECK-NEXT: %add = fadd double %res.026, %mul
111+
; CHECK-NEXT: %exitcond = icmp eq i64 %iv.next2, %call
112+
; CHECK-NEXT: br i1 %exitcond, label %for.cond.cleanup4, label %for.body5
113+
114+
; CHECK: invertentry: ; preds = %invertfor.body5.preheader
115+
; CHECK-NEXT: tail call void @free(i8* nonnull %malloccall)
116+
; CHECK-NEXT: tail call void @free(i8* nonnull %malloccall3)
117+
; CHECK-NEXT: ret void
118+
119+
; CHECK: invertfor.body5.preheader: ; preds = %invertfor.body5
120+
; CHECK-NEXT: %8 = icmp eq i64 %"iv'ac.0", 0
121+
; CHECK-NEXT: %_unwrap7 = getelementptr inbounds double*, double** %i0_malloccache, i64 %"iv'ac.0"
122+
; CHECK-NEXT: %forfree8 = load double*, double** %_unwrap7, align 8, !dereferenceable !8, !invariant.group !7
123+
; CHECK-NEXT: %9 = bitcast double* %forfree8 to i8*
124+
; CHECK-NEXT: tail call void @free(i8* nonnull %9)
125+
; CHECK-NEXT: br i1 %8, label %invertentry, label %incinvertfor.body5.preheader
126+
127+
; CHECK: incinvertfor.body5.preheader: ; preds = %invertfor.body5.preheader
128+
; CHECK-NEXT: %10 = add nsw i64 %"iv'ac.0", -1
129+
; CHECK-NEXT: br label %invertfor.cond.cleanup4
130+
131+
; CHECK: invertfor.cond.cleanup4: ; preds = %incinvertfor.body5.preheader, %for.cond.cleanup
132+
; CHECK-NEXT: %"add'de.0" = phi double [ 0.000000e+00, %for.cond.cleanup ], [ %23, %incinvertfor.body5.preheader ]
133+
; CHECK-NEXT: %"iv'ac.0" = phi i64 [ 9, %for.cond.cleanup ], [ %10, %incinvertfor.body5.preheader ]
134+
; CHECK-NEXT: %"arrayidx7'ipg_unwrap" = getelementptr inbounds double, double* %"out'", i64 %"iv'ac.0"
135+
; CHECK-NEXT: %11 = load double, double* %"arrayidx7'ipg_unwrap", align 8
136+
; CHECK-NEXT: store double 0.000000e+00, double* %"arrayidx7'ipg_unwrap", align 8
137+
; CHECK-NEXT: %12 = fadd fast double %"add'de.0", %11
138+
; CHECK-NEXT: %13 = getelementptr inbounds i64, i64* %_malloccache, i64 %"iv'ac.0"
139+
; CHECK-NEXT: %14 = load i64, i64* %13, align 8, !invariant.group !6
140+
; CHECK-NEXT: br label %invertfor.body5
141+
142+
; CHECK: invertfor.body5: ; preds = %incinvertfor.body5, %invertfor.cond.cleanup4
143+
; CHECK-NEXT: %"add'de.1" = phi double [ %12, %invertfor.cond.cleanup4 ], [ %23, %incinvertfor.body5 ]
144+
; CHECK-NEXT: %"iv1'ac.0" = phi i64 [ %14, %invertfor.cond.cleanup4 ], [ %24, %incinvertfor.body5 ]
145+
; CHECK-NEXT: %15 = getelementptr inbounds double*, double** %i0_malloccache, i64 %"iv'ac.0"
146+
; CHECK-NEXT: %16 = load double*, double** %15, align 8, !dereferenceable !8, !invariant.group !7
147+
; CHECK-NEXT: %17 = getelementptr inbounds double, double* %16, i64 %"iv1'ac.0"
148+
; CHECK-NEXT: %18 = load double, double* %17, align 8, !invariant.group !9
149+
; CHECK-NEXT: %m0diffei0 = fmul fast double %"add'de.1", %18
150+
; CHECK-NEXT: %m1diffei0 = fmul fast double %"add'de.1", %18
151+
; CHECK-NEXT: %19 = fadd fast double %m0diffei0, %m1diffei0
152+
; CHECK-NEXT: %"arrayidx'ipg_unwrap" = getelementptr inbounds double, double* %"data'", i64 %"iv1'ac.0"
153+
; CHECK-NEXT: %20 = load double, double* %"arrayidx'ipg_unwrap", align 8
154+
; CHECK-NEXT: %21 = fadd fast double %20, %19
155+
; CHECK-NEXT: store double %21, double* %"arrayidx'ipg_unwrap", align 8
156+
; CHECK-NEXT: %22 = icmp eq i64 %"iv1'ac.0", 0
157+
; CHECK-NEXT: %23 = select{{( fast)?}} i1 %22, double 0.000000e+00, double %"add'de.1"
158+
; CHECK-NEXT: br i1 %22, label %invertfor.body5.preheader, label %incinvertfor.body5
159+
160+
; CHECK: incinvertfor.body5: ; preds = %invertfor.body5
161+
; CHECK-NEXT: %24 = add nsw i64 %"iv1'ac.0", -1
162+
; CHECK-NEXT: br label %invertfor.body5
163+
; CHECK-NEXT: }

0 commit comments

Comments
 (0)