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