This project is read-only.

Debug.WriteLine[If]() and Console.WriteLine()

Jan 20, 2012 at 5:18 AM
Edited Jan 20, 2012 at 10:50 AM

This new feature will only work on CUDA 2.0+, And This isn't perfect, because...

#1 It doesn't translate String.Format() rules, you have to use printf() rules. It doesn't look like the decompiler knows all the types of each argument passed so there is not a generic way to convert {0} to a printf parameter, it requires the proper type, like %0$d or %0$s. So if you use printf rules you should be good when running on the gpu.

#2 Forced requirement that the first parameter to WriteLine() must be a string literal not a variable. So only use

Debug.WriteLine("My Message");
// or
Debug.WriteLine("The cordinates are %d, %d", point.x, point.y); 

Instead of something like

string msg = "The price is one BILLION dollars";
Debug.WriteLine(msg);
// or
Debug.WriteLine("The cordinates are " + point.x + ", " + point.y);

 

Here's the code I came up with...
CUDALanguage.cs > replace...

SpecialMethods.Add(new SpecialMember("Debug", null, new Func<MemberReferenceExpression, object, string>(CommentMeOut), false));
SpecialMethods.Add(new SpecialMember("Console", null, new Func<MemberReferenceExpression, object, string>(CommentMeOut), false));

with...

SpecialMethods.Add(new SpecialMember("Debug", "Write", new Func<MemberReferenceExpression, object, string>(TranslateToPrintF), false));
SpecialMethods.Add(new SpecialMember("Debug", "WriteIf", new Func<MemberReferenceExpression, object, string>(TranslateToPrintF), false));
SpecialMethods.Add(new SpecialMember("Debug", "WriteLine", new Func<MemberReferenceExpression, object, string>(TranslateToPrintF), false));
SpecialMethods.Add(new SpecialMember("Debug", "WriteLineIf", new Func<MemberReferenceExpression, object, string>(TranslateToPrintF), false));
SpecialMethods.Add(new SpecialMember("Debug", null, new Func<MemberReferenceExpression, object, string>(CommentMeOut), false));
SpecialMethods.Add(new SpecialMember("Console", "Write", new Func<MemberReferenceExpression, object, string>(TranslateToPrintF), false));
SpecialMethods.Add(new SpecialMember("Console", "WriteLine", new Func<MemberReferenceExpression, object, string>(TranslateToPrintF), false));
SpecialMethods.Add(new SpecialMember("Console", null, new Func<MemberReferenceExpression, object, string>(CommentMeOut), false));

CUDALanguage.cs > add function...

static string TranslateToPrintF(MemberReferenceExpression mre, object data)
{
    // TODO: detect architecture
    //if (CUDALanguage.Architecture != eArchitecture.sm_20)
    //    return CommentMeOut(mre, data);
    UseOptionalHeader(csSTDIO);
    string dbugwrite = string.Empty;
    dbugwrite = mre.TranslateToPrintF(data);
    return dbugwrite;
}

CUDALanguage.cs > add const after csCURAND_KERNAL const...

private const string csSTDIO = "stdio";

CUDALanguage.cs > add optional header after csCURAND_KERNAL code... 

OptionalHeaders.Add(new OptionalHeader(csSTDIO, @"#include <stdio.h>"));

CUDAOutputVisitor.cs > add function...

public static string TranslateToPrintF(this MemberReferenceExpression mre, object data)
{
    TextWriter output = new StringWriter();
    CUDAOutputVisitor visitor = new CUDAOutputVisitor(new TextWriterOutputFormatter(output), new CSharpFormattingOptions());
    var ex = data as InvocationExpression;
    if (ex == null)
        throw new ArgumentNullException("data as InvocationExpression");

    bool isWriteLine = ((MemberReferenceExpression)ex.Target).MemberName.StartsWith("WriteLine");
    bool hasIfCondition = ((MemberReferenceExpression)ex.Target).MemberName.EndsWith("If");
    List<Expression> arguments = ex.Arguments.ToList();
    int i = 0;

    if (hasIfCondition)
    {
        i = -1;
        output.Write("if(");
        arguments[0].AcceptVisitor(visitor, data);
        output.Write(") ");
    }

    output.Write("printf(");
    foreach (var arg in arguments)
    {
        if (i == -1) { /* skip it, it was an if condition */ }
        else if (i == 0)
        {
            if (!(arg is PrimitiveExpression))
                throw new CudafyLanguageException("When using Debug.Write" + (isWriteLine ? "Line" : "") + (hasIfCondition ? "If" : "") + "() the first parameter must be a string literal");

            string strFormat = arg.ToString();
            if (!strFormat.StartsWith("\""))
                throw new CudafyLanguageException("When using Debug.Write" + (isWriteLine ? "Line" : "") + "() the first parameter must be a string literal");

            if (hasIfCondition)
            {
                // Since debug if condition messages don't support parameters, escape any % to avoid ErrorUnknown
                strFormat = strFormat.Replace("%", "%%");
            }

            // NOTE: unless you know the type of each parameter passed to printf, 
            // then I don't see a good way to convert String.Format() rules to printf() rules
            //strFormat = CUDALanguage.TranslateStringFormat(strFormat);

            if (isWriteLine)
                strFormat = strFormat.Insert(strFormat.Length - 1, "\\n");

            output.Write(strFormat);
        }
        else if (hasIfCondition) {
            // Skip any other parameters if they are there, debug if conditions don't support formatting
        }
        else if (arg is ArrayCreateExpression)
        {
            var arrayCreateExpression = arg as ArrayCreateExpression;
            foreach (var arrayArg in arrayCreateExpression.Initializer.Elements)
            {
                output.Write(',');
                arrayArg.AcceptVisitor(visitor, data);
            }
        }
        else
        {
            output.Write(',');
            arg.AcceptVisitor(visitor, data);
        }
        i++;
    }
    output.Write(")");
    return output.ToString();
}

 

