This project is read-only.

Struct definition order in .cu

Dec 19, 2011 at 6:35 PM

The struct DependsOnB depends on struct B.

Therefore, I do this in my C# code:

var types = new[]{ typeof (B), typeof (DependsOnB )};

Followed by:

CudafyModule cudafyModule = CudafyTranslator.Cudafy(types);

However in the generated .cu code, struct DependsOnB seems to be placed before B, making DependsOnB complain about B not being defined :/.

Dec 19, 2011 at 7:59 PM
Edited Dec 20, 2011 at 2:02 AM

I managed to circumvent this bug by making this


/// <summary>
	/// Better than CudafyTranslator.Cudafy
	/// </summary>
	class SuperCudafier
	{
		/// <summary>
		/// Cudafies the types in the provided list, *in order*
		/// </summary>
		/// <param name="types"></param>
		/// <returns></returns>
		public static CudafyModule Cudafy(System.Type[] types)
		{
			CudafyModule complete = CudafyModule.TryDeserialize("complete");
			if (complete == null || !complete.TryVerifyChecksums())
			{
				complete = new CudafyModule();
				CudafyModule module = new CudafyModule();
				foreach (var type in types)
				{
					// Errors reported from (seems to originate) here (*) and their solution (->):
					// *NullReferenceException: "Object reference not set to an instance of an object."
					//   -> Constant device memory must be marked as 'public', not 'internal' or 'private', eg.
					//      public static Item[] DevPile = new Item[NumberOfItems];
					// *CudafyLanguageException: "Array create expressions are not supported."
					//   -> Struct to be cudafied cannot contain and instantiate regular array
					// *ArgumentNullException: "Value cannot be null. Parameter name: source"
					//   -> must not mark struct with [Cudafy] which has array that must be marked with [Cudafy], i.e. /do/
					//        public struct Container
					//      instead of
					//        [Cudafy] public struct Container
					//      when the struct contains
					//        [Cudafy] public static Item[] DevItems = new Item[Count];
					//   -> must not mark array with [Cudafy] is inside struct that must be marked with [Cudafy], i.e. /do/
					//        public static Item[] DevItems = new Item[Count];
					//      instead of
					//        [Cudafy] public static Item[] DevItems = new Item[Count];
					//      when the struct is marked as
					//        [Cudafy] public struct Container
					// *TargetException: "Non-static field requires a target."
					//   -> Did you forget to Cudafy or [Cudafy]-mark a struct which you send to the kernel,
					//      or any of the struct's subparts?
					// ---
					// Use (mouseover) 'name' to find the name of the type which contains the error! :)
					string name = type.Name;
					module =
						CudafyTranslator.Cudafy(ePlatform.Auto, eArchitecture.sm_12,
							                    cudaVersion: null, compile: false, types: type);
					// do NOT compile
					complete.CudaSourceCode += module.CudaSourceCode;
				}
				// move over all else to complete -- highly crucial
				module.CudaSourceCode = complete.CudaSourceCode;
				complete = module;

				// Errors reported from (seems to originate) here (*) and their solution (->):
				// *?: "the size of an array must be greater than zero"
				//   -> did this error appear upon second launching of the kernel?
				// *CudafyCompileException: "error: no operator "[]" matches these operands. operand types are: Item [ int ]"
				//   -> GPU constant memory definitions must be instantiated, i.e. like this:
				//      [Cudafy] public static Item[] DevPile = new Item[NumberOfItems];
				//      and not like this:
				//      [Cudafy] public static Item[] DevPile;
				// *CudafyCompileException: "error: attribute "__global__" does not apply here"
				//   -> Did you forget to Cudafy or [Cudafy]-mark a struct which you send to the kernel,
				//      or any of the struct's subparts?
				// *CudafyCompileException: "error: data member initializer is not allowed"
				//   -> 
				// *CudafyCompileException: "error: linkage specification is not allowed"
				//   -> Functions allocating on device are to be [CudafyIgnore]d or at least not [Cudafy]d.
				//      Also, it more eagerly generates __global__ function definitions when your functions
				//      return void, so make them return int instead
				// *CudafyCompileException: "error: identifier "Inner" is undefined"
				//   -> Forgot [Cudafy] attribute on the struct Inner?
				// *CudafyCompileException: "error: identifier "DevInnerLen0" is undefined"
				//   -> Cudafy does not seem to handle lengths of inner things...???
				//      Make your own length variable. Nah, nvm that.
				string compileCommandLine = complete.Compile(eGPUCompiler.CudaNvcc, deleteGeneratedCode: false);
				complete.Serialize("complete"); // we serialize after compiling
			}
			
			return complete;
		}
	}


Dec 21, 2011 at 10:51 AM

Can you supply source code of the structs you are cudafying?  A quick test here gave no issues.