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

Question

我正在尝试在方便的统一内存模型下实现函数与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;
}

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

Answer 1

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

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

Answer 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
$

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