Beispiel #1
0
        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();
        }
Beispiel #2
0
        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("}");
        }
Beispiel #3
0
        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);
        }