CUDA integration with C#

// cuda wrapper function
extern "C" int __declspec(dllexport) __stdcall SomeCalculationsCU
      (
      float *a_h,  // pointer to input array
      const unsigned int N, // input array size
      const unsigned int M, // kernel M parameter
      const int cuBlockSize = 512, // kernel block size (max 512)
      const int showErrors = 1 // show CUDA errors in console window
      )
{
      int tmp = PRINT_ERRORS;
      PRINT_ERRORS = showErrors;
 
      float *a_d; // pointer to device array
      size_t size = N * sizeof(float);
      int cuerr = 0; // no errors
      unsigned int timer = 0;
   
      cudaMalloc((void**)&a_d, size); // allocate array on device   
      cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
   
      int n_blocks = N / cuBlockSize + (N % cuBlockSize == 0 ? 0 : 1);
   
      cutCreateTimer(&timer); // from cutil.h
      cutStartTimer(timer);
      some_calculations <<<n_blocks, cuBlockSize>>> (a_d, N, M);  // kernel invocation
      cudaThreadSynchronize(); // by default kernel runs in parallel with CPU code
      cutStopTimer(timer);
   
      cuerr = checkCUDAError("cuda kernel");
 
      cudaMemcpy(a_h, a_d, size, cudaMemcpyDeviceToHost);  
      if(!cuerr) cuerr = checkCUDAError("cuda memcpy");
     
      sExecutionTime = cutGetTimerValue(timer);
     
      cudaFree(a_d);
      if(!cuerr) cuerr = checkCUDAError("cuda free");
 
      PRINT_ERRORS = tmp;
      return cuerr;
}

The most important in that function is extern "C" int _declspec(dllexport) _stdcall part which makes it visible outside dll. The calling convention (_stdcall here) must be specified because by default C functions are using _cdecl one and .NET platform invoke use CallingConvention.Winapi which is _stdcall. More info about Win32 calling conventions is here http://unixwiz.net/techtips/win32-callconv.html. Generally important is to use the same convention for caller and callee functions and not use _stdcall with variadic ones.

This function could be also the main() function in console application project so you can modify it and add some print results section if you want or just add main() function like this one:

int main(void)
{
      float *a_h;
      const unsigned int N = 2000;
      const unsigned int M = 10;
      const int cublocks = 256;

      size_t size = N * sizeof(float);
      a_h = (float*)malloc(size);
      for(unsigned int i = 0; i < N; i++) a_h[i] = (float)i;
 
      SomeCalculationsCU(a_h, N, M, cublocks, 1);
 
      printf("exec time = %f ms\n", sExecutionTime);
}

To run this as an application (not build dll) change Configuration Type to Application (.exe) in project properties -> Configuration Properties -> General section. This also allows you to use CUDA Profiler (located in C:\CUDA\cudaprof\bin\ directory).

File dllmain.cu has also example of variable available to access from the outside of dll:

// external variable example
extern "C" { float __declspec(dllexport) sExecutionTime = -1; }

// variable wrapper function
extern "C" float __declspec(dllexport) __stdcall GetExecutionTime()
{
      return sExecutionTime;
}

It is possible to access that kind of variable directly (will be shown in C# part) but this is much easier with wrapper functions approach (could works like get/set accessors).

Finally example kernel function:

// cuda kernel (internal)
__global__ void some_calculations(float *a, unsigned int N, unsigned int M)
{
      unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
      if (idx < N)
      {
            // note1: no need for shared memory here
            // note2: global memory access is coalesced
            //        (no structs, float only used)
 
            // do computations M times on each thread
            // to extend processor's time
            for(unsigned int i = 0; i < M; i++)
            {
                  // some easy arithmetics           
                  a[idx] = a[idx] * a[idx] * 0.1 - a[idx] - 10;
            }
      }
}

And CPU version for execution time comparison:

extern "C" void __declspec(dllexport) __stdcall SomeCalculationsCPU
      (
      float *a_h,
      const unsigned int N,
      const unsigned int M
      )

      unsigned int timer = 0;
      cutCreateTimer(&timer);
      cutStartTimer(timer);
      for(unsigned int i = 0; i < N; i++)
            for(unsigned int j = 0; j < M; j++)
                  *(a_h + i) = *(a_h + i) * *(a_h + i) * 0.1 - *(a_h + i) - 10;
      cutStopTimer(timer);
      sExecutionTime = cutGetTimerValue(timer);
}

Part 3: Using DLL in C# code

