cuda编程笔记（26）-- 核函数使用任务队列

最新推荐文章于 2025-10-28 23:37:51 发布

原创

最新推荐文章于 2025-10-28 23:37:51 发布 · 219 阅读

0 ·

CC 4.0 BY-SA版权

文章标签：

#笔记 #cuda

任务队列是常见的解耦方式。在CPU的多核编程里，经常有这种生产者消费者的场景：生产者线程往任务队列丢任务，消费者线程在任务队列取任务去干活。

在cuda的核函数里，似乎也可以这么设计，每个线程都去取任务干活。但是为了利用GPU的多线程的特性，可以将取任务的单位归属到block，一个block认领一个任务后，由block内的线程共同处理。这样即做到了任务队列解耦，还利用了GPU的多线程特性。

至于生产者消费者的异步模型，则可以通过cuda里的流来进行实现。

先看看自产自销的模式，熟悉任务队列的写法：

#ifndef __CUDACC__
#define __CUDACC__
#endif
#include <cuda_runtime.h>
#include <cstdio>

struct Task {//任务的数据结构由自己定义
    int id;
    int depth;
};

// ------------------- 全局任务队列定义 -------------------
struct TaskQueue {
    Task* tasks;
    int* head;
    int* tail;
    int  capacity;//现在的设计是固定容量，最多处理capacity个任务，而不是随时变化
};

__device__ TaskQueue gQueue;

// ------------------- 初始化队列 -------------------
__global__ void initQueue(Task* buf, int* head, int* tail, int capacity, int initTasks) {
    if (threadIdx.x == 0 && blockIdx.x == 0) {
        gQueue.tasks = buf;
        gQueue.head = head;
        gQueue.tail = tail;
        gQueue.capacity = capacity;