CUDA 常量内存问题：使用 cudaGetSymbolAddress 的设备符号无效

Question

我正在尝试在启动需要这些值的内核之前在我的 GPU 的常量内存上设置常量值。

我的代码（简化）：

__constant__ size_t con_N;

int main() 
{

    size_t N;
    size_t* dev_N = NULL;
    cudaError_t cudaStatus;

    //[...]

    cudaStatus = cudaGetSymbolAddress((void **)&dev_N, &con_N);
    if (cudaStatus != cudaSuccess) {
        cout<<"cudaGetSymbolAddress (dev_N) failed: "<<cudaGetErrorString(cudaStatus)<<endl;
    }

我打算之后cudaMemcpy我的N到dev_N。

但是，此时我在代码中得到的只是：

cudaGetSymbolAddress (dev_N) failed: invalid device symbol

我正在使用 CUDA 6.5，所以这不是引用符号问题，因为它是我迄今为止检查的大部分问答中的问题。

我尝试用con_N[1] 替换con_N（并在cudaGetSymbolAddress 参数中删除con_N 之前的&）：结果相同。

由于这个函数的原型是cudaGetSymbolAddress(void **devPtr , const void* symbol )，我猜它想得到我的符号地址。但是，我尝试使用 cudaStatus = cudaGetSymbolAddress((void **)&dev_N, (const void*) con_N); 并收到相同的消息。

当我删除 cudaGetSymbolAddress((void **)&dev_N, &con_N) 并直接使用 cudaMemcpyToSymbol(&con_N, &N, sizeof(size_t)) 时，我也会收到相同的错误消息。

恐怕我错过了一些重要的事情。任何帮助将不胜感激。

Answer 1

cudaGetSymbolAddress的正确用法是

cudaGetSymbolAddress((void **)&dev_N, con_N)

我将通过下面的简单示例来展示这一点。

如文档所述，该符号应实际驻留在设备上。因此，在 API 调用中使用 &con_N 似乎毫无意义，因为作为主机 API 的 cudaGetSymbolAddress，直接从主机访问驻留在设备上的某物的地址应该是不可能的。我不确定出现在 CUDA Runtime API 文档中的原型是否应该更好地阅读为 `

template<class T>
cudaError_t cudaGetSymbolAddress (void **devPtr, const T symbol)

用设备符号引用代替设备符号地址。

#include <stdio.h>

__constant__ int const_symbol;

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

/***************/
/* TEST KERNEL */
/***************/
__global__ void kernel() {

    printf("Address of symbol from device = %p\n", &const_symbol);

}

/********/
/* MAIN */
/********/
int main() 
{
    const int N = 16;
    int *pointer = NULL;

    gpuErrchk(cudaGetSymbolAddress((void**)&pointer, const_symbol));

    kernel<<<1,1>>>();

    printf("Address of symbol from host = %p\n", pointer);

    return 0;
}

Answer 2

在我看来，你的一行代码应该像下面这样固定。

cudaStatus = cudaGetSymbolAddress((void **)&dev_N, con_N);

希望对你有所帮助。