分配常量内存

Question

我正在尝试将模拟参数设置为常量内存，但是没有成功（CUDA.NET）。 cudaMemcpyToSymbol函数返回cudaErrorInvalidSymbol。 cudaMemcpyToSymbol中的第一个参数是字符串...它是符号名称吗？实际上，我不太理解它是如何被解析的。寻求任何帮助。

//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);

我的 .cu 文件包含以下内容：

__constant__ float param;

解决方案

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

- Vladimir

Answer 1

很遗憾，__ constant __必须与memcpy到符号的文件作用域相同，在你的情况下，你的__ constant __在一个单独的.cu文件中。

解决这个问题的简单方法是在你的.cu文件中提供一个包装函数，例如：

__constant__ float param;

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

// etc.
__global__ void ...

- Tom

---

感谢您的想法，对于cpp（--compile）是清晰的并且有效的。但是，在.net应用程序的后期构建事件中，对于nvcc myfile.cu --cubin则不适用。"(-cubin)将所有.cu/.ptx/.gpu输入文件编译为仅设备的.cubin文件。此步骤会丢弃每个.cu输入文件的主机代码。"因此，当我加载此模块时，它不包括主机函数。 - Vladimir

Answer 2

如果这个问题是实际存在的，您可以使用 cuModuleGetGlobal 然后像这样使用 cudaMemcpy：

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;
    }
}

在NVIDIA GPU Computing SDK 3.1和CUDA.NET 3.0中，您可以使用以下代码加载模块：
CUmodule module = cuda.LoadModule("MyCode.cubin");

- serg_s

---

感谢您的代码。这看起来是从驱动程序API访问常量内存的好方法...（例如ffpeg .c文件）...这本Google图书提到了使用您所做的那些函数：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

虽然那本书建议使用cuMemcpyHtoD和cuMemcpyDtoH函数进行复制。 - andrew pate

Answer 3

常量内存具有隐式的局部作用域链接。 请确保声明与使用它的文件在同一文件中。听起来你有两个文件。 可能还需要将param声明为数组（或者也可能不需要）。

- Anycorn