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; }