This is easy part with .NET platform invoke functionality however one thing is worth mention - accessing variables because we cannot use DllImport attribute on them. We have to find address of that variable and then marshal data.

using System.Runtime.InteropServices;

#region hard way to import variable from unmanaged dll
[DllImport("kernel32.dll", SetLastError = true, CharSet = CharSet.Ansi)]
internal static extern IntPtr GetProcAddress(IntPtr hModule, string procName);
 
[DllImport("kernel32.dll", SetLastError = true, CharSet = CharSet.Ansi)]
internal static extern IntPtr LoadLibrary(string lpszLib);
 
static float ReadsExecutionTime()
{
    IntPtr hdl = LoadLibrary("cudalib.dll");
    if (hdl != IntPtr.Zero)
    {
        IntPtr addr = GetProcAddress(hdl, "sExecutionTime");
        if (addr != IntPtr.Zero)
        {
            //int value = Marshal.ReadInt32(addr);      // for integer types
            float[] managedArray = new float[1];        // single value
            Marshal.Copy(addr, managedArray, 0, 1);     // for other types
            return managedArray[0];
        }
    }
    return 0;
}
#endregion
 
// easy way to import variable from unmanaged dll - make wrapper if you can
[DllImport("cudalib.dll", CharSet = CharSet.Ansi, SetLastError = true, CallingConvention = CallingConvention.StdCall)]
public static extern float GetExecutionTime();

Things are little bit easier with integer types because you may use Marshal class like in commented line above. For other types Marshal.Copy() method is a possible solution. I used one-element array of float values to retrieve my variable. Marshal.Copy() overloads gives you also possibility to copy data from managed source to unmanaged destination (the opposite direction).

As for other functions from dll, including CUDA calculations function:

[DllImport("cudalib.dll", CharSet = CharSet.Ansi, SetLastError = true, CallingConvention = CallingConvention.StdCall)]
public static extern int SomeCalculationsCU(float[] a_h, uint N, uint M, int cuBlockSize, int showErrors);
[DllImport("cudalib.dll", CharSet = CharSet.Ansi, SetLastError = true, CallingConvention = CallingConvention.StdCall)]
public static extern void SomeCalculationsCPU(float[] a_h, uint N, uint M);

Again, remember to set calling convention properly and that is all.

The example code shows also how to use Parallel.For() method (one of possible ways to do that). This is a .NET 4 feature. More information on parallel programming in the .NET framework here http://msdn.microsoft.com/en-us/library/dd460693.aspx and examples here http://code.msdn.microsoft.com/ParExtSamples

using System.Threading.Tasks;
using System.Diagnostics;
private static double ParallelForVersion(float[] farr3, uint N, uint M)
{
    Stopwatch stp = new Stopwatch();
    stp.Start();
    Parallel.For(0, N, i =>
    {
        for (uint j = 0; j < M; j++)
            farr3[i] = farr3[i] * farr3[i] * 0.1f - farr3[i] - 10;
    });
    stp.Stop();
    return stp.Elapsed.TotalMilliseconds;
}

This code does exactly the same calculations as CUDA version however as you will see the calculation results are not exactly the same because of different error bounds for CPU and GPU arithmetic. There is also use_fast_math option for nvcc complier which force it to use less accurate but faster functions like __sinf(x) instead of sinf(x) not every function has its counterpart prefixed with __ (i.e. there is no fast double precision version of sin(x)).

Code example

The attached example shows execution times of easy computations on every array's element for CUDA (GPU), unmanaged native code (CPU), managed .NET code (Sequential) and .NET parallel versions of the same function.

The configuration parameters which you can change are size of the array, number of loops for each thread (more computations), CUDA block size and number of test rounds for average value.

Block size should be less than array size to utilize more multiprocessors. Good values are 256 or 512(max for most CUDA enabled devices). You can use CUDA profiler from CUDA SDK to check occupancy (should be 1 or close to 1). Values which are not multiply of 16 (half warp) will cause uncoalesced global memory operations and decrease in performance.

CUDA kernel here is very simple one without any thread synchronization and shared memory usage. In fact all threads run totally independent which produce great performance when there is a lot of computations for each array element.

You can check results with less loops – like 1 or 2 to see that even for big arrays (tasks) simple computations are faster on CPU because of time needed for copy operations. The usage of parallel for() is determined mostly by array size and is generally ineffective on small ones.

Here are results for my laptop with very old Nvidia GPU with compute capability 1.1:

Please see attachment with complete code for x86 and x64 platforms.