函數成員作爲CUDA內核的參數

Question

我正在使用動態並行機制，我想創建一個模板內核，給出一個對象指針+成員函數指針執行函數。這是一個最小（未）工作實例中，具有-arch = compute_35 -dlink標誌編譯，

#include <iostream> 

struct A 
{ 
    int i; 
    __device__ void clear() 
    { 
     i = 0; 
    } 
}; 

template<typename Object, typename memberFunction> 
__global__ void generalKernel(Object* o, memberFunction f) 
{ 
    (o->*f)(); 
} 

template<typename Object, typename memberFunction> 
__device__ void executeFunction(Object* o, memberFunction f) 
{ 
    generalKernel<<<1,1>>>(o,f); 
    cudaDeviceSynchronize(); 
} 

__global__ void mainKernel(A* a) 
{ 
    executeFunction(a, &A::clear); 
} 

int main(int argc, char * argv[]) 
{ 
    A* a; 
    cudaMallocManaged(&a, sizeof(A)); 
    a->i = 1; 

    mainKernel<<<1,1>>>(a); 
    cudaDeviceSynchronize(); 

    std::cout << a->i << std::endl; 

    return EXIT_SUCCESS; 
} 

Answer 1

下面是一個簡單CUDA代碼，以顯示如何通過成員函數指針到內核。一切都在代碼中解釋。

#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); 
    } 
} 


// struct holds an 'int' type data memeber and '__device__' function member 
struct ST 
{ 
    int id; 
    __device__ void foo() 
    { 
     printf("value of id: %d\n",id); 
    } 
}; 

// creating an alias for our function pointer 
// since the function is a member of a struct, we add struct name and scope resolution 'ST::' 
// to signify as such 
typedef void (ST::*Fptr)(void); 

// templated kernel 
template<typename Object, typename memberFunction> 
__global__ void kernel(Object* o, memberFunction f) 
{ 
    (o->*f)(); 
} 

// declaring a __device__ function pointer, assigning it the address of 'ST::foo' 
// remember that this function pointer is also direclty accessible from the kernel 
__device__ Fptr fp = &ST::foo; 


int main(int argc, char** argv) 
{ 
    // declaring and initializing a host 'ST' object 
    ST h_st; 
    h_st.id = 10; 


    // device 'ST' object 
    ST* d_st; 
    // allocating device memory 
    gpuErrchk(cudaMalloc((void**)&d_st, sizeof(ST))); 
    // copying host data from host object to device object 
    gpuErrchk(cudaMemcpy(d_st, &h_st, sizeof(ST), cudaMemcpyHostToDevice)); 

    // declaring host side function pointer of type 'Fptr', which can be passed to kernel as argument 
    Fptr h_fptr; 
    // copying address of '__device__' function pointer to host side function pointer 
    gpuErrchk(cudaMemcpyFromSymbol(&h_fptr, fp, sizeof(Fptr))); 


    // passing arguments to kernel 
    kernel<<<1,1>>>(d_st,h_fptr); 

    // making sure no errors occured 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 

    // free device memory 
    gpuErrchk(cudaFree(d_st)); 

    return 0; 
}