デバイス関数へのポインタをカーネル関数の引数として渡すには？

Question

この質問は、CUDA C/C++プログラミングに関するものです。私はそれをかなりたくさん検索しようとしましたが、適切な質問を見つけられなかったので、答えもありませんでした。
I 1つの装置の機能、1つのカーネル関数と主な機能を有する：CUDA programming guideから

typedef float (*pfunc)(float arg); 

__device__ float dev_func(float arg) { 
    return arg * arg; 
} 

__global__ void ker_func(pfunc fnc) { 
    printf("%f\n", fnc(2)); 
} 

int main(void) { 
    pfunc fnc = dev_func; 
    //now how do I copy this pointer to device memory? 
    ker_func<<<1,1>>>(...); 
    return 0; 
} 

Answer 1

を：

ホスト・コードに取り込ま__global__関数のアドレスがデバイス・コードに使用することができません（カーネルを起動するなど）。同様に、デバイスコードで取り込まれた__global__関数のアドレスは、ホストコードで使用することはできません。

ホストコードに__device__機能のアドレスを使用することはできません。

グローバルに__device__関数ポインタを定義し、カーネルにそれを呼び出す：

だから次の2つのオプションをしました。

typedef float (*pfunc)(float arg); 

__device__ float dev_func(float arg) { 
    return arg * arg; 
} 

// create device function pointer here 
__device__ pfunc dev_func_ptr = dev_func; 

__global__ void ker_func() { 
    // call function through device function pointer 
    printf("%f\n", dev_func_ptr(2)); 
} 

あなたは、引数としてカーネルに関数ポインタを渡したい場合は、次の

#define gpuErrchk(val) \ 
    cudaErrorCheck(val, __FILE__, __LINE__, true) 
void cudaErrorCheck(cudaError_t err, char* file, int line, bool abort) 
{ 
    if(err != cudaSuccess) 
    { 
     printf("%s %s %d\n", cudaGetErrorString(err), file, line); 
     if(abort) exit(-1); 
    } 
} 

typedef float (*pfunc)(float arg); 

__device__ float dev_func(float arg) { 
    return arg * arg; 
} 

// create device function pointer here 
__device__ pfunc dev_func_ptr = dev_func; 

__global__ void ker_func(pfunc fnc) { 
    // call function through device function pointer 
    printf("%f\n", fnc(2)); 
} 


int main(int argc, char** argv) 
{ 
    // create a host function pointer 
    pfunc host_function_ptr; 
    // copy function pointer value from device to host 
    gpuErrchk(cudaMemcpyFromSymbol(&host_function_ptr, dev_func_ptr, sizeof(pfunc))); 
    // pass the copied function pointer in kernel 
    ker_func<<<1,1>>>(host_function_ptr); 

    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 

    return 0; 
}