Code generator bug

May 15, 2013 at 8:06 PM
C# kernel:

        [Cudafy]
        public static void UpdateTextureContents(GThread thread, Float4[,] pixels)
        {
            var i = thread.get_global_id(0);
            var j = thread.get_global_id(1);

            Float4 dummy = new Float4();

            dummy.x = 1f;
            dummy.y = 1f;
            dummy.z = 1f;
            dummy.w = 1f;

            pixels[i, j] = dummy;
        }
Generated CUDA source:
// CudafyOpenTK.MainWindow
extern "C" __global__ void UpdateTextureContents( Float4* pixels, int pixelsLen0, int pixelsLen1)
{
    int num = get_global_id(0);
    int num2 = get_global_id(1);
    Float4 float = Float4();
    float.x = 1.0f;
    float.y = 1.0f;
    float.z = 1.0f;
    float.w = 1.0f;
    pixels[(num) * pixelsLen1 + ( num2)] = float;
}
Which won't compile. Float4 should be converted to "float4" to match the CUDA datatype, and the variable is being renamed to "float" which doesn't work so well.

Suggestions before I start digging through the source? I managed to fix the Float4 problem by declaring a struct named "float4" (lower-case) in the global namespace,
public struct float4
{
    public float w;
    public float x;
    public float y;
    public float z;
}
and replacing "Float4" in the C# kernel with "float4". That just leaves the variable name issue.

I don't suppose there's an attribute or something that'll let me tell the system to keep the variable names as they are in the source? (I'm guessing those are lost before the IL inspector kicks in.)
May 15, 2013 at 9:04 PM
"I'm guessing those are lost before the IL inspector kicks in"
Yup, that's exactly what happens. I really wish there was a way around that.
btw, I wasn't aware that cudafy supported vector types. I was pretty sure it didn't.
May 15, 2013 at 9:22 PM
Edited May 15, 2013 at 9:33 PM
Yes, though it doesn't support swizzle operators and such. The structs are already available in the GASS.CUDA.Types namespace.

If I manually change the .cu output to:
__device__ int get_global_id(int dimension)
        {
            int result = 0;
            if (dimension == 0)
                result = blockIdx.x * blockDim.x + threadIdx.x;
            else if (dimension == 1)
                result = blockIdx.y * blockDim.y + threadIdx.y;
            else  if (dimension == 2)
                result = blockIdx.z * blockDim.z + threadIdx.z;
            return result;
        }

// CudafyOpenTK.MainWindow
extern "C" __global__ void UpdateTextureContents( float4* pixels, int pixelsLen0, int pixelsLen1);

// CudafyOpenTK.MainWindow
extern "C" __global__ void UpdateTextureContents( float4* pixels, int pixelsLen0, int pixelsLen1)
{
    int num = get_global_id(0);
    int num2 = get_global_id(1);
    float4 myvar = float4();
    myvar.x = 1.0f;
    myvar.y = 1.0f;
    myvar.z = 1.0f;
    myvar.w = 1.0f;
    pixels[(num) * pixelsLen1 + ( num2)] = myvar;
}
then NVCC compiles it just fine
(Well, there's an unresolved external for __tmainCRTStartup, but that's expected.)

(I just replaced Float4 with float4 and renamed the variable from "float" to "myvar")

The next question would be (and I'm actively tracing through the .Translator code at the moment trying to figure this out) how to ensure that generated var names aren't real types. I'm surprised that it thinks that "float" is a good identifier. (Not complaining - I, too, make mistakes from time to time!)

Edit: they're all already defined in the GASS.CUDA.Types namespace.
May 15, 2013 at 10:38 PM
Edited May 15, 2013 at 10:45 PM
And, here are the two problems.

In Cudafy.Translator\ICSharpCode.NRefactory\CSharp\Ast\Identifier.cs,
        public Identifier (string name, AstLocation location)
        {
            if (name == null)
                throw new ArgumentNullException("name");
            IsQuoted = name.StartsWith ("@");
            this.Name = IsQuoted ? name.Substring (1) : name;
            this.startLocation = location;
        }
In this case, the IL apparently names the variable "@float", which is valid C#. The method above, however, just drops the "@" and calls it "float." Bada-bing, BOOM.

The other is that the SpecialTypes collection doesn't know anything about vector types. I added this to Cudafy.Translator\CUDALanguage.cs:
            /// add exceptions for Float1..4, Int1..4, etc. CUDA vector types - CUDA.NET provides C# structs that are used during initial compilation.
            foreach(var cudaVectorType in new[] { "Float", "Int", "Long", "Short", "UChar", "UInt", "ULong", "UShort" })
                for(var i=1; i<=4; i++)
                {
                    var cSharpName = string.Format("{0}{1}", cudaVectorType, i);
                    var cudaName = cSharpName.ToLower();
                    SpecialTypes.Add(cSharpName, new SpecialTypeProps() { Name = cudaName });
                }
            
in CUDALanguage.InitializeCUDA(). That takes care of replacing "Float4" with "float4", as well as "IntN", "UCharN", and the other vector types that have C# struct equivelants defined.

Should I send you those two files, or can I have commit access? :) It's been awhile since I had to dig through compiler code.

