【Cuda 编程思想】理解CUDA中的线程协作

Mr.Lee jack

已于 2025-04-02 16:17:58 修改

阅读量247

点赞数 4

分类专栏： cuda 文章标签： c++ cuda

于 2025-04-02 14:52:49 首次发布

本文链接：https://blog.youkuaiyun.com/xzpdxz/article/details/146939396

版权

cuda 专栏收录该内容

4 篇文章

订阅专栏

理解CUDA中的线程协作

线程协作是CUDA编程中的一个核心概念，指的是线程之间相互配合完成任务的方式。在GPU上，成千上万的线程同时运行，如何让它们高效协同工作是提高性能的关键。

线程协作的基本形式

共享内存协作

共享内存是块内线程协作的主要媒介：

__global__ void sharedMemoryExample(float *input, float *output, int N) {
    // 声明共享内存
    __shared__ float sharedData[256];
    
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int tid = threadIdx.x;
    
    // 每个线程加载一个元素到共享内存
    if (idx < N) {
        sharedData[tid] = input[idx];
    }
    
    // 确保所有线程都完成了数据加载
    __syncthreads();
    
    // 现在所有线程都可以访问共享内存中的数据
    if (idx < N && tid > 0 && tid < blockDim.x - 1) {
        output[idx] = sharedData[tid-1] + sharedData[tid] + sharedData[tid+1];
    }
}

这个例子展示了线程如何通过共享内存交换数据，实现邻域计算。

同步点协作

__syncthreads()是块内线程协作的基本同步工具：

__global__ void syncExample(int *data) {
    __shared__ int result[256];
    int tid = threadIdx.x;
    
    // 第一阶段计算
    result[tid] = tid * 2;
    
    // 同步点：确保所有线程完成第一阶段
    __syncthreads();
    
    // 第二阶段：使用第一阶段的结果
    if (tid > 0) {
        data[blockIdx.x * blockDim.x + tid] = result[tid] + result[tid-1];
    }
}

原子操作协作

原子操作允许线程安全地修改共享数据：

__global__ void atomicExample(int *counter) {
    // 每个线程原子地增加计数器
    atomicAdd(counter, 1);
}

案例

构建并行数据结构（如队列、堆栈或树）时，需要原子操作来安全地更新指针和计数器

__global__ void parallelQueue(int *data, int *queue, int *queueSize, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < N && isValid(data[idx])) {
        // 原子地获取队列中的下一个位置,并将queueSize + 1
        int pos = atomicAdd(queueSize, 1);
        queue[pos] = data[idx];
    }
}

当需要从多个线程收集结果时，原子操作可以确保正确聚合

__global__ void findMinMax(float *data, float *minResult, float *maxResult, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    
    for (int i = idx; i < N; i += stride) {
        float val = data[i];
        // 原子地更新最小值
        atomicMin(minResult, val);
        // 原子地更新最大值
        atomicMax(maxResult, val);
    }
}

高级线程协作模式

动态工作分配

线程可以动态分配工作，确保负载均衡

__global__ void dynamicWorkDistribution(float *data, float *results, int N, int *workQueue, int *queueIndex) {
    int tid = threadIdx.x;
    int bid = blockIdx.x;
    int idx = bid * blockDim.x + tid;
    
    while (true) {
        // 原子地获取下一个工作项
        int workItem = atomicAdd(queueIndex, 1);
        
        if (workItem >= N) break; // 没有更多工作
        
        int dataIndex = workQueue[workItem];
        // 处理这个数据项
        results[workItem] = processData(data[dataIndex]);
    }
}

工作窃取模式

允许空闲线程"窃取"其他线程的工作

