GPUのCUDAカーネルスケジューラ

Question

私はCUDAカーネルスケジューラを書いています。スケジューラーはTaskポインターのベクトルを取得し、それらを実行に持ち込みます。ポインタは、任意のパラメータを持つカーネルをサポートするために、異なる型パラメータのKernelTaskオブジェクトを指しています。

スケジューラのCPU版とGPU版があります。 CPUのバージョンはうまく動作します。仮想関数Task::startを呼び出してカーネルを実行します。 GPUバージョンには3つの問題があります。

1. 仮想関数はCUDAでは使用できません。どのように私はダウンキャストなしでそれらを避けることができますか？
2. std :: getはホスト関数です。標準を実装する方法はありますか::自分自身でGPUを入手しますか？
3. （低優先度）KernelTaskのオブジェクトのサイズが異なるため、すべてで別々にコピーします。バッチコピーの方法はありますか？ここで

コードです：

// 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：CUDAで

（1）仮想関数：私は、次の例に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例外を取得：

__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

「仮想関数は、私がダウンして鋳造することなく、それらを避けることができますどのようにCUDAに許可されていません？」しかしhttp://docs.nvidia.com/cuda/cuda-c-programming-guide/#virtual-functions

：

あなたは両方の仮想__host__と__device__機能を有することができる

を仮想関数で__global__関数の引数としてクラスの オブジェクトを渡すことが許可されていません。

---

「のstd ::取得は、ホスト機能である。STDを実装する方法ですが:: GPUのために自分を取得しますか？関数ポインタについて http://thrust.github.io/doc/group__tuple.html

---

：

__global__関数のアドレスIは__host__と__device__実装の両方を持っている代わりにthrust::tupleを使用してお勧めしたい「

ホストコードで取られたデバイスコードで使用される （例えば、kエルネル）。

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