CUDA Kernel Scheduler auf GPU

Question

Ich schreibe einen CUDA Kernel Scheduler. Der Scheduler ruft einen Vektor von Task Zeigern ab und bringt sie zur Ausführung. Die Zeiger zeigen auf KernelTask Objekte mit unterschiedlichen Typparametern, um Kernel mit beliebigen Parametern zu unterstützen.

Es gibt eine CPU-Version des Schedulers und eine GPU-Version. Die CPU-Version funktioniert gut. Es ruft die virtuelle Funktion Task::start auf, um einen Kernel auszuführen. Die GPU-Version hat drei Probleme:

1. Virtuelle Funktionen sind in CUDA nicht zulässig. Wie kann ich sie ohne Casting vermeiden?
2. Std :: Get ist eine Host-Funktion. Gibt es eine Möglichkeit, std :: für die GPU zu implementieren?
3. (niedrige Priorität) Da KernelTask Objekte unterschiedlicher Größe sind, kopiere ich alle separat mit copyToGPU(). Gibt es eine Möglichkeit für das Kopieren im Stapelbetrieb? Hier

ist der Code:

// 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) Virtuelle Funktionen in CUDA: Ich erhalte eine Warp Illegal Address Ausnahme in scheduler im folgenden Beispiel:

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 funktioniert gut.

(3) Ich bin offen für Vorschläge.

(4) Wie übergebe ich einen Kernel-Funktionszeiger an einen Kernel? Ich erhalte eine Warp Misaligned Address Ausnahme in dem folgenden Beispiel: „? Virtuelle Funktionen in CUDA nicht erlaubt Wie kann ich sie vermeiden, ohne nach unten Gießen“

__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

Sie können sowohl virtuelle __host__ und __device__ Funktionen haben: http://docs.nvidia.com/cuda/cuda-c-programming-guide/#virtual-functions

jedoch:

Es ist nicht auf eine __global__ Funktion als Argument durchgelassen wird ein Objekt einer Klasse mit virtuellen Funktionen .

---

„std :: erhalten, ist eine Host-Funktion. Gibt es eine Möglichkeit std zu implementieren :: mich für die GPU bekommen? "

würde ich vorschlagen thrust::tuple stattdessen mit der sowohl eine __host__ und eine __device__ Umsetzung: http://thrust.github.io/doc/group__tuple.html

---

In Bezug auf Funktionszeiger:

Die Adresse einer __global__ Funktion Eingelesener Host-Code kann nicht im Gerätecode verwendet werden (z. B. um den k Ernel).

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