CUDA中统一内存的函数指针分配

Posted 2023-03-23

【中文标题】CUDA中统一内存的函数指针分配

【英文标题】：Assignment of function pointer with the unified memory in CUDA

【发布时间】：2017-07-06 21:59:23

【问题描述】：

我正在尝试在方便的统一内存模型下实现函数与CUDA的动态绑定。在这里，我们有一个结构体 Parameters，其中包含一个成员函数指针 void (*p_func)()。

#include <cstdio>

struct Parameters 
    void (*p_func)();
;

结构体由统一内存管理，我们将实际函数func_A分配给p_func。

__host__ __device__
void func_A() 
    printf("func_A is correctly invoked!\n");
    return;

当我们通过下面的代码时，问题就出现了：如果赋值1运行，即para->p_func = func_A，实际上设备和主机函数地址都是由函数地址在主人。相反，如果分配 2 运行，则地址都将成为设备 1。

__global__ void assign_func_pointer(Parameters* para) 
    para->p_func = func_A;


__global__ void run_on_device(Parameters* para) 
    printf("run on device with address %p\n", para->p_func);
    para->p_func();


void run_on_host(Parameters* para) 
    printf("run on host with address %p\n", para->p_func);
    para->p_func();


int main(int argc, char* argv[]) 

    Parameters* para;
    cudaMallocManaged(&para, sizeof(Parameters));

    // assignment 1, if we uncomment this section, p_func points to address at host
    para->p_func = func_A;
    printf("addr@host: %p\n", para->p_func);

    // assignment 2, if we uncomment this section, p_func points to address at device
    assign_func_pointer<<<1,1>>>(para); // 
    cudaDeviceSynchronize();
    printf("addr@device: %p\n", para->p_func);

    run_on_device<<<1,1>>>(para);
    cudaDeviceSynchronize();

    run_on_host(para);

    cudaFree(para);
    return 0;

现在的问题是，在统一内存模型下，设备和主机上的函数指针是否可以分别指向正确的函数地址？

【问题讨论】：

在你的结构体中放置两个函数指针，一个指向主机，一个指向设备，并根据上下文调度适当的函数。

@RobertCrovella

@RobertCrovella 如果一个变量不起作用，请添加另一个！是的，这是一个切实可行的解决方案。

【参考方案1】：

暂且不谈统一内存的技术问题，您的问题实际上是“一个变量能否同时具有两个不同的值？”答案显然是否定的。

更详细地说：CUDA 统一内存从根本上确保给定的托管分配在从主机和设备访问时将具有一致的值（在某些约束下）。你要求的是完全相反的，它显然不受支持。

【讨论】：

我想知道CUDA是否可以提供一种巧妙处理函数指针的机制，它不是一个普通的变量，而是一个特殊的变量，只能在特定于设备的条件下使用。正如@RobertCrovella 所提到的，现在似乎必须求助于两个函数指针来托管不同的地址。谢谢！

关键是函数指针一点都不特别。它们像任何其他指针一样是 juat 值。 CUDA 中没有自省，因此运行时无法知道值是函数指针

【参考方案2】：

通过对struct 定义的一些修改，可能会出现这样的情况：

$ cat t1288.cu
#include <cstdio>

struct Parameters 
    void (*p_hfunc)();
    void (*p_dfunc)();
    __host__ __device__
    void p_func()
      #ifdef __CUDA_ARCH__
      (*p_dfunc)();
      #else
      (*p_hfunc)();
      #endif
      
;

__host__ __device__
void func_A() 
    printf("func_A is correctly invoked!\n");
    return;


__global__ void assign_func_pointer(Parameters* para) 
    para->p_dfunc = func_A;


__global__ void run_on_device(Parameters* para) 
    printf("run on device\n"); // with address %p\n", para->p_dfunc);
    para->p_func();


void run_on_host(Parameters* para) 
    printf("run on host\n"); // with address %p\n", para->p_func);
    para->p_func();


int main(int argc, char* argv[]) 

    Parameters* para;
    cudaMallocManaged(&para, sizeof(Parameters));

    // assignment 1, if we uncomment this section, p_func points to address at host
    para->p_hfunc = func_A;
    printf("addr@host: %p\n", para->p_hfunc);

    // assignment 2, if we uncomment this section, p_func points to address at device
    assign_func_pointer<<<1,1>>>(para); //
    cudaDeviceSynchronize();
    printf("addr@device: %p\n", para->p_dfunc);

    run_on_device<<<1,1>>>(para);
    cudaDeviceSynchronize();
    run_on_host(para);

    cudaFree(para);
    return 0;

$ nvcc -arch=sm_35 -o t1288 t1288.cu
$ cuda-memcheck ./t1288
========= CUDA-MEMCHECK
addr@host: 0x402add
addr@device: 0x8
run on device
func_A is correctly invoked!
run on host
func_A is correctly invoked!
========= ERROR SUMMARY: 0 errors
$

我同意另一个答案，即即使使用托管内存，目前也不可能拥有一个在主机代码和设备代码中都能正常工作的数字函数指针。

【讨论】：

除了函数指针之外，结构体中引入的包装函数使解决方案更加优雅。很好的解释！