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??

Declare constant variable in my *.cu file
Set value (an array) from my *.cs file to that constant
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.

Using a string to identify a symbol (i.e. `cudaMemcpyToSymbol("sym",...)`) was deprecated in CUDA 4.1 and removed in 5.0, instead you should use the actual symbol (`cudaMemcpyToSymbol(sym,...)`). That doesn't answer the managedCUDA question, just an FYI. — Tom, Sep 13 '12 at 09:19

kunzmi · Accepted Answer · 2012-09-14T15:51:15.807

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.

ManagedCuda and __constant__ variables

1 Answers1

ManagedCuda and constant variables