ManagedCuda and __constant__ variables

Question

I am using ManagedCuda in C#, and I have one question that I can't find an answer to... maybe you can help me. I was reading that in C++ and CUDA you can declare a variable (which is an array) like:

__constant__ double myVar[X]; (this is supossed to hold an array of X elements)

and then later use this to set the value from host code:

cudaMemcpyToSymbol(myVar, &arrayFromHost[0], sizeof(arrayFromHost) * numElements, 
                   size_t(0),cudaMemcpyHostToDevice);

so now you can use something like:

__global__ void myFunction(double *res)
{
    *res = myVar[0] + 2.5;
}

using a value that was set in myVar from the host...

but in ManagedCuda I don't seem to be able to do that... how can I do that??

1. Declare constant variable in my *.cu file
2. Set value (an array) from my *.cs file to that constant
3. Use the value from constant 1. in a function inside that same *.cu file

(or a __device__ variable... I don't know... it will be a variable that will receive an array (with a unknown number of elements) the first time it is run, and from then on, the function will reference its values, but that variable will never change)

Right now I only declare a CudaDeviceVariable and I don't touch it ever again, but on my kernel I always have to send the DevicePointer, which I think makes it harder to understand when reading...

Right now it looks something like this:

myKernel.Run(staticData.DevicePointer, moreData.DevicePointer, 
             evenMoreData.DevicePointer, numberOfElementsWhichNeverChange,            
             moreStaticData.DevicePointer, myResults.DevicePointer)

I would like to skip the 3 parameters that have data that never changes and set it in another function like setData.Run(numElements, staticData, moreStaticData);
and use from constant or device variables in other funcions in my *.cu file.

Answer 1

myKernel has a method SetConstantVariable() doing exactly what you want to do. Just call it before you launch the kernel: In your *.cu file:

extern "C" 
{
    __constant__ double myConstVarInCuda[5];
    __global__ void myFunction(double *res) 
    { 
        *res = myConstVarInCuda[0] + 2.5; 
    } 
}

In C#:

 double[] myVarInCS = new double[] { 1.0, 2.0, 3.0, 4.0, 5.0 };    
    myKernel.SetConstantVariable("myConstVarInCuda", myVarInCS);
    myKernel.Run(...);

If you don't declare your Cuda code in an extern "C" scope, be aware that names get mangled. In that case you can look up the exact mangled name in the PTX code.