Allocate constant memory

Question

I'm trying to set my simulation params in constant memory but without luck (CUDA.NET). cudaMemcpyToSymbol function returns cudaErrorInvalidSymbol. The first parameter in cudaMemcpyToSymbol is string... Is it symbol name? actualy I don't understand how it could be resolved. Any help appreciated.

//init, load .cubin   
float[] arr = new float[1];
    arr[0] = 0.0f;
    int size = Marshal.SizeOf(arr[0]) * arr.Length;
    IntPtr ptr = Marshal.AllocHGlobal(size);
    Marshal.Copy(arr, 0, ptr, arr.Length);
    var error = CUDARuntime.cudaMemcpyToSymbol("param", ptr, 4, 0, cudaMemcpyKind.cudaMemcpyHostToDevice);

my .cu file contain

__constant__ float param;

Working solution

     cuda.LoadModule(Path.Combine(Environment.CurrentDirectory, "name.cubin"));            
 simParams = cuda.GetModuleGlobal("params");
 float[] parameters = new float[N]{...}             
 cuda.CopyHostToDevice<float>(simParams, parameters);

score 5 · Answer 1 · answered Mar 16 '10 at 08:00

5

Unfortunately the __ constant __ must be in the same file scope as the memcpy to the symbol, and in your case your __ constant __ is in a separate .cu file.

The simple way around this is to provide a wrapper function in your .cu file, for example:

__constant__ float param;

// Host function to set the constant
void setParam(float value)
{
  cudaMemcpyToSymbol("param", ptr, 4, 0, cudaMemcpyHostToDevice);
}

// etc.
__global__ void ...

answered Mar 16 '10 at 08:00

Tom

20,852
4
42
54

Thanks idea is clear and works for cpp (--compile). But not in case of nvcc myfile.cu --cubin in post build event of .net app. "(-cubin) Compile all .cu/.ptx/.gpu input files to device-only .cubin files. This step discards the host code for each .cu input file." So when I load this module it doesn't include host functions. – Vladimir Mar 16 '10 at 19:57

score 2 · Answer 2 · answered Jan 14 '11 at 06:59

If this question is actual you can use cuModuleGetGlobal and next cudaMemcpy like this:

private bool setValueToSymbol(CUmodule module, string symbol, int value)
{
    CUdeviceptr devPtr = new CUdeviceptr();
    uint lenBytes = 0;
    CUResult result = CUDADriver.cuModuleGetGlobal(ref devPtr, ref lenBytes, module, symbol);
    if (result == CUResult.Success)
    {
        int[] src = new int[] { value };
        cudaError error = CUDARuntime.cudaMemcpy(devPtr, src, lenBytes, cudaMemcpyKind.cudaMemcpyHostToDevice);
        if (error == cudaError.cudaSuccess)
            return true;
        else
            return false;
    }
    else
    {
        return false;
    }
}

where CUmodule module = cuda.LoadModule("MyCode.cubin"); This code works with NVIDIA GPU Computing SDK 3.1 and CUDA.NET 3.0.

Thanks for the code. This looks a good way to access constant memory from the driver API... (for example an ffpeg .c file) ... This google book mentions using those functions as you have done : https://books.google.co.uk/books?id=ynydqKP225EC&pg=PA157&lpg=PA157&dq=cuModuleGetGlobal+constant+memory&source=bl&ots=5l16S_QDyy&sig=ACfU3U20XxrwqFIn2XbkktsWzt_zuB8qJg&hl=en&sa=X&ved=2ahUKEwiu05uBz-_pAhW1RBUIHTDvCRkQ6AEwAHoECAsQAQ#v=onepage&q=cuModuleGetGlobal%20constant%20memory&f=false — andrew pate, Jun 07 '20 at 11:56
Although that book suggests cuMemcpyHtoD and cuMemcpyDtoH as the functions to use for the copy. — andrew pate, Jun 07 '20 at 12:15

score 1 · Answer 3 · answered Mar 16 '10 at 01:26

1

constant memory has implicit local scope linkage. make sure declaration is in the same file where you use it. it sounds like you have two files. may also have to declare param to array (or maybe not)

answered Mar 16 '10 at 01:26

Anycorn

50,217
42
167
261

Allocate constant memory

3 Answers3

Linked