__global__ void workStealing(Task *tasks, int numTasks) {
    __shared__ int blockTaskIndex;
    int tid = threadIdx.x;
    
    // 初始化块的任务索引
    if (tid == 0) {
        blockTaskIndex = blockIdx.x;
    }
    __syncthreads();
    
    while (true) {
        // 获取任务
        int myTask;
        if (tid == 0) {
            myTask = atomicAdd(&blockTaskIndex, 1);
        }
        // 广播任务ID给块内所有线程
        __shared__ int sharedTask;
        if (tid == 0) {
            sharedTask = myTask;
        }
        __syncthreads();
        myTask = sharedTask;
        
        if (myTask >= numTasks) {
            // 尝试从其他块窃取工作
            for (int otherBlock = 0; otherBlock < gridDim.x; otherBlock++) {
                if (otherBlock == blockIdx.x) continue;
                
                // 尝试窃取任务
                if (tid == 0) {
                    int stolenTask = atomicSub(&blockTaskIndices[otherBlock], 1);
                    if (stolenTask >= 0) {
                        sharedTask = stolenTask;
                    } else {
                        sharedTask = -1;
                    }
                }
                __syncthreads();
                
                if (sharedTask >= 0) {
                    // 执行窃取的任务
                    executeTask(tasks[sharedTask]);
                    break;
                }
            }
            
            // 如果没有任务可窃取，退出
            break;
        }
        
        // 执行分配的任务
        executeTask(tasks[myTask]);
    }
}

前缀和（并行扫描）

一个经典的线程协作算法，用于计算累积和：

__global__ void prefixSum(int *input, int *output, int N) {
    __shared__ int temp[2048]; // 假设块大小最大为1024
    int tid = threadIdx.x;
    int offset = 1;
    
    // 加载数据到共享内存
    temp[tid] = (tid < N) ? input[tid] : 0;
    __syncthreads();
    
    // 上扫描阶段
    for (int d = blockDim.x>>1; d > 0; d >>= 1) {
        __syncthreads();
        if (tid < d) {
            int ai = offset*(2*tid+1)-1;
            int bi = offset*(2*tid+2)-1;
            temp[bi] += temp[ai];
        }
        offset *= 2;
    }
    
    // 清除最后一个元素
    if (tid == 0) temp[blockDim.x - 1] = 0;
    
    // 下扫描阶段
    for (int d = 1; d < blockDim.x; d *= 2) {
        offset >>= 1;
        __syncthreads();
        if (tid < d) {
            int ai = offset*(2*tid+1)-1;
            int bi = offset*(2*tid+2)-1;
            int t = temp[ai];
            temp[ai] = temp[bi];
            temp[bi] += t;
        }
    }
    __syncthreads();
    
    // 写回结果
    if (tid < N) {
        output[tid] = temp[tid];
    }
}

线程协作的实际应用

数据分类与重组

__global__ void dataClassification(float *input, float *output, int N) {
    __shared__ int positive_count;
    __shared__ int negative_count;
    __shared__ int positive_indices[256];
    __shared__ int negative_indices[256];
    
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int tid = threadIdx.x;
    
    // 初始化共享计数器
    if (tid == 0) {
        positive_count = 0;
        negative_count = 0;
    }
    __syncthreads();
    
    // 第一阶段：分类
    bool is_positive = false;
    if (idx < N) {
        is_positive = (input[idx] > 0);
        // 原子操作确保线程安全地更新计数器和索引数组
        if (is_positive) {
            int pos = atomicAdd(&positive_count, 1);
            positive_indices[pos] = tid;
        } else {
            int pos = atomicAdd(&negative_count, 1);
            negative_indices[pos] = tid;
        }
    }
    __syncthreads();
    
    // 第二阶段：处理正值（线程协作，避免分支分化）
    for (int i = 0; i < positive_count; i++) {
        if (tid == positive_indices[i]) {
            // 所有处理正值的线程一起执行相同的代码路径
            output[idx] = sqrt(input[idx]) * sin(input[idx]);
        }
    }
    
    // 第三阶段：处理负值
    for (int i = 0; i < negative_count; i++) {
        if (tid == negative_indices[i]) {
            // 所有处理负值的线程一起执行相同的代码路径
            output[idx] = log(fabs(input[idx]) + 1);
        }
    }
}

这种方法的关键点：

使用共享内存记录不同类型数据的索引
通过原子操作安全地更新共享计数器
分阶段处理不同类型的数据，避免线程束分化
每个阶段内，执行相同操作的线程一起工作

协作式归约（求和）

__global__ void cooperativeReduction(float *input, float *output, int N) {
    __shared__ float sdata[256];
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 加载数据到共享内存
    sdata[tid] = (idx < N) ? input[idx] : 0;
    __syncthreads();
    
    // 执行归约（求和）
    for (int s = blockDim.x/2; s > 0; s >>= 1) {
        if (tid < s) {
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();
    }
    
    // 写回结果
    if (tid == 0) {
        output[blockIdx.x] = sdata[0];
    }
}