Passing structs

Sep 19, 2013 at 10:24 AM
I've been trying to pass a struct to the kernel, with no luck.
I've looked in the examples, but as far as I can tell they seem to be throwing the same error.

In CudafyExamples, ComplexNumbers.cs, upon calling CudafyTranslator.Cudafy I get this error:
CUDAFYSOURCETEMP.cu(4): error: identifier "a" is undefined
and if I look in the .cu file line 4 looks like this:
extern "C" global void complexAdd(* a, int aLen0, int aLen1, * b, int bLen0, int bLen1, * c, int cLen0, int cLen1);
which does seem to be missing the variable type.

What am I doing wrong?
Is there a cuda compute capability required for structs?
Coordinator
Sep 19, 2013 at 12:59 PM
"Works on my machine" as they say. Do you have the Target and Language set to Cuda?
Sep 19, 2013 at 5:28 PM
yes
Sep 19, 2013 at 10:28 PM
I've just upgraded to the latest version of Cudafy and this works.
I think I had a development version from when OpenCL was being implemented.
Jan 20, 2014 at 8:27 PM
Edited Jan 20, 2014 at 8:46 PM
Hello mcmillab or all else out there reading this post,

I am having difficulty passing a struct to a kernel as well. I am not 100% sure how to this is done in CUDAfy and after looking through the users manual and looking online, I have yet to have any luck. Below, I have attached what I have so far:

        public struct testStruct
        {
            public double a;
            public double b;
            public double c;

            public testStruct()
            {
                a = 20.0;
                b = 30.0;
                c = 0.00001;
            }
        }

        static void Main(string[] args)
        {

            testStruct testValues;

            try
            {
                CudafyModule km = CudafyTranslator.Cudafy(ePlatform.x64, eArchitecture.sm_20);
                _gpu = CudafyHost.GetDevice(eGPUType.Cuda);
                _gpu.LoadModule(km);

                testStruct dev_struct = _gpu.CopyToDevice(testValues);


                _gpu.Launch(500, 1).kernel(dev_testStruct);

                _gpu.Synchronize();

                _gpu.FreeAll();

            }

...
There error the compiler gives me is: The best overloaded method match for 'Cudafy.Host.GPGPU.CopyToDevice(string)' has some invalid arguments.

If someone could point me in the right direction, it would be greatly appreciated.

Thanks

SOLVED: I was able to find a good example within the CUDAfyExamples solution.
Coordinator
Jan 21, 2014 at 5:48 AM
Glad you found your answer. I will write a reply in case others look at this. You will need to put an [Cudafy] attribute on your struct and pass it as a type to the Cudafy method, along with the class containing your device method. When passing structs you will need to put them into an array before copying to device.
Jan 21, 2014 at 11:02 AM
Edited Jan 21, 2014 at 11:10 AM
Hi Nick,

Thank you for the reply. I actually am still struggling with this problem. I have attempted to follow the example within the CudafyExamples solution along with your instructions but I am getting the following error:

"Error: Could not find function 'kernel' in module." Below I have some of the source code:

static void Main(string[] args)
        {

            testStruct[] structArray = new testStruct[1];

            structArray[0].a = 20.0;
            structArray[0].b = 30.0;
            structArray[0].c = 0.0001;

            
            try
            {
                var km = CudafyModule.TryDeserialize();
                if (km == null || !km.TryVerifyChecksums())
                {

                    km = CudafyTranslator.Cudafy(ePlatform.Auto, eArchitecture.sm_20, typeof(testStruct));
                    km.Serialize();
                }

                _gpu = CudafyHost.GetDevice(CudafyModes.Target,0);
                _gpu.LoadModule(km);

                testStruct[] dev_structArray = _gpu.CopyToDevice(structArray);


                _gpu.Launch(1, 1).kernel(dev_structArray); 

                _gpu.FreeAll();

            }

            catch (CudafyLanguageException cle)
            {
                HandleException(cle);
            }
            catch (CudafyCompileException cce)
            {
                HandleException(cce);
            }
            catch (CudafyHostException che)
            {
                HandleException(che);
            }           

        }
       
       
        [Cudafy]
        public struct testStruct
        {
            public double a;
            public double b;
            public double c;
        }

        [Cudafy]
        public static void kernel(GThread thread, testStruct[] testValues)
        {
                int x = thread.blockIdx.x;
                testValues[0].c = x*20;
                testValues[0].a = x*30;

        }

        ...
I am not sure what this error means. Could you please point me in the right direction as how to fix this problem?

Thanks,

Joe
Coordinator
Jan 21, 2014 at 1:32 PM
It means the function called kernel has not be cudafied. When you explicitly name the types you want to cudafy then you must include ALL types. The magic cudafy overload method - no types specified - only cudafies the calling type. So add the type in which Main is a method (probably Program).
Jan 23, 2014 at 11:21 AM
Thanks again Nick for the response. Your suggestion worked for my previous problem.

I now am attempting copy a struct that resides within a different class than my GPU kernel. When I CUDAfy the class where my kernel code lives and where I call the struct, I get the exception: '(11): error: attribute "global" does not apply here'

When I CUDAfy the class where my struct lives, I get the same error as before: 'Error: Cannot find function 'kernel' in module.' I experimented a bit and attempted to CUDAfy both classes (the one holding my struct and the other which holds my kernel and calls the struct) and when I do that I get a different exception that says 'Additional information: An item with the same key has already been added.' which I believe is not an exception from CUDAfy. Might you have any suggestions on how I can resolve this issue?

Thank you again,

Joe