You have to use the vector types when writing color textures from CUDA. My initial OpenTK test used a height field texture (1 channel float - basically a 2D float array), so it worked fine, but when you're drawing color data you have to use the vector types.
May 15, 2013 at 10:47 PM
"they're all already defined in the GASS.CUDA.Types namespace."
Yes, that would be expected, for interop.
What I meant regarding the lack of support in cudafy was that the IL parser makes no use of vector types whatsoever. You know, giving meaning to "+" operators on vectors, supporting the typical cuda math functions when applied on vectors, etc.
I'm not complaining, mind you - I never use cuda vectors.
Now, the var naming in generated code, that's something that bothers me. There's no solution, I'm very aware of that, but still it irks me; it's unaesthetic, heh. It also makes it harder to use nsight.
May 15, 2013 at 10:51 PM
"The method above, however, just drops the "@" and calls it "float." Bada-bing, BOOM."
Ah, indeed.
May 15, 2013 at 10:51 PM
Edited May 15, 2013 at 10:52 PM
Yeah, CUDA doesn't define vector operators (as far as I can tell), just the datatypes, which are effectively structs. In GLSL and HLSL you can "swizzle" on vec4/float4, etc,. to write to multiple components at the same time, but I have a sneaking suspicion that it gets compiled down to individual component operations. (Then again, the GPU gods have surprised me before)...

This wouldn't even have come up if I hadn't been converting the OpenTK interop test I wrote to something "simpler" - IE, a real texture generator instead of a height field demo. Yay.
May 15, 2013 at 11:04 PM
Correction: the "@" case in Identifier.cs never executes. Argh. Still looking.

The naïve (and possibly crappy) solution would be to track variable names that are also types in a dictionary, generate a different version (prepend some underscores or something), etc.

I have to fix this on my end to proceed, would be happy to contribute the changes back, but this is pretty low-level so I'm not sure how comfortable y'all are with that.
May 15, 2013 at 11:43 PM
Edited May 16, 2013 at 12:14 AM
"This wouldn't even have come up if I hadn't been converting the OpenTK interop test I wrote to "
That's the universe telling you how you shouldn't be helpful to others.

"Yeah, CUDA doesn't define vector operators "
yes, it doesn't. I meant that only as helpful feature of cudafy which would translate as "standard" operations in CUDA C. OpenCL does a lot of smart stuff with vectors, which could justify all this syntax sugar.

Remember to match your float4 struct's StuctLayoutArrtibute with the allignment used by CUDA's vector.

Anyway, why don't you simply "cudafy" your float4 struct and use that instead of trying to make do with GASS structs (which are there only for interop)? You could even call it "myFloat4" to make your bug go away. Once the kernel returns your myFloat4[] array, you can just pin it onto a intPtr and do whatever else you need.
May 16, 2013 at 2:42 AM
That fixes the type problem, but not the variable name. Still looking for that one.

Sent from my iPhone

Coordinator
May 16, 2013 at 11:18 AM
Vector support is on the to-do list. Handling OpenCL, too, means we need to be extra careful. Can either of you draw up a table with the equivalent OpenCL and CUDA names for the various vectors?
May 16, 2013 at 1:10 PM
I'll get such a table ready, together with aditional info.
May 16, 2013 at 2:27 PM
Thanks, all.

Sent from my iPhone