Указатели на функции С CUDA функций устройства

Question

Я хотел бы использовать указатели на функции в моем коде Cuda C++, как показано ниже,

typedef __device__ void customFunc(const char*, uint64_t, char*, const uint64_t);

который, что я после. Его эквивалент без «__device__» работает отлично.

Поддерживаются ли указатели на функции в Cuda?

Edit:

Я специально заинтересован в том, чтобы использовать __device__ функции, как функции указателей на функции __device__

Answer 1

Там нет никакой магии, связанные с использованием функций устройства указателей в коде устройства. Он функционально и синтаксически идентичен стандарту C++.

Например:

#include <cstdio> 

typedef int (*ufunc)(int args); 

__device__ int f1(int x) 
{ 
    int res = 2*x; 
    printf("f1 arg = %d, res = %d\n", x, res); 
    return res; 
} 

__device__ int f2(int x, int y, ufunc op) 
{ 
    int res = x + op(y); 
    printf("f2 arg = %d, %d, res = %d\n", x, y, res); 
    return res; 
} 


__global__ void kernel(int *z) 
{ 

    int x = threadIdx.x; 
    int y = blockIdx.x; 
    int tid = threadIdx.x + blockDim.x * blockIdx.x; 

    z[tid] = f2(x, y, &f1); 
} 

int main() 
{ 
    const int nt = 4, nb = 4; 
    int* a_d; 
    cudaMalloc(&a_d, sizeof(float) * nt *nb); 

    kernel<<<nb, nt>>>(a_d); 
    cudaDeviceSynchronize(); 
    cudaDeviceReset(); 

    return 0; 
} 
#include <cstdio> 

typedef int (*bfunc)(int args); 

__device__ int f1(int x) 
{ 
    int res = 2*x; 
    printf("f1 arg = %d, res = %d\n", x, res); 
    return res; 
} 

__device__ int f2(int x, int y, bfunc op) 
{ 
    int res = x + f1(y); 
    printf("f2 arg = %d, %d, res = %d\n", x, y, res); 
    return res; 
} 


__global__ void kernel(int *z) 
{ 

    int x = threadIdx.x; 
    int y = blockIdx.x; 
    int tid = threadIdx.x + blockDim.x * blockIdx.x; 

    z[tid] = f2(x, y, &f1); 
} 

int main() 
{ 
    const int nt = 4, nb = 4; 
    int* a_d; 
    cudaMalloc(&a_d, sizeof(float) * nt *nb); 

    kernel<<<nb, nt>>>(a_d); 
    cudaDeviceSynchronize(); 
    cudaDeviceReset(); 

    return 0; 
} 

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

Важно иметь в виду, что в CUDA это не законно включить CUDA спецификаторов (__device__, __constant__, __global__ и т.д.) в определениях типа. Каждый экземпляр переменной имеет спецификатор как часть его определения.