0

我最近尝试使用函数指针在我的应用程序中动态定义几个处理阶段，在 sm_30 上运行。

在这里发布代码会很困难，因为涉及到许多不同的文件和函数，但基本上，我从 Cuda Toolkit 5.0 中包含的示例开始。

我分配了一个设备函数缓冲区，在其中复制了一个设备函数指针，正如在示例中定义的那样，这要归功于与 DeviceToDevice 复制类型一起使用的 cudaMemcpyfromsymbolAsync。

我的设备指针在 .cu.h 中定义如下：

//device function pointer model
typedef void (*func)(structGpuArgument*);

//Declaring a function
__device__ void gpuFunc1(structGpuArgument* arg1);

在其他地方我有一个 .cu ，其中包含包含以下代码的先前声明：

//get the actual function pointer
__device__ func gpuFuncPtr = gpuFunc1;

//Buffer to store a list of function pointer
func* pFuncDevBuffer;
cudaMalloc(&pFuncDevBuffer,NB_FUNC*sizeof(func));

//copy the actual function pointer (symbol) to the list buffer 
cudaMemcpyFromSymbolAsync( pFuncDevBuffer+i ,gpuFuncPtr,sizeof(func),0,cudaMemcpyDeviceToDevice,stream)

//Launch the kernel that will use the functions
kernel_test<<<1,10,0,stream>>>(pFuncDevBuffer)
...

//defining the kernel that uses pointer buffer
__global__ void kernel_test(func* pFuncDevBuffer)
{
   printf("func address : %p\n",pFuncDevBuffer[0]);
   pFuncDevBuffer[0](NULL);
}

//defining the function pointed by the function pointer
__device__ void gpuFunc1(structGpuArgument* arg1)
{
     do_something;
}

事实上，只要将设备函数缓冲区作为参数的全局内核定义在与函数及其指针相同的文件中，一切都会正常工作。然后内核可以打印出函数的地址（0x4）并毫无问题地执行其代码我不使用单独编译。

当在程序的同一个实例中，在别处定义的第二个内核在参数中采用完全相同的函数指针缓冲区时，它可以为函数指针（0x4）打印出完全相同的内存地址，但如果它试图执行它，它在 cuda-memcheck 的 0x00000000 处发出非法指令失败。之后任何其他 cuda API 调用都会冻结，我需要重新启动计算机（我的 gpu 不支持通过 cuda-smi 重置）。

我想知道以这种方式使用函数指针是否存在已知问题，即使用在其他文件中定义的函数指针缓冲区，但共享相同的函数指针定义。

此外，如果在段错误后重置设备而不重新启动整个系统，它可以帮助我在调试应用程序时节省时间。

谢谢您的帮助

4

0 回答 0