Cuda에서 표현 템플릿을 사용하여 람다 식 작성하기

Question

CUDA 커널에 임의의 함수를 전달해야하는 CUDA 응용 프로그램을 개발 중입니다. 가능한 상황마다 함수 포인터를 선언하고 커널에 넘겨 주면 너무 많은 번거 로움 (> 50 개의 다른 함수)이 될 것이고 그 중 모두는 sin(x)/y과 같은 기본 함수로 구성되어 있으므로 최소한의 람다 CUDA 커널에 대한 표현 기능. C++ 11 기능은 아직까지는 장치 코드에 의해 지원되지 않기 때문에 (필자가 아는 한) 관련 정보를 온라인으로 찾지 못했기 때문에 표현 템플릿을 직접 가르치고 커널에 전달할 간단한 람다 표현 규칙을 구현하기로 결정했습니다. .

다음 코드는 NVCC에서 컴파일되고 잘 실행되는 최소 구현 코드입니다. 그러나이 경로를 따라 가면 나는 1 개의 변수로만 함수를 구현할 수 있습니다. 내 코드를 확장하여 sin(_x) + _y과 같은 함수 합성을 처리 할 수있는 방법이 있습니까?

미리 감사드립니다.

#include<math.h> 

#ifdef __CUDACC__ 
#define HOST_DEVICE __host__ __device__ 
#else 
#define HOST_DEVICE 
#endif 

struct Id {}; 

template <typename Op, typename Left, typename Right> 
struct BinaryOp 
{ 
    Left left; 
    Right right; 
    HOST_DEVICE BinaryOp(Left t1, Right t2) : left(t1), right(t2) {} 

    HOST_DEVICE double operator() (double x) { 
     return Op::apply(left(x), right(x)); 
    } 
}; 

template <typename Op, typename Arg> 
struct UnaryOp 
{ 
    Arg arg; 
    HOST_DEVICE UnaryOp(Arg t1) : arg(t1) {} 

    HOST_DEVICE double operator() (double x) { 
     return Op::apply(arg(x)); 
    } 
}; 

template <> 
struct UnaryOp<Id, double> 
{ 
    HOST_DEVICE UnaryOp() {} 
    HOST_DEVICE double operator() (double x) { 
     return x; 
    } 
}; 

struct Sin 
{ 
    HOST_DEVICE static double apply(double x) { 
     return sin(x); 
    } 
}; 

struct Plus 
{ 
    HOST_DEVICE static double apply(double a, double b) { 
     return a + b; 
    } 
}; 

template <typename Left, typename Right> 
BinaryOp<Plus, Left, Right> operator+ (Left lhs, Right rhs) { 
    return BinaryOp<Plus, Left, Right>(lhs, rhs); 
} 

template <typename Arg> 
UnaryOp<Sin, Arg> _sin(Arg arg) { 
    return UnaryOp<Sin, Arg>(arg); 
} 

template <class T> 
__global__ void test(T func, double x) { 
    printf("%e\n", func(x)); 
} 

int main() 
{ 
    UnaryOp<Id, double> _x; 
    double x = 1.0; 
    test<<<1, 1>>>(_sin(_x) + _x, x); 
    cudaDeviceSynchronize(); // Needed or the host will return before kernel is finished 
    return 0; 
} 

Answer 1

그래서 질문을하고 해킹 한 후 해결책을 찾았습니다. 못생긴지만 그것은 나 자신을 위해 일합니다. 다음은 최대 3 개의 자유 변수를 지원하는 수정 된 코드입니다. 더 많은 변수를 하드 코딩 할 수 있지만 현재 프로젝트가 필요하지 않습니다.

#include<math.h> 

#ifdef __CUDACC__ 
#define HOST_DEVICE __host__ __device__ 
#else 
#define HOST_DEVICE 
#endif 

struct Id {}; 

template <typename Op, typename Left, typename Right> 
struct BinaryOp 
{ 
    Left left; 
    Right right; 
    HOST_DEVICE BinaryOp(Left t1, Right t2) : left(t1), right(t2) {} 

    HOST_DEVICE double operator() (double x1, double x2 = 0.0, double x3 = 0.0) { 
     return Op::apply(left(x1, x2, x3), right(x1, x2, x3)); 
    } 
}; 

template <typename Op, typename Arg> 
struct UnaryOp 
{ 
    Arg arg; 
    HOST_DEVICE UnaryOp(Arg t1) : arg(t1) {} 

    HOST_DEVICE double operator() (double x1, double x2 = 0.0, double x3 = 0.0) { 
     return Op::apply(arg(x1, x2, x3)); 
    } 
}; 

template <int argnum> 
struct Var 
{ 
    HOST_DEVICE Var() {} 
    HOST_DEVICE double operator() (double x1, double x2 = 0.0, double x3 = 0.0) { 
     if (1 == argnum) return x1; 
     else if (2 == argnum) return x2; 
     else return x3; 
    } 
}; 

struct Sin 
{ 
    HOST_DEVICE static double apply(double x) { 
     return sin(x); 
    } 
}; 

struct Plus 
{ 
    HOST_DEVICE static double apply(double a, double b) { 
     return a + b; 
    } 
}; 

template <typename Left, typename Right> 
BinaryOp<Plus, Left, Right> operator+ (Left lhs, Right rhs) { 
    return BinaryOp<Plus, Left, Right>(lhs, rhs); 
} 

template <typename Arg> 
UnaryOp<Sin, Arg> _sin(Arg arg) { 
    return UnaryOp<Sin, Arg>(arg); 
} 

template <class T> 
__global__ void test(T func, double x, double y, double z = 0.0) { 
    printf("%e\n", func(x, y)); 
} 

Var<1> _x; 
Var<2> _y; 

int main() 
{ 
    test<<<1, 1>>>(_sin(_x) + _y, 1.0, 2.0); 
    cudaDeviceSynchronize(); // Needed or the host will return before kernel is finished 
    return 0; 
} 

이것은 분명히 추한 해킹입니다. 람다 식은 double (또는 double으로 변환 할 수있는 형식)에서만 작동합니다. 그러나 나는 그 순간을 극복 할 방법을 상상할 수 없다. NVCC는 더 이상 이런 종류의 해킹이 필요하지 않도록 C++ 11 기능을 곧 지원할 수 있기를 바랍니다.

누군가가 나에게 라이브러리 또는 더 나은 해킹 방법을 제시 할 수 있다면 큰 도움이 될 것입니다. 어떤 도움을 주셔서 감사합니다!