I knew there had to be a way debug on the device without using Nsight. I hope this helps other folks trying to debug some problematic code.

Ohhhhh almost forgot, one last important thing to know if your not getting any messages printed out to the console is that you may need to call gpu.Synchronize() to get the messages to appear.

Jan 20, 2012 at 10:16 AM

This looks quite good even with the restrictions.  It was a planned feature for quite some time, but there were some troublesome issues.  I also believed (possibly wrongly) that you also need to make a call in the host code to enable this functionality.

Another issue is indeed not translating the place holders in the string.  If you run using the emulator you may get some general bad news when formatting takes place.  I therefore hesitate before adding it to the release code.  Do you see a solution?  Do you not have access to the type of each argument passed to String.Format?

Jan 20, 2012 at 10:35 AM

There wasn't a call that I had to make in the host code, only thing was to set the arch to sm_20 (cuda 2.0+). CUDA actually calls printf() on the host automatically so it's printed to the console.

If your running in the emulator there shouldn't be much problem besides that you would need to use String.Format rules instead of printf() rules. So {0} instead of %d. There wouldn't be an error, just your msg wouldn't print the way you intended if you left printf() rules when running in the emulator.

A solution for those who use the emulator could be to define a compile condition symbol like EMULATOR to avoid having to modify the debug statements if you switch back and forth. You could then do

#if EMULATOR
Debug.WriteLine("My message {0}", "is here");
#else
Debug.WriteLine("My message %s", "is here");
#endif

I think this feature is key to new folks like myself, it's a pain going from c# exception handling to C no exception handling. I was having to comment out code and use temp arguments to see what was going wrong, and I was using random numbers in my function so I wasn't running into the problems in emulation, because emulation was too slow for me to wait around for it to hit a bad random number.

I'm actually modifying this at the moment to also support WriteIf and WriteLineIf. An additional tip to users is if you compile in Release mode instead of debug, the entire debug statement does not even get generated in the cuda code just like c# releases.

Jan 20, 2012 at 10:52 AM

I just updated the code in the first post to now support Debug.WriteIf(condition, msg); and Debug.WriteLineIf(condition, msg);

Jan 20, 2012 at 11:32 AM

Nick,

Just to clarify the part about not knowing the type of the parameters. In C# when you do {0} you could easily just convert any arguments to a string by using .ToString() and inserting it there without needing the type. However printf() will error if you use %s for a numeric parameter, you must use %d. I started writing the code for the conversion originally thats when I realized the flags for printf() must identify the type each parameter is.

When testing Debug.WriteLine("thread {0}", thread.threadIdx.x); I would need to know that thread.threadIdx.x resolves as an integer to replace {0} with %d. When debugging that resolved as UnknownType.

I also later found out that I dont think the printf function supports %0$d which 0$ would be the parameter position used to repeat a paramter, the cuda docs don't list that as an option, but it also said the printf does not validate the flags, because ultimately the formatting is done by whatever the host supports



Anyway, I hope you put Debug.WriteLine[If](), Debug.Assert(), and local array initalization posts in the release code for all to use.

Jan 21, 2012 at 8:08 PM

So whats your verdict nick?

Jan 23, 2012 at 8:05 AM

It looks good.  However what happens if architecture is not 2.0+ ?  Possible solution will be to use your new behaviour  for Fermi and default to the existing behaviour for pre-Fermi.  What do you think?

Jan 23, 2012 at 8:13 AM

That code already does that, if you fill in this TODO

    // TODO: detect architecture
    //if (CUDALanguage.Architecture != eArchitecture.sm_20)
    //    return CommentMeOut(mre, data);

 I was thinking in the Cudafy method you could set a static property on CUDALanguage then just uncomment that code... however != sm_20 is incorrect, it should be >= sm_20 since there is a sm_21 but its not defined in cudafy.

I was thinking of making one more check on this, I think maybe I overlooked the annotations for the parameters from Mono.Cecil, but I'm pretty sure when I looked it only had a place holder like 1A-74 or something like that.

Anyway I'm working on many advancements to cudafy right now that are supported by cuda 2.0+, but I'm thinking of making the Cudafy Translator extendable and making a more open source base for these modifications, and would complement cudafy but not require you to implement


Jan 27, 2012 at 10:23 AM

Compute capability is tested for 2.x.

Committed and working nicely!  Thank you.