CUDA Kernel Scheduler на GPU

Question

Я пишу планировщик ядра CUDA. Планировщик получает вектор указателей Task и приводит их к исполнению. Указатели указывают на KernelTask объекты различных параметров типа для поддержки ядер с произвольными параметрами.

Существует версия процессора Планировщика и версия графического процессора. Версия CPU отлично работает. Он вызывает виртуальную функцию Task::start для выполнения ядра. Версия GPU имеет три проблемы:

1. Виртуальные функции в CUDA запрещены. Как я могу избежать их без каста?
2. std :: get - это хост-функция. Есть ли способ реализовать 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(); 
}; 

---

EDIT:

(1) Виртуальные функции в CUDA: Я получаю Warp Illegal Address исключение в scheduler в следующем примере:

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? "

я предлагаю использовать 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