장치 측 CUDA 가상 함수를 구현하는 방법은 무엇입니까?

Question

CUDA는 가상 함수가있는 클래스가 커널 함수로 전달되는 것을 허용하지 않습니다. 이 제한 사항에 대한 해결 방법이 있습니까?

커널 기능 내에서 다형성을 사용하고 싶습니다.

감사합니다.

Answer 1

로버트 Crovella의 코멘트의 가장 중요한 부분은 다음과 같습니다

객체는 단순히 디바이스에 생성 될 필요가있다.

그런 점을 염두에 두면서, 나는 추상적 인 class Function을 가진 상황을 다루었으며, 그 다음에 다른 기능과 평가를 캡슐화 한 일부 구현을 다루었습니다.

class Function 
{ 
public: 
    __device__ Function() {} 
    __device__ virtual ~Function() {} 
    __device__ virtual void Evaluate(const real* __restrict__ positions, real* fitnesses, const SIZE_TYPE particlesCount) const = 0; 
}; 

class FunctionRsj : public Function 
{ 
private: 
    SIZE_TYPE m_DimensionsCount; 
    SIZE_TYPE m_PointsCount; 
    real* m_Y; 
    real* m_X; 
public: 
    __device__ FunctionRsj(const SIZE_TYPE dimensionsCount, const SIZE_TYPE pointsCount, real* configFileData) 
     : m_DimensionsCount(dimensionsCount), 
      m_PointsCount(pointsCount), 
      m_Y(configFileData), 
      m_X(configFileData + pointsCount) {} 

    __device__ ~FunctionRsj() 
    { 
     // m_Y points to the beginning of the config 
     // file data, use it for destruction as this 
     // object took ownership of configFilDeata. 
     delete[] m_Y; 
    } 

    __device__ void Evaluate(const real* __restrict__ positions, real* fitnesses, const SIZE_TYPE particlesCount) const 
    { 
     // Implement evaluation of FunctionRsj here. 
    } 
}; 

__global__ void evaluate_fitnesses(
    const real* __restrict__ positions, 
    real* fitnesses, 
    Function const* const* __restrict__ function, 
    const SIZE_TYPE particlesCount) 
{ 
    // This whole kernel is just a proxy as kernels 
    // cannot be member functions. 
    (*function)->Evaluate(positions, fitnesses, particlesCount); 
} 

__global__ void create_function(
    Function** function, 
    SIZE_TYPE dimensionsCount, 
    SIZE_TYPE pointsCount, 
    real* configFileData) 
{ 
    // It is necessary to create object representing a function 
    // directly in global memory of the GPU device for virtual 
    // functions to work correctly, i.e. virtual function table 
    // HAS to be on GPU as well. 
    if (threadIdx.x == 0 && blockIdx.x == 0) 
    { 
     (*function) = new FunctionRsj(dimensionsCount, pointsCount, configFileData); 
    } 
} 

__global__ void delete_function(Function** function) 
{ 
    delete *function; 
} 

int main() 
{ 
    // Lets just assume d_FunctionConfigData, d_Positions, 
    // d_Fitnesses are arrays allocated on GPU already ... 

    // Create function. 
    Function** d_Function; 
    cudaMalloc(&d_Function, sizeof(Function**)); 
    create_function<<<1, 1>>>(d_Function, 10, 10, d_FunctionConfigData); 

    // Evaluate using proxy kernel. 
    evaluate_fitnesses<<< 
     m_Configuration.GetEvaluationGridSize(), 
     m_Configuration.GetEvaluationBlockSize(), 
     m_Configuration.GetEvaluationSharedMemorySize()>>>(
     d_Positions, 
     d_Fitnesses, 
     d_Function, 
     m_Configuration.GetParticlesCount()); 

    // Delete function object on GPU. 
    delete_function<<<1, 1>>>(d_Function); 
}