private static StrideConstraint BuildAddStrideConstraint(GPUVerifier verifier, Expr e, StrideConstraint lhsc, StrideConstraint rhsc)
{
if (lhsc is EqStrideConstraint && rhsc is EqStrideConstraint)
{
return(new EqStrideConstraint(e));
}
if (lhsc is EqStrideConstraint && rhsc is ModStrideConstraint)
{
return(BuildAddStrideConstraint(verifier, e, rhsc, lhsc));
}
if (lhsc is ModStrideConstraint && rhsc is EqStrideConstraint)
{
var lhsmc = (ModStrideConstraint)lhsc;
var rhsec = (EqStrideConstraint)rhsc;
return(new ModStrideConstraint(lhsmc.Mod, verifier.IntRep.MakeAdd(lhsmc.ModEq, rhsec.Eq)));
}
if (lhsc is ModStrideConstraint && rhsc is ModStrideConstraint)
{
var lhsmc = (ModStrideConstraint)lhsc;
var rhsmc = (ModStrideConstraint)rhsc;
if (lhsmc.Mod == rhsmc.Mod)
{
return(new ModStrideConstraint(lhsmc.Mod, verifier.IntRep.MakeAdd(lhsmc.ModEq, rhsmc.ModEq)));
}
}
return(Bottom(verifier, e));
}
public LiteralIndexVisitor(GPUVerifier verifier)
{
foreach (var v in verifier.KernelArrayInfo.GetPrivateArrays())
{
this.LiteralIndexedArrays[v.Name] = new HashSet <string>();
}
}
internal HashSet <Variable> FindWrittenGroupSharedArrays(GPUVerifier verifier)
{
// We add any group-shared array that may be written to or accessed atomically
// in the region.
//
// We also add any group-shared array that may be written to by an asynchronous
// memory copy somewhere in the kernel. This is because asynchronous copies can
// cross barriers. Currently we are very conservative about this.
HashSet <Variable> result = new HashSet <Variable>();
foreach (var v in verifier.KernelArrayInfo.GetGroupSharedArrays(false))
{
if (verifier.ArraysAccessedByAsyncWorkGroupCopy[AccessType.WRITE].Contains(v.Name))
{
result.Add(v);
}
}
foreach (var m in Blocks.Select(Item => Item.Cmds).SelectMany(Item => Item).OfType <CallCmd>()
.Select(Item => Item.Proc.Modifies).SelectMany(Item => Item))
{
// m is a variable modified by a call in the barrier interval
Variable v;
if (verifier.TryGetArrayFromAccessHasOccurred(GVUtil.StripThreadIdentifier(m.Name), AccessType.WRITE, out v) ||
verifier.TryGetArrayFromAccessHasOccurred(GVUtil.StripThreadIdentifier(m.Name), AccessType.ATOMIC, out v))
{
if (verifier.KernelArrayInfo.GetGroupSharedArrays(false).Contains(v))
{
result.Add(v);
}
}
}
return(result);
}
private static bool IsDisjunctionOfPredicates(Expr guard)
{
if (!(guard is NAryExpr))
{
return(false);
}
NAryExpr nary = (NAryExpr)guard;
if (nary.Args.Count() != 2)
{
return(false);
}
if (!(nary.Fun is BinaryOperator))
{
return(false);
}
BinaryOperator binOp = (BinaryOperator)nary.Fun;
if (binOp.Op != BinaryOperator.Opcode.Or)
{
return(false);
}
if (!(nary.Args[0] is IdentifierExpr && nary.Args[1] is IdentifierExpr))
{
return(false);
}
return(GPUVerifier.IsPredicate(GVUtil.StripThreadIdentifier(
((IdentifierExpr)nary.Args[0]).Name)) &&
GPUVerifier.IsPredicate(GVUtil.StripThreadIdentifier(
((IdentifierExpr)nary.Args[1]).Name)));
}
private static Expr MaybeExtractGuard(GPUVerifier verifier, Implementation impl, Block b)
{
if (b.Cmds.Count() > 0)
{
var a = b.Cmds[0] as AssumeCmd;
if (a != null && QKeyValue.FindBoolAttribute(a.Attributes, "partition"))
{
if (a.Expr is IdentifierExpr)
{
return(verifier.varDefAnalysesRegion[impl].DefOfVariableName(((IdentifierExpr)a.Expr).Name));
}
else if (a.Expr is NAryExpr)
{
var nary = (NAryExpr)a.Expr;
if (nary.Fun is UnaryOperator &&
(nary.Fun as UnaryOperator).Op == UnaryOperator.Opcode.Not &&
nary.Args[0] is IdentifierExpr)
{
var d = verifier.varDefAnalysesRegion[impl].DefOfVariableName(((IdentifierExpr)(a.Expr as NAryExpr).Args[0]).Name);
if (d == null)
{
return(null);
}
else
{
return(Expr.Not(d));
}
}
}
}
}
return(null);
}
public VariableDualiser(int id, GPUVerifier verifier, string procName)
{
this.id = id;
this.verifier = verifier;
this.uniformityAnalyser = verifier.UniformityAnalyser;
this.procName = procName;
this.quantifiedVars = new HashSet<Variable>();
}
internal LiteralIndexVisitor(GPUVerifier Verifier)
{
this.LiteralIndexedArrays = new Dictionary <string, HashSet <string> >();
foreach (var v in Verifier.KernelArrayInfo.GetPrivateArrays())
{
this.LiteralIndexedArrays[v.Name] = new HashSet <string>();
}
}
private LoopInvariantGenerator(GPUVerifier verifier, Implementation impl)
{
this.verifier = verifier;
this.impl = impl;
invariantGenerationRules = new List <InvariantGenerationRule>();
invariantGenerationRules.Add(new PowerOfTwoInvariantGenerator(verifier));
}
public AbstractHoudiniTransformation(GPUVerifier verifier)
{
this.verifier = verifier;
this.candidates = verifier.Program.TopLevelDeclarations.OfType <Constant>()
.Where(item => QKeyValue.FindBoolAttribute(item.Attributes, "existential"))
.Select(item => item.Name);
this.counter = 0;
this.existentialFunctions = new List <Declaration>();
}
public ThreadInstantiator(
Expr instantiationExpr, int thread, GPUVerifier verifier, string procName)
{
this.instantiationExpr = instantiationExpr;
this.thread = thread;
this.verifier = verifier;
this.uni = verifier.UniformityAnalyser;
this.procName = procName;
}
public static void EstablishDisabledLoops(GPUVerifier verifier, Implementation impl)
{
foreach (var region in verifier.RootRegion(impl).SubRegions())
{
if (!AccessesGlobalArrayOrUnsafeBarrier(region, verifier))
{
verifier.AddRegionWithLoopInvariantsDisabled(region);
}
}
}
public Expr MaybeBuildPredicate(GPUVerifier verifier, Expr e)
{
var msc = this as ModStrideConstraint;
if (msc != null && !msc.IsBottom())
{
Expr modEqExpr = Expr.Eq(verifier.IntRep.MakeModPow2(e, msc.mod), verifier.IntRep.MakeModPow2(msc.modEq, msc.mod));
return(modEqExpr);
}
return(null);
}
internal static bool AccessesGlobalArrayOrUnsafeBarrier(IRegion region, GPUVerifier verifier)
{
// Heuristic to establish whether to speculate loop invariants for a specific loop
// based on the commands that occur int the loop.
foreach (Cmd c in region.Cmds())
{
if (AccessesGlobalArrayOrUnsafeBarrier(c, verifier))
{
return(true);
}
}
return(false);
}
public override Variable VisitVariable(Variable node)
{
if ((!(node is Constant) && !SkipDualiseVariable(node as Variable)) ||
GPUVerifier.IsThreadLocalIdConstant(node) ||
GPUVerifier.IsGroupIdConstant(node))
{
node.TypedIdent = DualiseTypedIdent(node);
node.Name = node.Name + "$" + id;
return(node);
}
return(base.VisitVariable(node));
}
private static void GenerateCandidateForLoopBounds(GPUVerifier verifier, Implementation impl, IRegion region)
{
HashSet <Variable> loopCounters = new HashSet <Variable>();
HashSet <Variable> modifiedVariables = region.GetModifiedVariables();
// Get the partition variables associated with the header
HashSet <Variable> partitionVars = region.PartitionVariablesOfRegion();
foreach (Variable v in partitionVars)
{
// Find the expression which defines a particular partition variable.
// Visit the expression and select any variable in the mod set of the loop.
// We assume that any variable satisfying these conditions is a loop counter
Expr partitionDefExpr = verifier.VarDefAnalysesRegion[impl].DefOfVariableName(v.Name);
if (partitionDefExpr == null) // multiple definitions or no definition
{
continue;
}
var visitor = new VariablesOccurringInExpressionVisitor();
visitor.Visit(partitionDefExpr);
foreach (Variable variable in visitor.GetVariables())
{
if (modifiedVariables.Contains(variable))
{
loopCounters.Add(variable);
}
}
}
foreach (Variable loopCounter in loopCounters)
{
foreach (Block preheader in region.PreHeaders())
{
foreach (AssignCmd cmd in preheader.Cmds.Where(x => x is AssignCmd).Reverse <Cmd>())
{
var lhss = cmd.Lhss.Where(x => x is SimpleAssignLhs);
foreach (var lhsRhs in lhss.Zip(cmd.Rhss))
{
if (lhsRhs.Item1.DeepAssignedVariable.Name == loopCounter.Name)
{
verifier.AddCandidateInvariant(region, verifier.IntRep.MakeSle(new IdentifierExpr(loopCounter.tok, loopCounter), lhsRhs.Item2), "loopBound");
verifier.AddCandidateInvariant(region, verifier.IntRep.MakeSge(new IdentifierExpr(loopCounter.tok, loopCounter), lhsRhs.Item2), "loopBound");
verifier.AddCandidateInvariant(region, verifier.IntRep.MakeUle(new IdentifierExpr(loopCounter.tok, loopCounter), lhsRhs.Item2), "loopBound");
verifier.AddCandidateInvariant(region, verifier.IntRep.MakeUge(new IdentifierExpr(loopCounter.tok, loopCounter), lhsRhs.Item2), "loopBound");
}
}
}
}
}
}
public static void PreInstrument(GPUVerifier verifier, Implementation impl)
{
foreach (var region in verifier.RootRegion(impl).SubRegions())
{
if (verifier.RegionHasLoopInvariantsDisabled(region))
{
continue;
}
GenerateCandidateForReducedStrengthStrideVariables(verifier, impl, region);
GenerateCandidateForNonNegativeGuardVariables(verifier, impl, region);
GenerateCandidateForNonUniformGuardVariables(verifier, impl, region);
GenerateCandidateForLoopBounds(verifier, impl, region);
GenerateCandidateForEnabledness(verifier, impl, region);
GenerateCandidateForEnablednessWhenAccessingSharedArrays(verifier, impl, region);
}
}
private static void GenerateCandidateForReducedStrengthStrideVariables(GPUVerifier verifier, Implementation impl, IRegion region)
{
var rsa = verifier.ReducedStrengthAnalysesRegion[impl];
var regionId = region.Identifier();
foreach (string iv in rsa.StridedInductionVariables(regionId))
{
var sc = rsa.GetStrideConstraint(iv, regionId);
Variable ivVariable = impl.LocVars.Where(item => item.Name == iv).First();
var ivExpr = new IdentifierExpr(Token.NoToken, ivVariable);
var ivPred = sc.MaybeBuildPredicate(verifier, ivExpr);
if (ivPred != null)
{
verifier.AddCandidateInvariant(region, ivPred, "loopCounterIsStrided");
}
}
}
private static void GenerateCandidateForNonNegativeGuardVariables(GPUVerifier verifier, Implementation impl, IRegion region)
{
HashSet <Variable> partitionVars = region.PartitionVariablesOfHeader();
HashSet <Variable> nonnegVars = new HashSet <Variable>();
var formals = impl.InParams.Select(x => x.Name);
var modset = GetModifiedVariables(region).Select(x => x.Name);
Regex pattern = new Regex(@"\bBV\d*_((SLE)|(SLT)|(SGE)|(SGT))\b");
foreach (var v in partitionVars)
{
var expr = verifier.varDefAnalysesRegion[impl].DefOfVariableName(v.Name);
if (!(expr is NAryExpr))
{
continue;
}
var nary = expr as NAryExpr;
if (!pattern.Match(nary.Fun.FunctionName).Success)
{
continue;
}
var visitor = new VariablesOccurringInExpressionVisitor();
visitor.Visit(nary);
nonnegVars.UnionWith(
visitor.GetVariables().Where(
x => x.Name.StartsWith("$") &&
!formals.Contains(x.Name) &&
modset.Contains(x.Name) &&
x.TypedIdent.Type.IsBv
)
);
}
foreach (var v in nonnegVars)
{
int BVWidth = v.TypedIdent.Type.BvBits;
// REVISIT: really we only want to guess for /integer/ variables.
if (BVWidth >= 8)
{
var inv = verifier.IntRep.MakeSle(verifier.Zero(BVWidth), new IdentifierExpr(v.tok, v));
verifier.AddCandidateInvariant(region, inv, "guardNonNeg");
}
}
}
public override Expr VisitIdentifierExpr(IdentifierExpr node)
{
Debug.Assert(!(node.Decl is Formal));
if (GPUVerifier.IsThreadLocalIdConstant(node.Decl))
{
Debug.Assert(node.Decl.Name.Equals(GPUVerifier._X.Name));
return(InstantiationExprs.Item1.Clone() as Expr);
}
if (node.Decl is Constant ||
verifier.KernelArrayInfo.GetGroupSharedArrays(true).Contains(node.Decl) ||
verifier.KernelArrayInfo.GetGlobalArrays(true).Contains(node.Decl))
{
return(base.VisitIdentifierExpr(node));
}
Console.WriteLine("Expression " + node + " is not valid as part of a barrier invariant: it cannot be instantiated by arbitrary threads.");
Console.WriteLine("Check that it is not a thread local variable, or a thread local (rather than __local or __global) array.");
Console.WriteLine("In particular, if you have a local variable called tid, which you initialise to e.g. get_local_id(0), this will not work:");
Console.WriteLine(" you need to use get_local_id(0) directly.");
Environment.Exit(1);
return(null);
}
private static bool IsDisjunctionOfPredicates(Expr guard)
{
NAryExpr nary = guard as NAryExpr;
if (nary == null || nary.Args.Count() != 2)
{
return(false);
}
BinaryOperator binOp = nary.Fun as BinaryOperator;
if (binOp == null || binOp.Op != BinaryOperator.Opcode.Or)
{
return(false);
}
if (!(nary.Args[0] is IdentifierExpr && nary.Args[1] is IdentifierExpr))
{
return(false);
}
return(GPUVerifier.IsPredicate(Utilities.StripThreadIdentifier(((IdentifierExpr)nary.Args[0]).Name)) &&
GPUVerifier.IsPredicate(Utilities.StripThreadIdentifier(((IdentifierExpr)nary.Args[1]).Name)));
}
private void AddBarrierDivergenceCandidates(HashSet <Variable> localVars, Implementation impl, IRegion region)
{
if (!verifier.ContainsBarrierCall(region) && !GPUVerifyVCGenCommandLineOptions.WarpSync)
{
return;
}
Expr guard = region.Guard();
if (guard != null && verifier.UniformityAnalyser.IsUniform(impl.Name, guard))
{
return;
}
if (IsDisjunctionOfPredicates(guard))
{
string loopPredicate = ((guard as NAryExpr).Args[0] as IdentifierExpr).Name;
loopPredicate = loopPredicate.Substring(0, loopPredicate.IndexOf('$'));
// Int type used here, but it doesn't matter as we will print and then re-parse the program
var uniformEnabledPredicate = Expr.Eq(
new IdentifierExpr(Token.NoToken, new LocalVariable(Token.NoToken, new TypedIdent(Token.NoToken, loopPredicate + "$1", Type.Int))),
new IdentifierExpr(Token.NoToken, new LocalVariable(Token.NoToken, new TypedIdent(Token.NoToken, loopPredicate + "$2", Type.Int))));
verifier.AddCandidateInvariant(region, uniformEnabledPredicate, "loopPredicateEquality");
verifier.AddCandidateInvariant(region, Expr.Imp(verifier.ThreadsInSameGroup(), uniformEnabledPredicate), "loopPredicateEquality");
Dictionary <string, int> assignmentCounts = GetAssignmentCounts(impl);
HashSet <string> alreadyConsidered = new HashSet <string>();
foreach (var v in localVars)
{
string lv = Utilities.StripThreadIdentifier(v.Name);
if (alreadyConsidered.Contains(lv))
{
continue;
}
alreadyConsidered.Add(lv);
if (verifier.UniformityAnalyser.IsUniform(impl.Name, v.Name))
{
continue;
}
if (GPUVerifier.IsPredicate(lv))
{
continue;
}
if (!assignmentCounts.ContainsKey(lv) || assignmentCounts[lv] <= 1)
{
continue;
}
if (!verifier.ContainsNamedVariable(region.GetModifiedVariables(), lv))
{
continue;
}
AddPredicatedEqualityCandidateInvariant(region, loopPredicate, new LocalVariable(Token.NoToken, new TypedIdent(Token.NoToken, lv, Type.Int)));
}
}
}
internal ThreadPairInstantiator(GPUVerifier verifier, Expr InstantiationExpr1, Expr InstantiationExpr2, int Thread)
{
this.verifier = verifier;
this.InstantiationExprs = new Tuple <Expr, Expr>(InstantiationExpr1, InstantiationExpr2);
this.Thread = Thread;
}
private bool InstantiationExprIsThreadId()
{
return((InstantiationExpr is IdentifierExpr) &&
((IdentifierExpr)InstantiationExpr).Decl.Name.Equals(GPUVerifier.MakeThreadId("X", Thread).Name));
}
internal void DualiseKernel()
{
List <Declaration> NewTopLevelDeclarations = new List <Declaration>();
// This loop really does have to be a "for(i ...)" loop. The reason is
// that dualisation may add additional functions to the program, which
// get put into the program's top level declarations and also need to
// be dualised.
var decls = verifier.Program.TopLevelDeclarations.ToList();
for (int i = 0; i < UpdateDeclarationsAndCountTotal(decls); i++)
{
Declaration d = decls[i];
if (d is Axiom)
{
VariableDualiser vd1 = new VariableDualiser(1, null, null);
VariableDualiser vd2 = new VariableDualiser(2, null, null);
Axiom NewAxiom1 = vd1.VisitAxiom(d.Clone() as Axiom);
Axiom NewAxiom2 = vd2.VisitAxiom(d.Clone() as Axiom);
NewTopLevelDeclarations.Add(NewAxiom1);
// Test whether dualisation had any effect by seeing whether the new
// axioms are syntactically indistinguishable. If they are, then there
// is no point adding the second axiom.
if (!NewAxiom1.ToString().Equals(NewAxiom2.ToString()))
{
NewTopLevelDeclarations.Add(NewAxiom2);
}
continue;
}
if (d is Procedure)
{
DualiseProcedure(d as Procedure);
NewTopLevelDeclarations.Add(d);
continue;
}
if (d is Implementation)
{
DualiseImplementation(d as Implementation);
NewTopLevelDeclarations.Add(d);
continue;
}
if (d is Variable && ((d as Variable).IsMutable ||
GPUVerifier.IsThreadLocalIdConstant(d as Variable) ||
(GPUVerifier.IsGroupIdConstant(d as Variable) && !GPUVerifyVCGenCommandLineOptions.OnlyIntraGroupRaceChecking)))
{
var v = d as Variable;
if (v.Name.Contains("_NOT_ACCESSED_") || v.Name.Contains("_ARRAY_OFFSET"))
{
NewTopLevelDeclarations.Add(v);
continue;
}
if (QKeyValue.FindBoolAttribute(v.Attributes, "atomic_usedmap"))
{
NewTopLevelDeclarations.Add(v);
continue;
}
if (verifier.KernelArrayInfo.GetGlobalArrays(true).Contains(v))
{
NewTopLevelDeclarations.Add(v);
continue;
}
if (verifier.KernelArrayInfo.GetGroupSharedArrays(true).Contains(v))
{
if (!GPUVerifyVCGenCommandLineOptions.OnlyIntraGroupRaceChecking)
{
Variable newV = new GlobalVariable(Token.NoToken, new TypedIdent(Token.NoToken,
v.Name, new MapType(Token.NoToken, new List <TypeVariable>(),
new List <Microsoft.Boogie.Type> {
Microsoft.Boogie.Type.GetBvType(1)
},
v.TypedIdent.Type)));
newV.Attributes = v.Attributes;
NewTopLevelDeclarations.Add(newV);
}
else
{
NewTopLevelDeclarations.Add(v);
}
continue;
}
NewTopLevelDeclarations.Add(new VariableDualiser(1, null, null).VisitVariable((Variable)v.Clone()));
if (!QKeyValue.FindBoolAttribute(v.Attributes, "race_checking"))
{
NewTopLevelDeclarations.Add(new VariableDualiser(2, null, null).VisitVariable((Variable)v.Clone()));
}
continue;
}
NewTopLevelDeclarations.Add(d);
}
verifier.Program.TopLevelDeclarations = NewTopLevelDeclarations;
}
public KernelDualiser(GPUVerifier verifier)
{
this.verifier = verifier;
BarrierInvariantDescriptors = new List <BarrierInvariantDescriptor>();
}
private void MakeDual(List <Cmd> cs, Cmd c)
{
if (c is CallCmd)
{
CallCmd Call = c as CallCmd;
if (QKeyValue.FindBoolAttribute(Call.Proc.Attributes, "barrier_invariant"))
{
// There may be a predicate, and there must be an invariant expression and at least one instantiation
Debug.Assert(Call.Ins.Count >= (2 + (verifier.uniformityAnalyser.IsUniform(Call.callee) ? 0 : 1)));
var BIDescriptor = new UnaryBarrierInvariantDescriptor(
verifier.uniformityAnalyser.IsUniform(Call.callee) ? Expr.True : Call.Ins[0],
Expr.Neq(Call.Ins[verifier.uniformityAnalyser.IsUniform(Call.callee) ? 0 : 1],
verifier.Zero(1)),
Call.Attributes,
this, procName, verifier);
for (var i = 1 + (verifier.uniformityAnalyser.IsUniform(Call.callee) ? 0 : 1); i < Call.Ins.Count; i++)
{
BIDescriptor.AddInstantiationExpr(Call.Ins[i]);
}
BarrierInvariantDescriptors.Add(BIDescriptor);
return;
}
if (QKeyValue.FindBoolAttribute(Call.Proc.Attributes, "binary_barrier_invariant"))
{
// There may be a predicate, and there must be an invariant expression and at least one pair of
// instantiation expressions
Debug.Assert(Call.Ins.Count >= (3 + (verifier.uniformityAnalyser.IsUniform(Call.callee) ? 0 : 1)));
var BIDescriptor = new BinaryBarrierInvariantDescriptor(
verifier.uniformityAnalyser.IsUniform(Call.callee) ? Expr.True : Call.Ins[0],
Expr.Neq(Call.Ins[verifier.uniformityAnalyser.IsUniform(Call.callee) ? 0 : 1],
verifier.Zero(1)),
Call.Attributes,
this, procName, verifier);
for (var i = 1 + (verifier.uniformityAnalyser.IsUniform(Call.callee) ? 0 : 1); i < Call.Ins.Count; i += 2)
{
BIDescriptor.AddInstantiationExprPair(Call.Ins[i], Call.Ins[i + 1]);
}
BarrierInvariantDescriptors.Add(BIDescriptor);
return;
}
if (GPUVerifier.IsBarrier(Call.Proc))
{
// Assert barrier invariants
foreach (var BIDescriptor in BarrierInvariantDescriptors)
{
QKeyValue SourceLocationInfo = BIDescriptor.GetSourceLocationInfo();
cs.Add(BIDescriptor.GetAssertCmd());
var vd = new VariableDualiser(1, verifier.uniformityAnalyser, procName);
if (GPUVerifyVCGenCommandLineOptions.BarrierAccessChecks)
{
foreach (Expr AccessExpr in BIDescriptor.GetAccessedExprs())
{
var Assert = new AssertCmd(Token.NoToken, AccessExpr, MakeThreadSpecificAttributes(SourceLocationInfo, 1));
Assert.Attributes = new QKeyValue(Token.NoToken, "barrier_invariant_access_check",
new List <object> {
Expr.True
}, Assert.Attributes);
cs.Add(vd.VisitAssertCmd(Assert));
}
}
}
}
List <Expr> uniformNewIns = new List <Expr>();
List <Expr> nonUniformNewIns = new List <Expr>();
for (int i = 0; i < Call.Ins.Count; i++)
{
if (verifier.uniformityAnalyser.knowsOf(Call.callee) && verifier.uniformityAnalyser.IsUniform(Call.callee, verifier.uniformityAnalyser.GetInParameter(Call.callee, i)))
{
uniformNewIns.Add(Call.Ins[i]);
}
else if (!verifier.OnlyThread2.Contains(Call.callee))
{
nonUniformNewIns.Add(new VariableDualiser(1, verifier.uniformityAnalyser, procName).VisitExpr(Call.Ins[i]));
}
}
for (int i = 0; i < Call.Ins.Count; i++)
{
if (
!(verifier.uniformityAnalyser.knowsOf(Call.callee) && verifier.uniformityAnalyser.IsUniform(Call.callee, verifier.uniformityAnalyser.GetInParameter(Call.callee, i))) &&
!verifier.OnlyThread1.Contains(Call.callee))
{
nonUniformNewIns.Add(new VariableDualiser(2, verifier.uniformityAnalyser, procName).VisitExpr(Call.Ins[i]));
}
}
List <Expr> newIns = uniformNewIns;
newIns.AddRange(nonUniformNewIns);
List <IdentifierExpr> uniformNewOuts = new List <IdentifierExpr>();
List <IdentifierExpr> nonUniformNewOuts = new List <IdentifierExpr>();
for (int i = 0; i < Call.Outs.Count; i++)
{
if (verifier.uniformityAnalyser.knowsOf(Call.callee) && verifier.uniformityAnalyser.IsUniform(Call.callee, verifier.uniformityAnalyser.GetOutParameter(Call.callee, i)))
{
uniformNewOuts.Add(Call.Outs[i]);
}
else
{
nonUniformNewOuts.Add(new VariableDualiser(1, verifier.uniformityAnalyser, procName).VisitIdentifierExpr(Call.Outs[i].Clone() as IdentifierExpr) as IdentifierExpr);
}
}
for (int i = 0; i < Call.Outs.Count; i++)
{
if (!(verifier.uniformityAnalyser.knowsOf(Call.callee) && verifier.uniformityAnalyser.IsUniform(Call.callee, verifier.uniformityAnalyser.GetOutParameter(Call.callee, i))))
{
nonUniformNewOuts.Add(new VariableDualiser(2, verifier.uniformityAnalyser, procName).VisitIdentifierExpr(Call.Outs[i].Clone() as IdentifierExpr) as IdentifierExpr);
}
}
List <IdentifierExpr> newOuts = uniformNewOuts;
newOuts.AddRange(nonUniformNewOuts);
CallCmd NewCallCmd = new CallCmd(Call.tok, Call.callee, newIns, newOuts);
NewCallCmd.Proc = Call.Proc;
NewCallCmd.Attributes = Call.Attributes;
if (NewCallCmd.callee.StartsWith("_LOG_ATOMIC"))
{
QKeyValue curr = NewCallCmd.Attributes;
if (curr.Key.StartsWith("arg"))
{
NewCallCmd.Attributes = new QKeyValue(Token.NoToken, curr.Key, new List <object>(new object[] { Dualise(curr.Params[0] as Expr, 1) }), curr.Next);
}
for (curr = NewCallCmd.Attributes; curr.Next != null; curr = curr.Next)
{
if (curr.Next.Key.StartsWith("arg"))
{
curr.Next = new QKeyValue(Token.NoToken, curr.Next.Key, new List <object>(new object[] { Dualise(curr.Next.Params[0] as Expr, 1) }), curr.Next.Next);
}
}
}
else if (NewCallCmd.callee.StartsWith("_CHECK_ATOMIC"))
{
QKeyValue curr = NewCallCmd.Attributes;
if (curr.Key.StartsWith("arg"))
{
NewCallCmd.Attributes = new QKeyValue(Token.NoToken, curr.Key, new List <object>(new object[] { Dualise(curr.Params[0] as Expr, 2) }), curr.Next);
}
for (curr = NewCallCmd.Attributes; curr.Next != null; curr = curr.Next)
{
if (curr.Next.Key.StartsWith("arg"))
{
curr.Next = new QKeyValue(Token.NoToken, curr.Next.Key, new List <object>(new object[] { Dualise(curr.Next.Params[0] as Expr, 2) }), curr.Next.Next);
}
}
}
cs.Add(NewCallCmd);
if (GPUVerifier.IsBarrier(Call.Proc))
{
foreach (var BIDescriptor in BarrierInvariantDescriptors)
{
foreach (var Instantiation in BIDescriptor.GetInstantiationCmds())
{
cs.Add(Instantiation);
}
}
BarrierInvariantDescriptors.Clear();
}
}
else if (c is AssignCmd)
{
AssignCmd assign = c as AssignCmd;
var vd1 = new VariableDualiser(1, verifier.uniformityAnalyser, procName);
var vd2 = new VariableDualiser(2, verifier.uniformityAnalyser, procName);
List <AssignLhs> lhss1 = new List <AssignLhs>();
List <AssignLhs> lhss2 = new List <AssignLhs>();
List <Expr> rhss1 = new List <Expr>();
List <Expr> rhss2 = new List <Expr>();
foreach (var pair in assign.Lhss.Zip(assign.Rhss))
{
if (pair.Item1 is SimpleAssignLhs &&
verifier.uniformityAnalyser.IsUniform(procName,
(pair.Item1 as SimpleAssignLhs).AssignedVariable.Name))
{
lhss1.Add(pair.Item1);
rhss1.Add(pair.Item2);
}
else
{
lhss1.Add(vd1.Visit(pair.Item1.Clone() as AssignLhs) as AssignLhs);
lhss2.Add(vd2.Visit(pair.Item1.Clone() as AssignLhs) as AssignLhs);
rhss1.Add(vd1.VisitExpr(pair.Item2.Clone() as Expr));
rhss2.Add(vd2.VisitExpr(pair.Item2.Clone() as Expr));
}
}
Debug.Assert(lhss1.Count > 0);
cs.Add(new AssignCmd(Token.NoToken, lhss1, rhss1));
if (lhss2.Count > 0)
{
cs.Add(new AssignCmd(Token.NoToken, lhss2, rhss2));
}
}
else if (c is HavocCmd)
{
HavocCmd havoc = c as HavocCmd;
Debug.Assert(havoc.Vars.Count() == 1);
HavocCmd newHavoc;
newHavoc = new HavocCmd(havoc.tok, new List <IdentifierExpr>(new IdentifierExpr[] {
(IdentifierExpr)(new VariableDualiser(1, verifier.uniformityAnalyser, procName).VisitIdentifierExpr(havoc.Vars[0].Clone() as IdentifierExpr)),
(IdentifierExpr)(new VariableDualiser(2, verifier.uniformityAnalyser, procName).VisitIdentifierExpr(havoc.Vars[0].Clone() as IdentifierExpr))
}));
cs.Add(newHavoc);
}
else if (c is AssertCmd)
{
AssertCmd a = c as AssertCmd;
if (QKeyValue.FindBoolAttribute(a.Attributes, "sourceloc") ||
QKeyValue.FindBoolAttribute(a.Attributes, "block_sourceloc") ||
QKeyValue.FindBoolAttribute(a.Attributes, "array_bounds"))
{
// This is just a location marker, so we do not dualise it
cs.Add(new AssertCmd(Token.NoToken, new VariableDualiser(1, verifier.uniformityAnalyser, procName).VisitExpr(a.Expr.Clone() as Expr),
(QKeyValue)a.Attributes.Clone()));
}
else
{
var isUniform = verifier.uniformityAnalyser.IsUniform(procName, a.Expr);
cs.Add(MakeThreadSpecificAssert(a, 1));
if (!GPUVerifyVCGenCommandLineOptions.AsymmetricAsserts && !ContainsAsymmetricExpression(a.Expr) && !isUniform)
{
cs.Add(MakeThreadSpecificAssert(a, 2));
}
}
}
else if (c is AssumeCmd)
{
AssumeCmd ass = c as AssumeCmd;
if (QKeyValue.FindStringAttribute(ass.Attributes, "captureState") != null)
{
cs.Add(c);
}
else if (QKeyValue.FindBoolAttribute(ass.Attributes, "backedge"))
{
AssumeCmd newAss = new AssumeCmd(c.tok, Expr.Or(new VariableDualiser(1, verifier.uniformityAnalyser, procName).VisitExpr(ass.Expr.Clone() as Expr),
new VariableDualiser(2, verifier.uniformityAnalyser, procName).VisitExpr(ass.Expr.Clone() as Expr)));
newAss.Attributes = ass.Attributes;
cs.Add(newAss);
}
else if (QKeyValue.FindBoolAttribute(ass.Attributes, "atomic_refinement"))
{
// Generate the following:
// havoc v$1, v$2;
// assume !_USED[offset$1][v$1];
// _USED[offset$1][v$1] := true;
// assume !_USED[offset$2][v$2];
// _USED[offset$2][v$2] := true;
Expr variable = QKeyValue.FindExprAttribute(ass.Attributes, "variable");
Expr offset = QKeyValue.FindExprAttribute(ass.Attributes, "offset");
List <Expr> offsets = (new int[] { 1, 2 }).Select(x => new VariableDualiser(x, verifier.uniformityAnalyser, procName).VisitExpr(offset.Clone() as Expr)).ToList();
List <Expr> vars = (new int[] { 1, 2 }).Select(x => new VariableDualiser(x, verifier.uniformityAnalyser, procName).VisitExpr(variable.Clone() as Expr)).ToList();
IdentifierExpr arrayref = new IdentifierExpr(Token.NoToken, verifier.FindOrCreateUsedMap(QKeyValue.FindStringAttribute(ass.Attributes, "arrayref"), vars[0].Type));
foreach (int i in (new int[] { 0, 1 }))
{
AssumeCmd newAss = new AssumeCmd(c.tok, Expr.Not(new NAryExpr(Token.NoToken, new MapSelect(Token.NoToken, 1),
new List <Expr> {
new NAryExpr(Token.NoToken, new MapSelect(Token.NoToken, 1),
new List <Expr> {
arrayref, offsets[i]
}),
vars[i]
})));
cs.Add(newAss);
var lhs = new MapAssignLhs(Token.NoToken, new MapAssignLhs(Token.NoToken, new SimpleAssignLhs(Token.NoToken, arrayref),
new List <Expr> {
offsets[i]
}), new List <Expr> {
vars[i]
});
AssignCmd assign = new AssignCmd(c.tok,
new List <AssignLhs> {
lhs
},
new List <Expr> {
Expr.True
});
cs.Add(assign);
}
}
else
{
var isUniform = verifier.uniformityAnalyser.IsUniform(procName, ass.Expr);
AssumeCmd newAss = new AssumeCmd(c.tok, new VariableDualiser(1, verifier.uniformityAnalyser, procName).VisitExpr(ass.Expr.Clone() as Expr));
if (!ContainsAsymmetricExpression(ass.Expr) && !isUniform)
{
newAss.Expr = Expr.And(newAss.Expr, new VariableDualiser(2, verifier.uniformityAnalyser, procName).VisitExpr(ass.Expr.Clone() as Expr));
}
newAss.Attributes = ass.Attributes;
cs.Add(newAss);
}
}
else
{
Debug.Assert(false);
}
}
private static bool AccessesGlobalArrayOrUnsafeBarrier(Cmd c, GPUVerifier verifier)
{
var stateToCheck = verifier.KernelArrayInfo;
if (c is CallCmd)
{
// Speculate invariants if we see atomics, async_work_group_copy, and
// wait_group_events, which relate to race checking
CallCmd call = c as CallCmd;
if (QKeyValue.FindBoolAttribute(call.Attributes, "atomic"))
{
return(true);
}
if (QKeyValue.FindBoolAttribute(call.Attributes, "async_work_group_copy"))
{
return(true);
}
if (QKeyValue.FindBoolAttribute(call.Attributes, "wait_group_events"))
{
return(true);
}
// Speculate invariants if we see an unsafe barrier,
// which we need to check for barrier divergence
if (GPUVerifier.IsBarrier(call.Proc) &&
!QKeyValue.FindBoolAttribute(call.Proc.Attributes, "safe_barrier"))
{
return(true);
}
// Speculate invariants if we see a call to a procedure that has a non-local array
// or constant array in its modset
List <Variable> vars = new List <Variable>();
call.AddAssignedVariables(vars);
foreach (Variable v in vars)
{
if (stateToCheck.GetGlobalAndGroupSharedArrays(false).Contains(v))
{
return(true);
}
if (stateToCheck.GetConstantArrays().Contains(v))
{
return(true);
}
}
}
// Speculate invariants if race instrumentation or a constant write
// instrumentation will occur
if (c is AssignCmd)
{
AssignCmd assign = c as AssignCmd;
ReadCollector rc = new ReadCollector(stateToCheck);
foreach (var rhs in assign.Rhss)
{
rc.Visit(rhs);
}
foreach (var access in rc.NonPrivateAccesses)
{
// Ignore disabled arrays
if (stateToCheck.GetGlobalAndGroupSharedArrays(false).Contains(access.V))
{
// Ignore read-only arrays (whether or not they are disabled)
if (!stateToCheck.GetReadOnlyGlobalAndGroupSharedArrays(true).Contains(access.V))
{
return(true);
}
}
}
foreach (var lhsRhs in assign.Lhss.Zip(assign.Rhss))
{
WriteCollector wc = new WriteCollector(stateToCheck);
wc.Visit(lhsRhs.Item1);
if (wc.FoundNonPrivateWrite())
{
// Ignore disabled arrays
if (stateToCheck.GetGlobalAndGroupSharedArrays(false).Contains(wc.GetAccess().V))
{
return(true);
}
}
}
foreach (var lhsRhs in assign.Lhss.Zip(assign.Rhss))
{
ConstantWriteCollector cwc = new ConstantWriteCollector(stateToCheck);
cwc.Visit(lhsRhs.Item1);
if (cwc.FoundWrite())
{
// Ignore disabled arrays
if (stateToCheck.GetGlobalAndGroupSharedArrays(false).Contains(cwc.GetAccess().V))
{
return(true);
}
}
}
}
// Speculate invariants if we see an assert that is not a sourceloc or
// block_sourceloc assert; such asserts is likely user supplied.
if (c is AssertCmd)
{
AssertCmd assertion = c as AssertCmd;
if (!QKeyValue.FindBoolAttribute(assertion.Attributes, "sourceloc") &&
!QKeyValue.FindBoolAttribute(assertion.Attributes, "block_sourceloc") &&
!assertion.Expr.Equals(Expr.True))
{
return(true);
}
}
// Speculate invariants if we see an assume that is not a partition; such
// an assume is likely user supplied.
if (c is AssumeCmd)
{
AssumeCmd assumption = c as AssumeCmd;
if (!QKeyValue.FindBoolAttribute(assumption.Attributes, "partition"))
{
return(true);
}
}
return(false);
}
private static void GenerateCandidateForEnablednessWhenAccessingSharedArrays(GPUVerifier verifier, Implementation impl, IRegion region)
{
Block header = region.Header();
if (verifier.UniformityAnalyser.IsUniform(impl.Name, header))
{
return;
}
var cfg = Program.GraphFromImpl(impl);
Dictionary <Block, HashSet <Block> > controlDependence = cfg.ControlDependence();
controlDependence.TransitiveClosure();
cfg.ComputeLoops();
List <Expr> guards = new List <Expr>();
foreach (var b in controlDependence.Keys.Where(item => controlDependence[item].Contains(region.Header())))
{
foreach (var succ in cfg.Successors(b).Where(item => cfg.DominatorMap.DominatedBy(header, item)))
{
var guard = MaybeExtractGuard(verifier, impl, succ);
if (guard != null)
{
guards.Add(guard);
break;
}
}
}
if (guards.Count == 0)
{
return;
}
IEnumerable <Variable> readVariables;
IEnumerable <Variable> writtenVariables;
GetReadAndWrittenVariables(region, out readVariables, out writtenVariables);
foreach (var v in readVariables.Where(item => verifier.KernelArrayInfo.GetGlobalAndGroupSharedArrays(false).Contains(item) &&
!verifier.KernelArrayInfo.GetReadOnlyGlobalAndGroupSharedArrays(true).Contains(item)))
{
foreach (var g in guards)
{
verifier.AddCandidateInvariant(
region,
Expr.Imp(Expr.Ident(verifier.FindOrCreateAccessHasOccurredVariable(v.Name, AccessType.READ)), g),
"accessOnlyIfEnabledInEnclosingScopes",
"do_not_predicate");
}
}
foreach (var v in writtenVariables.Where(item => verifier.KernelArrayInfo.GetGlobalAndGroupSharedArrays(false).Contains(item)))
{
foreach (var g in guards)
{
verifier.AddCandidateInvariant(
region,
Expr.Imp(Expr.Ident(verifier.FindOrCreateAccessHasOccurredVariable(v.Name, AccessType.WRITE)), g),
"accessOnlyIfEnabledInEnclosingScopes",
"do_not_predicate");
}
}
}
private static void GenerateCandidateForNonUniformGuardVariables(GPUVerifier verifier, Implementation impl, IRegion region)
{
if (!verifier.ContainsBarrierCall(region) && !GPUVerifyVCGenCommandLineOptions.WarpSync)
{
return;
}
HashSet <Variable> partitionVars = region.PartitionVariablesOfHeader();
HashSet <Variable> guardVars = new HashSet <Variable>();
var formals = impl.InParams.Select(x => x.Name);
var modset = region.GetModifiedVariables().Select(x => x.Name);
foreach (var v in partitionVars)
{
Expr expr = verifier.VarDefAnalysesRegion[impl].DefOfVariableName(v.Name);
if (expr == null)
{
continue;
}
var visitor = new VariablesOccurringInExpressionVisitor();
visitor.Visit(expr);
guardVars.UnionWith(
visitor.GetVariables().Where(x => x.Name.StartsWith("$") &&
!formals.Contains(x.Name) && modset.Contains(x.Name) &&
!verifier.UniformityAnalyser.IsUniform(impl.Name, x.Name) &&
x.TypedIdent.Type.IsBv && (x.TypedIdent.Type.BvBits % 8 == 0)));
}
List <AssignCmd> assignments = new List <AssignCmd>();
foreach (Block b in region.PreHeaders())
{
foreach (AssignCmd c in b.Cmds.Where(x => x is AssignCmd))
{
assignments.Add(c);
}
}
foreach (var v in guardVars)
{
foreach (AssignCmd c in assignments)
{
foreach (var a in c.Lhss.Zip(c.Rhss))
{
var lhs = a.Item1;
var rhs = a.Item2;
if (!(lhs is SimpleAssignLhs))
{
continue;
}
var sLhs = (SimpleAssignLhs)lhs;
var theVar = sLhs.DeepAssignedVariable;
if (theVar.Name == v.Name)
{
var sub = verifier.IntRep.MakeSub(new IdentifierExpr(Token.NoToken, v), rhs as Expr);
List <Expr> args = new List <Expr>();
args.Add(sub);
Function otherbv = verifier.FindOrCreateOther(sub.Type);
var inv = Expr.Eq(sub, new NAryExpr(Token.NoToken, new FunctionCall(otherbv), args));
verifier.AddCandidateInvariant(region, inv, "guardMinusInitialIsUniform");
var groupInv = Expr.Imp(verifier.ThreadsInSameGroup(), inv);
verifier.AddCandidateInvariant(region, groupInv, "guardMinusInitialIsUniform");
}
}
}
}
}
public static void PostInstrument(GPUVerifier verifier, Implementation impl)
{
new LoopInvariantGenerator(verifier, impl).PostInstrument();
}
