Назначение указателя функции с унифицированной памятью в CUDA

Question

Я пытаюсь реализовать динамическое связывание функций с CUDA под удобной унифицированной моделью памяти. Здесь у нас есть структура Параметры, содержащие элемент, указатель функции 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 работает, то есть, пара-> p_func = func_A, оба устройства и хост-адрес функции фактически назначается адрес функции на хост. В отличие от этого, если назначение 2 работает, адреса становятся одним из устройств.

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

Я согласен с другой ответ, что в настоящее время не представляется возможным даже с управляемой памяти, чтобы иметь единый числовой указатель на функцию который корректно работает как в коде хоста, так и в коде устройства.