GPU의 CUDA 커널 스케줄러

Question

저는 CUDA 커널 스케줄러를 쓰고 있습니다. 스케줄러는 Task 포인터의 벡터를 가져 와서 실행시킵니다. 포인터는 KernelTask 다른 유형 매개 변수의 객체를 가리키며 임의의 매개 변수가있는 커널을 지원합니다.

CPU 버전의 스케줄러와 GPU 버전이 있습니다. CPU 버전은 정상적으로 작동합니다. 커널을 실행하기 위해 가상 함수 Task::start을 호출합니다. GPU 버전에는 세 가지 문제가 있습니다.

1. 가상 함수는 CUDA에서 사용할 수 없습니다. 다운 캐스팅없이 어떻게 피할 수 있습니까?
2. std :: get는 호스트 함수입니다. 표준을 구현할 수있는 방법이 있나요 :: 자신을 GPU 용으로 만들 수 있습니까?
3. (낮은 우선 순위) KernelTask 개체의 크기가 다르기 때문에 copyToGPU()을 사용하여 모든 개체를 별도로 복사합니다. 일괄 복사 방법이 있습니까?

   // see http://stackoverflow.com/questions/7858817/unpacking-a-tuple-to-call-a-matching-function-pointer 
   template<int ...> 
   struct seq { }; 
   
   template<int N, int ...S> 
   struct gens : gens<N-1, N-1, S...> { }; 
   
   template<int ...S> 
   struct gens<0, S...> { 
       typedef seq<S...> type; 
   }; 
   
   class Task { 
   private: 
       bool visited; 
       bool reached; 
   protected: 
       std::vector<std::shared_ptr<Task>> dependsOn; 
       Task(); 
   public: 
       Task **d_dependsOn = NULL; 
       int d_dependsOnSize; 
       Task *d_self = NULL; 
   
       int streamId; 
       int id; 
       cudaStream_t stream; 
   
       virtual void copyToGPU() = 0; 
       virtual void start() = 0; 
       virtual void d_start() = 0; 
       virtual ~Task() {} 
       void init(); 
       void addDependency(std::shared_ptr<Task> t); 
       cudaStream_t dfs(); 
   }; 
   
   template<typename... Args> 
   class KernelTask : public Task { 
   private: 
       std::tuple<Args...> params; 
       dim3 threads; 
       dim3 blocks; 
       void (*kfp)(Args...); 
   
       template<int ...S> 
       void callFunc(seq<S...>) { 
        // inserting task into stream 
        this->kfp<<<this->blocks, this->threads, 0, this->stream>>>(std::get<S>(params) ...); 
        checkCudaErrors(cudaGetLastError()); 
   
        if (DEBUG) printf("Task %d: Inserting Task in Stream.\n", this->id); 
       } 
   
       template<int ...S> 
       __device__ void d_callFunc(seq<S...>) { 
        // inserting task into stream 
        this->kfp<<<this->blocks, this->threads, 0, this->stream>>>(std::get<S>(params) ...); 
   
        if (DEBUG) printf("Task %d: Inserting Task in Stream.\n", this->id); 
       } 
   
       KernelTask(int id, void (*kfp)(Args...), std::tuple<Args...> params, dim3 threads, dim3 blocks); 
   
   public: 
       ~KernelTask(); 
       void copyToGPU(); 
   
       void start() override { 
        callFunc(typename gens<sizeof...(Args)>::type()); 
       } 
   
       __device__ void d_start() override { 
        d_callFunc(typename gens<sizeof...(Args)>::type()); 
       } 
   
       static std::shared_ptr<KernelTask<Args...>> create(int id, void (*kfp)(Args...), std::tuple<Args...> params, dim3 threads, dim3 blocks); 
   }; 
   
   class Scheduler { 
   private: 
       std::vector<std::shared_ptr<Task>> tasks; 
   public: 
       Scheduler(std::vector<std::shared_ptr<Task>> &tasks) { 
        this->tasks = tasks; 
       } 
   
       void runCPUScheduler(); 
       void runGPUScheduler(); 
   }; 
   

   ---

   편집 :

   (1) CUDA에서 가상 함수 : 내가 얻을 다음 예제 scheduler에서 Warp Illegal Address 예외 :

   다음

코드입니다

struct Base { 
    __host__ __device__ virtual void start() = 0; 
    virtual ~Base() {} 
}; 

struct Derived : Base { 
    __host__ __device__ void start() override { 
     printf("In start\n"); 
    } 
}; 

__global__ void scheduler(Base *c) { 
    c->start(); 
} 

int main(int argc, char **argv) { 
    Base *c = new Derived(); 
    Base *d_c; 
    checkCudaErrors(cudaMalloc(&d_c, sizeof(Derived))); 
    checkCudaErrors(cudaMemcpy(d_c, c, sizeof(Derived), cudaMemcpyHostToDevice)); 

    c->start(); 
    scheduler<<<1,1>>>(d_c); 

    checkCudaErrors(cudaFree(d_c)); 

    return 0; 
} 

(2) thrust::tuple 잘 처리됩니다.

(3) 제안 사항이 있습니다.

(4) 커널 함수 포인터를 커널에 전달하려면 어떻게합니까? 나는 다음과 같은 예에서 Warp Misaligned Address 예외를 얻을 : ".? 가상 기능은 내가 아래로 캐스팅 않고 그들을 피할 수있는 방법 CUDA에서 허용되지 않습니다"

__global__ void baz(int a, int b) { 
    printf("%d + %d = %d\n", a, b, a+b); 
} 

void schedulerHost(void (*kfp)(int, int)) { 
    kfp<<<1,1>>>(1,2); 
} 

__global__ void schedulerDevice(void (*kfp)(int, int)) { 
    kfp<<<1,1>>>(1,2); 
} 

int main(int argc, char **argv) { 
    schedulerHost(&baz); 
    schedulerDevice<<<1,1>>>(&baz); 
    return 0; 
} 

Answer 1

을

당신은 모두 가상 __host__ 및 __device__ 기능을 가질 수 있습니다 그러나 http://docs.nvidia.com/cuda/cuda-c-programming-guide/#virtual-functions

: __global__ 함수에 인수로 가상 함수가있는 클래스의 개체를 전달할 수 없습니다

을 .

---

"표준 : 얻을 호스트 기능입니다. 표준을 구현하는 방법이 있나요 :: 자신은 GPU 얻을? "

I 대신 thrust::tuple을 사용하여 제안이라고하는 갖는다 모두 __host__ 및 __device__ 구현 : http://thrust.github.io/doc/group__tuple.html

---

대하여 함수 포인터하십시오 __global__ 기능

어드레스 호스트 코드에서 가져온 것이 기기 코드에 사용 된 일 수 없습니다 (예 : k ernel).

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#function-pointers