public HlGraphEntry(HighLevel.HlGraph HlGraph, List <HighLevel.ArgumentLocation> fromInclusiveLocation, List <HighLevel.ArgumentLocation> toExclusiveLocation)
{
m_HlGraph = HlGraph;
m_fromInclusiveLocation = fromInclusiveLocation;
m_toExclusiveLocation = toExclusiveLocation;
}
public static bool Do(HlGraph Graph)
{
DataFlow <LocationList> .Result DfResult = DoDF(Graph);
List <Location> DeadDefinitionList = new List <Location>();
bool Changed = false;
foreach (BasicBlock BasicBlock in Graph.BasicBlocks)
{
DeadDefinitionList.Clear();
LocationList List = DfResult.ExitState[BasicBlock];
if (List == null)
{
System.Diagnostics.Debug.Assert(BasicBlock == Graph.CanonicalExitBlock);
}
else
{
DeadDefinitionList.AddRange(List);
}
Changed |= ForBlock(BasicBlock, ref DeadDefinitionList);
}
return(Changed);
}
internal static HlGraphEntry ConstructRelatedHlGraphEntry(MethodInfo Method, HighLevel.HlGraph ParentGraph, HlGraphCache RelatedGraphCache)
{
HlGraphEntry CacheEntry = ConstructHlGraphEntry(Method, 0, false, "Cil2OpenCL_Sub_Seq{0}");
RelatedGraphCache.SetValue(IntPtr.Zero, Method, CacheEntry);
ParentGraph.RelatedGraphs[Method] = CacheEntry;
GenerateOpenCLSource(CacheEntry);
return(CacheEntry);
}
private static string getOpenCLSource(HighLevel.HlGraph HLgraph)
{
using (StringWriter Srcwriter = new StringWriter())
{
OpenCLInterop.WriteOpenCL(HLgraph, Srcwriter);
string OpenClSource = Srcwriter.ToString();
if (DumpCode > 2)
{
System.Console.WriteLine(OpenClSource);
}
return(OpenClSource);
}
}
public InvokeContext(HighLevel.HlGraph HLgraph, int RandomSeed)
{
m_ValueTypeMap = HLgraph.ValueTypeMap;
m_Arguments = new List <InvokeArgument>(HLgraph.Arguments.Count);
for (int i = 0; i < HLgraph.Arguments.Count; i++)
{
System.Diagnostics.Debug.Assert(HLgraph.Arguments[i].Index == i);
m_Arguments.Add(null);
}
if (!object.ReferenceEquals(HLgraph.RandomSeedArgument, null))
{
PutArgument(HLgraph.RandomSeedArgument, RandomSeed);
}
}
private static void DoInvoke(int[] WorkSize, object Target, HlGraphEntry CacheEntry, InvokeContext ctx, OpenCLNet.Device device)
{
HighLevel.HlGraph HLgraph = CacheEntry.HlGraph;
foreach (KeyValuePair <FieldInfo, HighLevel.ArgumentLocation> Entry in HLgraph.StaticFieldMap)
{
ctx.PutArgument(Entry.Value, Entry.Key.GetValue(null));
}
SetArguments(ctx, Target, HLgraph.RootPathEntry);
/*
* foreach (KeyValuePair<FieldInfo, HighLevel.ArgumentLocation> Entry in HLgraph.ThisFieldMap)
* {
* ctx.PutArgument(Entry.Value, Entry.Key.GetValue(Target));
* }
* foreach (KeyValuePair<FieldInfo, Dictionary<FieldInfo, HighLevel.ArgumentLocation>> Entry in HLgraph.OuterThisFieldMap) {
* object RealThis = Entry.Key.GetValue(Target);
* foreach (KeyValuePair<FieldInfo, HighLevel.ArgumentLocation> SubEntry in Entry.Value) {
* ctx.PutArgument(SubEntry.Value, SubEntry.Key.GetValue(RealThis));
* }
* }*/
foreach (KeyValuePair <HighLevel.ArgumentLocation, HighLevel.ArrayInfo> Entry in HLgraph.MultiDimensionalArrayInfo)
{
System.Diagnostics.Debug.Assert(Entry.Key.Index >= 0 && Entry.Key.Index < ctx.Arguments.Count);
InvokeArgument BaseArrayArg = ctx.Arguments[Entry.Key.Index];
System.Diagnostics.Debug.Assert(BaseArrayArg != null && BaseArrayArg.Value != null && BaseArrayArg.Value.GetType() == Entry.Key.DataType);
System.Diagnostics.Debug.Assert(Entry.Key.DataType.IsArray && Entry.Key.DataType.GetArrayRank() == Entry.Value.DimensionCount);
System.Diagnostics.Debug.Assert(BaseArrayArg.Value is Array);
Array BaseArray = (System.Array)BaseArrayArg.Value;
long BaseFactor = 1;
for (int Dimension = 1; Dimension < Entry.Value.DimensionCount; Dimension++)
{
int ThisDimensionLength = BaseArray.GetLength(Entry.Value.DimensionCount - 1 - (Dimension - 1));
BaseFactor *= ThisDimensionLength;
ctx.PutArgument(Entry.Value.ScaleArgument[Dimension], (int)BaseFactor);
}
}
ctx.Complete();
OpenCLInterop.CallOpenCLNet(WorkSize, CacheEntry, ctx, HLgraph, device);
}
/// <summary>
/// Perform data-flow analysis to get dead definitions for each basic block.
/// </summary>
/// <param name="Graph">HlGraph.</param>
/// <returns>Data flow results.</returns>
/// <remarks>Dead definitions is an all-paths reverse data flow problem described by the
/// following equations:
///
/// <code>
/// InitFn(n) := ( (n == CanonicalExit) ? ALL_STACK | ALL_LOCAL | ALL_ARG : TOP, TOP )
/// NodeFn(n, l) := l + n.DAD - n.UEU
/// EdgeFn(s, t, l) := l
/// MergeFn(l, r) := (l & r)
/// </code>
/// </remarks>
public static DataFlow <LocationList> .Result DoDF(HlGraph Graph)
{
DataFlow <LocationList> .Result DfResult = DataFlow <LocationList> .Reverse(
Graph.BasicBlocks,
delegate(BasicBlock BasicBlock, out LocationList Entry, out LocationList Exit)
{
if (Graph.CanonicalExitBlock == BasicBlock)
{
Entry = new LocationList();
for (int i = 0; i < Graph.m_MaxStack; i++)
{
Entry.Add(new StackLocation(i));
}
foreach (LocalVariableLocation LocVar in Graph.LocalVariables)
{
Entry.Add(LocVar);
}
foreach (ArgumentLocation Arg in Graph.Arguments)
{
Entry.Add(Arg);
}
Exit = null;
}
else
{
Entry = Exit = null;
}
},
delegate(BasicBlock BasicBlock, LocationList DeadDefinitionList)
{
return(BeforeInstruction(BasicBlock, 0, DeadDefinitionList));
},
delegate(BasicBlock Source, BasicBlock Target, LocationList DeadDefinitionList)
{
return(DeadDefinitionList);
},
LocationList.MergeInlineAnd
);
return(DfResult);
}
internal static void WriteOpenCL(HighLevel.HlGraph HLgraph, Type StructType, TextWriter writer)
{
if (StructType == null)
{
throw new ArgumentNullException("StructType");
}
else if (!StructType.IsValueType)
{
throw new ArgumentException(string.Format("Unable to generate OpenCL code for non-ValueType '{0}'", StructType.FullName));
}
writer.WriteLine();
writer.WriteLine("// OpenCL structure definition for type '{0}'", StructType.FullName);
writer.WriteLine("typedef {0} {{", GetOpenClType(HLgraph, StructType));
FieldInfo[] Fields = StructType.GetFields();
foreach (FieldInfo Field in Fields)
{
writer.WriteLine("\t{0} {1};", GetOpenClType(HLgraph, Field.FieldType), Field.Name);
}
writer.WriteLine("}} t_{0};", HLgraph.ValueTypeMap[StructType]);
}
public abstract List <HighLevel.Instruction> GetHighLevel(HighLevel.HlGraph Context);
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);
}
public override string ToString()
{
StringBuilder String = new StringBuilder();
int i = 0;
if (!IsStaticCall)
{
if (SubNodes.Count == 0)
{
String.Append("(???).");
}
else
{
String.Append(SubNodes[0].ToString());
String.Append(".");
}
i++;
}
string Name = object.ReferenceEquals(HlGraph, null) ? OpenClAliasAttribute.Get(MethodInfo) : HlGraph.GetOpenClFunctionName(MethodInfo);
if (Name == null)
{
Name = MethodInfo.Name;
}
String.Append(Name);
String.Append("(");
bool IsFirst = true;
for (; i < SubNodes.Count; i++)
{
if (IsFirst)
{
IsFirst = false;
}
else
{
String.Append(", ");
}
String.Append(SubNodes[i].ToString());
}
String.Append(")");
return(String.ToString());
}
private static void GenerateOpenCLSource(HlGraphEntry CacheEntry)
{
// Non-kernel methods include just their own code, but kernel methods include everything required
if (!CacheEntry.HlGraph.IsKernel || CacheEntry.HlGraph.RelatedGraphs.Count == 0)
{
CacheEntry.Source = getOpenCLSource(CacheEntry.HlGraph);
return;
}
// No recursion allowed, so we can get away with a topological sort of all involved functions
// with no prototypes beforehand.
// The following code has been adapted from
// http://www.logarithmic.net/pfh-files/blog/01208083168/sort.py
// "Tarjan's algorithm and topological sorting implementation in Python" by Paul Harrison
// Step 1: get list of all involved methods
List <HighLevel.HlGraph> Nodes = new List <HighLevel.HlGraph>();
Nodes.Add(CacheEntry.HlGraph);
for (int i = 0; i < Nodes.Count; i++)
{
HighLevel.HlGraph Entry = Nodes[i];
foreach (HlGraphEntry SubEntry in Entry.RelatedGraphs.Values)
{
if (!Nodes.Contains(SubEntry.HlGraph))
{
Nodes.Add(SubEntry.HlGraph);
}
}
}
// Step 2: topological sort
Dictionary <HighLevel.HlGraph, int> count = new Dictionary <HighLevel.HlGraph, int>();
foreach (HighLevel.HlGraph Current in Nodes)
{
count[Current] = 0;
}
foreach (HighLevel.HlGraph Current in Nodes)
{
foreach (HlGraphEntry SuccessorEntry in Current.RelatedGraphs.Values)
{
count[SuccessorEntry.HlGraph]++;
}
}
List <HighLevel.HlGraph> Ready = new List <HighLevel.HlGraph>();
List <HighLevel.HlGraph> Result = new List <HighLevel.HlGraph>(Nodes.Count);
Ready.Add(CacheEntry.HlGraph);
while (Ready.Count > 0)
{
HighLevel.HlGraph Current = Ready[Ready.Count - 1];
Ready.RemoveAt(Ready.Count - 1);
Result.Add(Current);
System.Diagnostics.Debug.Assert(count[Current] == 0);
foreach (HlGraphEntry Successor in Current.RelatedGraphs.Values)
{
count[Successor.HlGraph]--;
if (count[Successor.HlGraph] == 0)
{
Ready.Add(Successor.HlGraph);
}
}
}
// Step 3: check for recursions. If there is any strongly-connected component, count[s]
// will never reach zero, so the Ready list runs empty without all functions being
// inserted into the Result list.
if (Result.Count != Nodes.Count)
{
throw new InvalidOperationException("Unable to compute topological sort of functions. Recursions are not supported.");
}
// Generate code for all HlGraphs in the Result list, in reverse order
Result.Reverse();
CacheEntry.Source = getOpenCLSource(Result);
}
public InvokeContext(HighLevel.HlGraph HLgraph)
: this(HLgraph, GetRandomSeed())
{
}
public static string GetOpenClType(HighLevel.HlGraph HlGraph, Type DataType)
{
return(InnerGetOpenClType(HlGraph, DataType));
}
private static string InnerGetOpenClType(HighLevel.HlGraph HlGraph, Type DataType)
{
if (DataType == typeof(void))
{
return("void");
}
else if (DataType == typeof(sbyte))
{
return("char");
}
else if (DataType == typeof(byte))
{
return("uchar");
}
else if (DataType == typeof(short))
{
return("short");
}
else if (DataType == typeof(ushort))
{
return("ushort");
}
else if (DataType == typeof(int) || DataType == typeof(IntPtr) || DataType == typeof(bool))
{
return("int");
}
else if (DataType == typeof(uint) || DataType == typeof(UIntPtr))
{
return("uint");
}
else if (DataType == typeof(long))
{
return("long");
}
else if (DataType == typeof(ulong))
{
return("ulong");
}
else if (DataType == typeof(float))
{
return("float");
}
else if (DataType == typeof(double))
{
return("double");
}
else if (DataType.IsByRef)
{
return(InnerGetOpenClType(HlGraph, DataType.GetElementType()) + "*");
}
else if (DataType.IsArray)
{
return(InnerGetOpenClType(HlGraph, DataType.GetElementType()) + "*");
}
else if (object.ReferenceEquals(DataType, RandomStateDataType))
{
return("uint2");
}
else if (DataType.IsValueType && DataType.BaseType == typeof(System.ValueType) && HlGraph != null && HlGraph.ValueTypeMap != null)
{
string Name;
if (!HlGraph.ValueTypeMap.TryGetValue(DataType, out Name))
{
HlGraph.ValueTypeMap[DataType] = Name = "genstruct_" + HlGraph.ValueTypeMap.Count;
}
return("struct " + Name);
}
else
{
throw new ArgumentException(string.Format("Sorry, data type '{0}' cannot be mapped to OpenCL.", DataType));
}
}
internal static void CallOpenCLNet(int[] WorkSize, HlGraphEntry CacheEntry, InvokeContext ctx, HighLevel.HlGraph HLgraph, OpenCLNet.Device device)
{
// We can invoke the kernel using the arguments from ctx now :)
if (device == null)
{
device = GetFirstGpu();
if (device == null)
{
device = GetFirstCpu();
}
if (device == null)
{
throw new ArgumentNullException("device");
}
}
OpenCLNet.Platform Platform = device.Platform;
OpenCLNet.Context context;
OpenCLNet.Program program;
lock (CacheEntry)
{
context = CacheEntry.Context;
if (context == null)
{
IntPtr[] properties = new IntPtr[]
{
new IntPtr((long)OpenCLNet.ContextProperties.PLATFORM), Platform.PlatformID,
IntPtr.Zero,
};
context = CacheEntry.Context = Platform.CreateContext(properties, new OpenCLNet.Device[] { device }, null, IntPtr.Zero);
}
program = CacheEntry.Program;
if (program == null)
{
StringBuilder source = new StringBuilder();
source.Append(GetOpenCLSourceHeader(Platform, device, CacheEntry));
source.Append(CacheEntry.Source);
source.Append(GetOpenCLSourceFooter(Platform, device));
program = context.CreateProgramWithSource(source.ToString());
try
{
program.Build();
}
catch (Exception ex)
{
string err = program.GetBuildLog(device);
throw new Exception(err, ex);
}
CacheEntry.Program = program;
}
}
using (CallContext CallContext = new CallContext(context, device, OpenCLNet.CommandQueueProperties.PROFILING_ENABLE, program.CreateKernel(HLgraph.MethodName)))
{
OpenCLNet.CommandQueue commandQueue = CallContext.CommandQueue;
for (int i = 0; i < ctx.Arguments.Count; i++)
{
ctx.Arguments[i].WriteToKernel(CallContext, i);
}
//OpenCLNet.Event StartEvent, EndEvent;
//commandQueue.EnqueueMarker(out StartEvent);
IntPtr[] GlobalWorkSize = new IntPtr[WorkSize.Length];
for (int i = 0; i < WorkSize.Length; i++)
{
GlobalWorkSize[i] = new IntPtr(WorkSize[i]);
}
commandQueue.EnqueueNDRangeKernel(CallContext.Kernel, (uint)GlobalWorkSize.Length, null, GlobalWorkSize, null);
for (int i = 0; i < ctx.Arguments.Count; i++)
{
ctx.Arguments[i].ReadFromKernel(CallContext, i);
}
commandQueue.Finish();
//commandQueue.EnqueueMarker(out EndEvent);
//commandQueue.Finish();
//ulong StartTime, EndTime;
//StartEvent.GetEventProfilingInfo(OpenCLNet.ProfilingInfo.QUEUED, out StartTime);
//EndEvent.GetEventProfilingInfo(OpenCLNet.ProfilingInfo.END, out EndTime);
}
}
internal static void WriteOpenCL(HighLevel.HlGraph HLgraph, TextWriter writer)
{
writer.WriteLine();
writer.WriteLine("// OpenCL source for {2} method '{0}' of type '{1}'", HLgraph.MethodBase.ToString(), HLgraph.MethodBase.DeclaringType.ToString(), HLgraph.IsKernel ? "kernel" : "related");
writer.WriteLine("{2}{0} {1}(", GetOpenClType(HLgraph, ((MethodInfo)HLgraph.MethodBase).ReturnType), HLgraph.MethodName, HLgraph.IsKernel ? "__kernel " : string.Empty);
for (int i = 0; i < HLgraph.Arguments.Count; i++)
{
HighLevel.ArgumentLocation Argument = HLgraph.Arguments[i];
string AttributeString = string.Empty;
if ((Argument.Flags & HighLevel.LocationFlags.IndirectRead) != 0)
{
AttributeString += "/*[in";
}
if ((Argument.Flags & HighLevel.LocationFlags.IndirectWrite) != 0)
{
if (AttributeString == string.Empty)
{
AttributeString += "/*[out";
}
else
{
AttributeString += ",out";
}
}
if (AttributeString != string.Empty)
{
AttributeString += "]*/ ";
}
if ((Argument.DataType.IsArray || Argument.DataType.IsPointer || Argument.DataType.IsByRef) &&
((Argument.Flags & Hybrid.MsilToOpenCL.HighLevel.LocationFlags.PointerLocal) == 0))
{
AttributeString += "__global ";
}
writer.WriteLine("\t{0}{1} {2}{3}", AttributeString, GetOpenClType(HLgraph, Argument.DataType), Argument.Name, i + 1 < HLgraph.Arguments.Count ? "," : string.Empty);
}
writer.WriteLine(")");
writer.WriteLine("/*");
writer.WriteLine(" Generated by CIL2OpenCL");
writer.WriteLine("*/");
writer.WriteLine("{");
foreach (HighLevel.LocalVariableLocation LocalVariable in HLgraph.LocalVariables)
{
string AttributeString = string.Empty;
if ((LocalVariable.Flags & HighLevel.LocationFlags.Read) != 0)
{
if (AttributeString == string.Empty)
{
AttributeString += "/*[";
}
else
{
AttributeString += ",";
}
AttributeString += "read";
}
if ((LocalVariable.Flags & HighLevel.LocationFlags.Write) != 0)
{
if (AttributeString == string.Empty)
{
AttributeString += "/*[";
}
else
{
AttributeString += ",";
}
AttributeString += "write";
}
if ((LocalVariable.Flags & HighLevel.LocationFlags.IndirectRead) != 0)
{
if (AttributeString == string.Empty)
{
AttributeString += "/*[";
}
else
{
AttributeString += ",";
}
AttributeString += "deref_read";
}
if ((LocalVariable.Flags & HighLevel.LocationFlags.IndirectWrite) != 0)
{
if (AttributeString == string.Empty)
{
AttributeString += "/*[";
}
else
{
AttributeString += ",";
}
AttributeString += "deref_write";
}
if (AttributeString == string.Empty)
{
AttributeString = "/*UNUSED*/ // ";
}
else
{
AttributeString += "]*/ ";
}
if ((LocalVariable.Flags & Hybrid.MsilToOpenCL.HighLevel.LocationFlags.PointerGlobal) != 0)
{
AttributeString += "__global ";
}
writer.WriteLine("\t{0}{1} {2};", AttributeString, GetOpenClType(HLgraph, LocalVariable.DataType), LocalVariable.Name);
}
HighLevel.BasicBlock FallThroughTargetBlock = HLgraph.CanonicalStartBlock;
for (int i = 0; i < HLgraph.BasicBlocks.Count; i++)
{
HighLevel.BasicBlock BB = HLgraph.BasicBlocks[i];
if (BB == HLgraph.CanonicalEntryBlock || BB == HLgraph.CanonicalExitBlock)
{
continue;
}
if (FallThroughTargetBlock != null && FallThroughTargetBlock != BB)
{
writer.WriteLine("\tgoto {0};", FallThroughTargetBlock.LabelName);
}
FallThroughTargetBlock = null;
writer.WriteLine();
if (BB.LabelNameUsed)
{
writer.WriteLine("{0}:", BB.LabelName);
}
else
{
writer.WriteLine("//{0}: (unreferenced block label)", BB.LabelName);
}
foreach (HighLevel.Instruction Instruction in BB.Instructions)
{
writer.WriteLine("\t{0}", Instruction.ToString());
}
if (BB.Successors.Count == 0)
{
writer.WriteLine("\t// End of block is unreachable");
}
else if (BB.Successors[0] == HLgraph.CanonicalExitBlock)
{
writer.WriteLine("\t// End of block is unreachable/canonical routine exit");
}
else
{
FallThroughTargetBlock = BB.Successors[0];
}
}
writer.WriteLine("}");
}
private static void WriteCode(HighLevel.HlGraph HLgraph, TextWriter writer)
{
writer.WriteLine("// begin {0}", HLgraph.MethodBase);
if (HLgraph.MethodBase.IsConstructor)
{
writer.Write("constructor {0}::{1} (", ((System.Reflection.ConstructorInfo)HLgraph.MethodBase).DeclaringType, HLgraph.MethodBase.Name);
}
else
{
writer.Write("{0} {1}(", ((MethodInfo)HLgraph.MethodBase).ReturnType, HLgraph.MethodBase.Name);
}
for (int i = 0; i < HLgraph.Arguments.Count; i++)
{
if (i > 0)
{
writer.Write(", ");
}
HighLevel.ArgumentLocation Argument = HLgraph.Arguments[i];
string AttributeString = string.Empty;
if ((Argument.Flags & HighLevel.LocationFlags.IndirectRead) != 0)
{
AttributeString += "__deref_read ";
}
if ((Argument.Flags & HighLevel.LocationFlags.IndirectWrite) != 0)
{
AttributeString += "__deref_write ";
}
writer.Write("{0}{1} {2}", AttributeString, Argument.DataType, Argument.Name);
}
writer.WriteLine(") {");
foreach (HighLevel.LocalVariableLocation LocalVariable in HLgraph.LocalVariables)
{
writer.WriteLine("\t{0} {1};", LocalVariable.DataType, LocalVariable.Name);
}
for (int i = 0; i < HLgraph.BasicBlocks.Count; i++)
{
HighLevel.BasicBlock BB = HLgraph.BasicBlocks[i];
if (BB == HLgraph.CanonicalEntryBlock || BB == HLgraph.CanonicalExitBlock)
{
continue;
}
writer.WriteLine();
writer.WriteLine("{0}:", BB.LabelName);
foreach (HighLevel.Instruction Instruction in BB.Instructions)
{
writer.WriteLine("\t{0}", Instruction.ToString());
}
if (BB.Successors.Count == 0)
{
writer.WriteLine("\t// unreachable code");
}
else if (i + 1 == HLgraph.BasicBlocks.Count || HLgraph.BasicBlocks[i + 1] != BB.Successors[0])
{
if (BB.Successors[0] == HLgraph.CanonicalExitBlock)
{
writer.WriteLine("\t// to canonical routine exit");
}
else
{
writer.WriteLine("\tgoto {0};", BB.Successors[0].LabelName);
}
}
}
writer.WriteLine("}");
writer.WriteLine("// end");
writer.WriteLine();
}
private void TraverseTree(Node Node, bool TopIsDef, bool ArrayDef)
{
if (Node == null)
{
return;
}
if (TopIsDef)
{
if (Node.NodeType == NodeType.Location)
{
Location Location = ((LocationNode)Node).Location;
if (!DefinedLocations.Contains(Location))
{
DefinedLocations.Add(Location);
}
}
else if (Node.NodeType == NodeType.ArrayAccess)
{
TraverseTree(Node, false, true);
}
else if (Node.NodeType == NodeType.AddressOf)
{
TraverseTree(((AddressOfNode)Node).SubNodes[0], true, false);
TraverseTree(((AddressOfNode)Node).SubNodes[0], false, false);
}
else if (Node.NodeType == NodeType.Deref)
{
TraverseTree(((DerefNode)Node).SubNodes[0], false, false);
}
else if (Node.NodeType == NodeType.InstanceField)
{
// Do nothing here. ConvertForOpenCl will convert this to an argument location,
// so regular analysis treats it as such at a later pass
}
else if (Node.NodeType == NodeType.NamedField)
{
TraverseTree(((NamedFieldNode)Node).SubNodes[0], true, false);
}
else
{
System.Diagnostics.Debugger.Break();
}
}
else
{
if (Node.NodeType == NodeType.Location)
{
Location Location = ((LocationNode)Node).Location;
if (!UsedLocations.Contains(Location))
{
UsedLocations.Add(Location);
}
}
else
{
HlGraphEntry RelatedGraphEntry;
if (Node.NodeType == NodeType.ArrayAccess && Node.SubNodes.Count > 0 && Node.SubNodes[0].NodeType == NodeType.Location)
{
Location ArrayLocation = ((LocationNode)Node.SubNodes[0]).Location;
if (ArrayDef && !IndirectDefinedLocations.Contains(ArrayLocation))
{
IndirectDefinedLocations.Add(ArrayLocation);
}
else if (!ArrayDef && !IndirectUsedLocations.Contains(ArrayLocation))
{
IndirectUsedLocations.Add(ArrayLocation);
}
}
else if (Node.NodeType == NodeType.Call && Node.SubNodes.Count > 0 && !object.ReferenceEquals(RelatedGraphs, null) &&
RelatedGraphs.TryGetValue(((CallNode)Node).MethodInfo, out RelatedGraphEntry))
{
// Call to other generated method. Propagate indirect usage flags
HlGraph HlGraph = RelatedGraphEntry.HlGraph;
for (int i = 0; i < Math.Min(Node.SubNodes.Count, HlGraph.Arguments.Count); i++)
{
ArgumentLocation Argument = HlGraph.Arguments[i];
if ((Argument.Flags & (LocationFlags.IndirectRead | LocationFlags.IndirectWrite)) != 0)
{
Node SubNode = Node.SubNodes[i];
if (SubNode.NodeType == NodeType.Location)
{
Location ArrayLocation = ((LocationNode)SubNode).Location;
if ((Argument.Flags & LocationFlags.IndirectRead) != 0 && !IndirectUsedLocations.Contains(ArrayLocation))
{
IndirectUsedLocations.Add(ArrayLocation);
}
if ((Argument.Flags & LocationFlags.IndirectWrite) != 0 && !IndirectDefinedLocations.Contains(ArrayLocation))
{
IndirectDefinedLocations.Add(ArrayLocation);
}
}
}
}
}
foreach (Node SubNode in Node.SubNodes)
{
TraverseTree(SubNode, (SubNode.NodeType == NodeType.AddressOf) ? true : false, false);
}
}
}
}
/// <summary>
/// Perform intra-block copy propagation.
/// </summary>
/// <param name="Graph">HlGraph.</param>
/// <returns>TRUE if any changes have occured, FALSE otherwise.</returns>
public static bool Do(HlGraph Graph)
{
bool Changed = false;
CpState State = new CpState();
Dictionary <Location, LocDef> InsertedDict = new Dictionary <Location, LocDef>();
List <Location> RemoveKeys = new List <Location>();
List <KeyValuePair <Location, LocDef> > NewDefs = new List <KeyValuePair <Location, LocDef> >();
//
// Get the list of undefined locations at the end of each basic block
//
DataFlow <LocationList> .Result DfResult = DeadAssignmentElimination.DoDF(Graph);
//
// Iterate through basic blocks
//
foreach (BasicBlock BB in Graph.BasicBlocks)
{
//
// Iterate through instructions
//
State.Defs.Clear();
for (int InstructionIndex = 0; InstructionIndex < BB.Instructions.Count; InstructionIndex++)
{
Instruction Instruction = BB.Instructions[InstructionIndex];
//
// Query lists of used and defined locations for the current instruction. If it uses
// a location that is present in the CpState, the definition can only be pasted if:
//
// 1) The definition is not of the form "x := f(..., x, ...)"; i.e., the location
// does not depend on an equally-named one at the defining site,
//
// OR
//
// 2) The definition is of the form "x := f(..., x, ...)", and the definition is dead after
// the current instruction.
//
// NOTE: Pasting the definition in case (2) duplicates evaluation of "f" at the current
// instruction, as dead assignment elimination will not be able to remove the original
// defining site due to location x being used here. Therefore, in this case, we need to
// remove the instruction itself if the definition is pasted, but we can only do so if the
// value is not used afterwards.
//
LocationUsage LocationUsage = LocationUsage.ForInstruction(Instruction);
foreach (Location Location in LocationUsage.UsedLocations)
{
LocDef LocDef;
if (State.Defs.TryGetValue(Location, out LocDef) && LocDef.SelfRef)
{
LocationList DeadList = DeadAssignmentElimination.BeforeInstruction(BB, InstructionIndex + 1, DfResult.ExitState[BB]);
if (!DeadList.Contains(Location))
{
State.Defs.Remove(Location);
}
}
}
//
// Rewrite trees
//
Instruction.Argument = ForTree(BB, Instruction, State, Instruction.Argument, false, ref Changed, InsertedDict);
Instruction.Result = ForTree(BB, Instruction, State, Instruction.Result, true, ref Changed, InsertedDict);
//
// Change CpState to reflect any newly defined locations
//
foreach (Location Location in LocationUsage.DefinedLocations)
{
State.Defs.Remove(Location);
if (Location.LocationType == LocationType.CilStack || Location.LocationType == LocationType.LocalVariable)
{
System.Diagnostics.Debug.Assert(Instruction.InstructionType == InstructionType.Assignment);
if (Instruction.Result != null && Instruction.Result.NodeType == NodeType.Location)
{
LocDef NewLocDef = new LocDef(BB, Instruction, Location, Instruction.Argument);
NewDefs.Add(new KeyValuePair <Location, LocDef>(Location, NewLocDef));
}
}
//
// Remove other definition entries that have been invalidated by redefining one of their inputs
//
foreach (KeyValuePair <Location, LocDef> Entry in State.Defs)
{
if (Entry.Value.UsedLocations.Contains(Location))
{
RemoveKeys.Add(Entry.Key);
}
}
}
//
// Remove and add entries to CpState as necessary
//
if (RemoveKeys != null && RemoveKeys.Count != 0)
{
foreach (Location RemoveLocation in RemoveKeys)
{
State.Defs.Remove(RemoveLocation);
}
RemoveKeys.Clear();
}
if (NewDefs != null && NewDefs.Count != 0)
{
foreach (KeyValuePair <Location, LocDef> Entry in NewDefs)
{
State.Defs.Add(Entry.Key, Entry.Value);
}
NewDefs.Clear();
}
//
// Remove instructions at defining sites that have become obsolete. This is only
// true for CilStack locations, as they are the only ones guaranteed to be consumed
// at most once. All other obsolete instructions are taken care of by dead code
// elimination.
//
// BUGBUG: this fails for the DUP instruction...
//
if (InsertedDict.Count != 0 && LocationUsage.UsedLocations.Count != 0)
{
foreach (Location Location in LocationUsage.UsedLocations)
{
LocDef RemoveLocDef;
if (Location.LocationType == LocationType.CilStack && InsertedDict.TryGetValue(Location, out RemoveLocDef))
{
bool OK = RemoveLocDef.BasicBlock.Instructions.Remove(RemoveLocDef.Instruction);
//System.Diagnostics.Debug.Assert(OK);
System.Diagnostics.Debug.Assert(object.ReferenceEquals(RemoveLocDef.BasicBlock, BB));
if (object.ReferenceEquals(RemoveLocDef.BasicBlock, BB) && OK)
{
System.Diagnostics.Debug.Assert(object.ReferenceEquals(BB.Instructions[InstructionIndex - 1], Instruction));
InstructionIndex--;
}
}
}
}
InsertedDict.Clear();
}
}
return(Changed);
}
