private static Node ForTree(BasicBlock BB, Instruction Instruction, CpState State, Node Node, bool TopIsDef, ref bool Changed, Dictionary <Location, LocDef> InsertedDict) { if (Node == null) { return(null); } if (!TopIsDef && Node.NodeType == NodeType.Location) { LocationNode LocationNode = (LocationNode)Node; LocDef LocDef; if (LocationNode.Location.LocationType == LocationType.CilStack && State.Defs.TryGetValue(LocationNode.Location, out LocDef)) { InsertedDict[LocationNode.Location] = LocDef; Changed = true; Node = LocDef.Node; } } else { for (int i = 0; i < Node.SubNodes.Count; i++) { Node SubNode = Node.SubNodes[i]; SubNode = ForTree(BB, Instruction, State, SubNode, (SubNode.NodeType == NodeType.AddressOf) ? true : false, ref Changed, InsertedDict); Node.SubNodes[i] = SubNode; } } return(Node); }
private static HlGraphEntry ConstructHlGraphEntry(MethodInfo Method, int GidParamCount, bool IsKernel, string NameTemplate) { TextWriter writer = System.Console.Out; string MethodName = string.Format(NameTemplate, HlGraphSequenceNumber++); HighLevel.HlGraph HLgraph = new HighLevel.HlGraph(Method, MethodName); HLgraph.IsKernel = IsKernel; HLgraph.ValueTypeMap = m_ValueTypeMap; if (!IsKernel && Method.DeclaringType.IsValueType && ((Method.CallingConvention & CallingConventions.HasThis) != 0)) { System.Diagnostics.Debug.Assert(HLgraph.Arguments.Count > 0); System.Diagnostics.Debug.Assert(HLgraph.Arguments[0].DataType.IsByRef && HLgraph.Arguments[0].DataType.GetElementType() == Method.DeclaringType); HLgraph.KeepThis = true; } if (DumpCode > 3) { WriteCode(HLgraph, writer); } // Optimize it (just some copy propagation and dead assignment elimination to get rid of // CIL stack accesses) HLgraph.Optimize(); if (DumpCode > 4) { WriteCode(HLgraph, writer); } // Convert all expression trees into something OpenCL can understand HLgraph.ConvertForOpenCl(subGraphCache); System.Diagnostics.Debug.Assert(HLgraph.KeepThis || !HLgraph.HasThisParameter); // Change the real first arguments (the "int"s of the Action<> method) to local variables // which get their value from OpenCL's built-in get_global_id() routine. // NOTE: ConvertArgumentToLocal removes the specified argument, so both calls need to specify // an ArgumentId of zero!!! List <HighLevel.LocalVariableLocation> IdLocation = new List <HighLevel.LocalVariableLocation>(); for (int i = 0; i < GidParamCount; i++) { IdLocation.Add(HLgraph.ConvertArgumentToLocal(0)); } // Add fromInclusive and toExclusive as additional parameters List <HighLevel.ArgumentLocation> StartIdLocation = new List <HighLevel.ArgumentLocation>(); List <HighLevel.ArgumentLocation> EndIdLocation = new List <HighLevel.ArgumentLocation>(); for (int i = 0; i < GidParamCount; i++) { StartIdLocation.Add(HLgraph.InsertArgument(i * 2 + 0, "fromInclusive" + i, typeof(int), false)); EndIdLocation.Add(HLgraph.InsertArgument(i * 2 + 1, "toExclusive" + i, typeof(int), false)); } // "i0 = get_global_id(0) + fromInclusive0;" for (int i = 0; i < GidParamCount; i++) { HLgraph.CanonicalStartBlock.Instructions.Insert(i, new HighLevel.AssignmentInstruction( new HighLevel.LocationNode(IdLocation[i]), new HighLevel.AddNode( new HighLevel.CallNode(typeof(OpenClFunctions).GetMethod("get_global_id", new Type[] { typeof(uint) }), new HighLevel.IntegerConstantNode(i)), new HighLevel.LocationNode(StartIdLocation[i]) ) ) ); } // "if (i0 >= toExclusive0) return;" if (GidParamCount > 0) { HighLevel.BasicBlock ReturnBlock = null; foreach (HighLevel.BasicBlock BB in HLgraph.BasicBlocks) { if (BB.Instructions.Count == 1 && BB.Instructions[0].InstructionType == HighLevel.InstructionType.Return) { ReturnBlock = BB; break; } } if (ReturnBlock == null) { ReturnBlock = new HighLevel.BasicBlock("CANONICAL_RETURN_BLOCK"); ReturnBlock.Instructions.Add(new HighLevel.ReturnInstruction(null)); HLgraph.BasicBlocks.Add(ReturnBlock); } ReturnBlock.LabelNameUsed = true; for (int i = 0; i < GidParamCount; i++) { HLgraph.CanonicalStartBlock.Instructions.Insert(GidParamCount + i, new HighLevel.ConditionalBranchInstruction( new HighLevel.GreaterEqualsNode( new HighLevel.LocationNode(IdLocation[i]), new HighLevel.LocationNode(EndIdLocation[i]) ), ReturnBlock ) ); } } // Create the argument to pass the random seed, if necessary if (!object.ReferenceEquals(HLgraph.RandomStateLocation, null) && HLgraph.RandomStateLocation.LocationType == HighLevel.LocationType.LocalVariable) { // This can only happen for kernels. All nested routines get a pointer to // the kernel's rnd_state instead System.Diagnostics.Debug.Assert(HLgraph.IsKernel); System.Diagnostics.Debug.Assert(object.ReferenceEquals(HLgraph.RandomSeedArgument, null)); HLgraph.RandomSeedArgument = HLgraph.CreateArgument("rnd_seed", typeof(uint), false); if (GidParamCount > 0) { HighLevel.Node LocalSeed = null; for (int i = 0; i < GidParamCount; i++) { if (LocalSeed == null) { LocalSeed = new HighLevel.LocationNode(IdLocation[i]); } else { LocalSeed = new HighLevel.AddNode( new HighLevel.LocationNode(IdLocation[i]), new HighLevel.MulNode( LocalSeed, new HighLevel.IntegerConstantNode(0x10000) /* TODO: what is a good factor here ??? */ ) ); } } HLgraph.CanonicalStartBlock.Instructions.Add(new HighLevel.AssignmentInstruction( new HighLevel.NamedFieldNode(new HighLevel.LocationNode(HLgraph.RandomStateLocation), "x", typeof(uint)), new HighLevel.AddNode( LocalSeed, new HighLevel.IntegerConstantNode(1) ) ) ); } else { HLgraph.CanonicalStartBlock.Instructions.Add(new HighLevel.AssignmentInstruction( new HighLevel.NamedFieldNode(new HighLevel.LocationNode(HLgraph.RandomStateLocation), "x", typeof(uint)), new HighLevel.IntegerConstantNode(1) ) ); } HLgraph.CanonicalStartBlock.Instructions.Add(new HighLevel.AssignmentInstruction( new HighLevel.NamedFieldNode(new HighLevel.LocationNode(HLgraph.RandomStateLocation), "y", typeof(uint)), new HighLevel.LocationNode(HLgraph.RandomSeedArgument) )); // Perform TWO warmup rounds, so our not-so-random start states // get a chance to really inflict changes to all 32 bits of // generated random numbers. HLgraph.CanonicalStartBlock.Instructions.Add(new HighLevel.AssignmentInstruction( null, new HighLevel.CallNode( typeof(OpenClFunctions).GetMethod("rnd"), new HighLevel.AddressOfNode( new HighLevel.LocationNode(HLgraph.RandomStateLocation) ) ) ) ); HLgraph.CanonicalStartBlock.Instructions.Add(new HighLevel.AssignmentInstruction( null, new HighLevel.CallNode( typeof(OpenClFunctions).GetMethod("rnd"), new HighLevel.AddressOfNode( new HighLevel.LocationNode(HLgraph.RandomStateLocation) ) ) ) ); } if (DumpCode > 5) { WriteCode(HLgraph, writer); } // Update location usage information HLgraph.AnalyzeLocationUsage(); // Finally, add the graph to the cache HlGraphEntry CacheEntry = new HlGraphEntry(HLgraph, StartIdLocation, EndIdLocation); return(CacheEntry); }