GPU 上的 CUDA 内核调度程序

Question

我正在编写一个 CUDA 内核调度程序。调度程序获取Task 指针向量并将它们带入执行。指针指向不同类型参数的KernelTask 对象，以支持具有任意参数的内核。

调度程序有 CPU 版本和 GPU 版本。 CPU版本工作正常。它调用虚函数Task::start 来执行一个内核。 GPU版本存在三个问题：

1. CUDA 中不允许使用虚拟函数。我怎样才能在不向下投射的情况下避免它们？
2. std::get 是一个主机函数。有没有办法为 GPU 实现 std::get 自己？
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 异常：

__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 中不允许使用虚拟函数。如何在不向下转换的情况下避免它们？”

您可以同时拥有虚拟__host__ 和__device__ 函数：http://docs.nvidia.com/cuda/cuda-c-programming-guide/#virtual-functions

但是：

不允许作为参数传递给__global__ 函数 具有虚函数的类的对象。

---

"std::get 是一个宿主函数。有没有办法为 GPU 实现 std::get 自己？"

我建议使用 thrust::tuple 代替它，它同时具有 __host__ 和 __device__ 实现： http://thrust.github.io/doc/group__tuple.html

---

关于函数指针：

主机代码中的__global__函数的地址不能 用于设备代码（例如启动内核）。

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