public override Expr VisitIdentifierExpr(IdentifierExpr node)
{
if (node.Decl is Formal)
{
return new IdentifierExpr(
node.tok, new Formal(node.tok, DualiseTypedIdent(node.Decl), ((Formal)node.Decl).InComing));
}
if (!(node.Decl is Constant) && !SkipDualiseVariable(node.Decl as Variable))
{
return new IdentifierExpr(node.tok, new LocalVariable(node.tok, DualiseTypedIdent(node.Decl)));
}
if (verifier.IsThreadLocalIdConstant(node.Decl))
{
return new IdentifierExpr(node.tok, new Constant(node.tok, DualiseTypedIdent(node.Decl)));
}
if (verifier.IsGroupIdConstant(node.Decl))
{
if (GPUVerifyVCGenCommandLineOptions.OnlyIntraGroupRaceChecking)
{
return node;
}
return new IdentifierExpr(node.tok, new Constant(node.tok, DualiseTypedIdent(node.Decl)));
}
return node;
}
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(InstantiationExpr.Clone() as Expr);
}
if (node.Decl is Constant ||
QKeyValue.FindBoolAttribute(node.Decl.Attributes, "global") ||
QKeyValue.FindBoolAttribute(node.Decl.Attributes, "group_shared") ||
(Uni != null && Uni.IsUniform(ProcName, node.Decl.Name)))
{
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);
}
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));
}
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);
}
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;
}