CUDA多线程

一、基础

1. ‌线程块与线程索引‌
   CUDA线程以‌线程块（Thread Block）‌为基本执行单元，每个线程块内包含多个线程，通过threadIdx、blockIdx等内置变量定位线程位置。线程块在GPU上并行执行，同一块内的线程可通过共享内存高效通信‌。

2. ‌块内同步：__syncthreads()‌
   用于同步线程块内所有线程的执行进度，确保共享内存的写入在所有线程可见后继续执行后续操作。例如，在累加共享内存时需两次同步：第一次清空内存，第二次汇总结果‌。

   __shared__ int sh_arr;
   if (threadIdx.x < 10) sh_arr[threadIdx.x] = 0;
   __syncthreads(); // 同步清空操作
   atomicAdd(&sh_arr[tid], 1);
   __syncthreads(); // 同步累加操作

3. ‌全局同步限制‌
   CUDA默认不支持跨线程块的全局同步，需通过‌原子操作‌或‌协作组（Cooperative Groups）‌实现更复杂的同步逻辑。协作组允许自定义线程组粒度（如warp、网格级），并通过sync()方法同步‌。

二、‌核心机制 

• ‌主上下文共享（运行时API）‌：
  • 同一进程的线程共享设备的主上下文（由cudaSetDevice隐式创建）。
  • 线程首次调用CUDA运行时API时自动附加到主上下文。

    #include <thread>
    #include <cuda_runtime.h>
    
    void thread_task(int device_id) {
        cudaSetDevice(device_id);  // 切换设备（主上下文自动附加）
        float *d_data;
        cudaMalloc(&d_data, 1024);  // 共享主上下文资源
        // ... 执行内核或内存操作
    }
    
    int main() {
        std::thread t1(thread_task, 0);  // 线程1使用设备0
        std::thread t2(thread_task, 1);  // 线程2使用设备1
        t1.join();
        t2.join();
        return 0;
    }

• ‌线程局部上下文（驱动API）‌：
  • 使用cuCtxCreate + cuCtxPushCurrent显式创建独立上下文。
  • 每个线程维护独立的上下文堆栈。
    ‌1）独立上下文模式‌
    每个线程创建独立上下文，避免资源竞争，但需注意设备资源限制（如显存、流处理器占用）。

    // 线程函数
    void thread_func() {
        CUcontext ctx;
        cuCtxCreate(&ctx, 0, device);  // 每个线程独立创建上下文
        cuCtxPushCurrent(ctx);          // 设为当前上下文
        
        // 执行CUDA操作（如内存分配、内核启动）
        
        cuCtxPopCurrent(&ctx);          // 弹出上下文
        cuCtxDestroy(ctx);              // 销毁上下文
    }

    2）共享上下文模式‌
    主线程创建上下文，其他线程通过同步机制安全共享。

    // 全局变量
    CUcontext g_ctx;
    std::mutex g_ctx_mutex;
    
    // 主线程初始化
    cuCtxCreate(&g_ctx, 0, device);
    
    // 子线程操作
    void thread_func() {
        std::lock_guard<std::mutex> lock(g_ctx_mutex);  // 互斥锁保护
        cuCtxPushCurrent(g_ctx);        // 压入共享上下文
        
        // 执行CUDA操作
        
        CUcontext popped;
        cuCtxPopCurrent(&popped);       // 弹出上下文
    }

三、同步策略

1. ‌主机端多线程同步‌

   • ‌锁保护非线程安全API‌：CUDA上下文操作（如cudaMemcpy、cudaMalloc）需通过互斥锁（如std::mutex）保护，避免多线程竞争导致未定义行为‌。
   • ‌独立设备上下文‌：多进程或多线程可各自绑定独立GPU设备（cudaSetDevice()），减少资源争用‌。

2. ‌设备端高效同步‌

   • ‌原子操作‌：通过atomicAdd、atomicExch等函数实现全局变量的线程安全更新，避免数据竞争‌。
   • ‌流与事件同步‌：
     • ‌流内顺序执行‌：同一CUDA流中的操作按提交顺序执行，通过cudaStreamSynchronize()等待流完成‌。
     • ‌事件跨流协调‌：使用cudaEventRecord()记录事件，通过cudaStreamWaitEvent()实现流间依赖‌。

四、性能优化

1. ‌最小化同步开销‌

   • 避免频繁调用__syncthreads()，仅在关键数据依赖处同步‌。
   • 使用异步操作（如cudaMemcpyAsync）重叠数据传输与计算，隐藏延迟‌。

2. ‌共享内存优化‌

   • 共享内存访问需对齐以减少bank冲突，提升访存效率‌。
   • 利用共享内存缓存重复访问的全局数据，降低延迟‌。

3. ‌协助组‌

   • 协作组提供更灵活的线程同步机制，支持动态线程组划分（如warp、块、网格级）：

     #include <cooperative_groups.h>
     using namespace cooperative_groups;
     
     grid_group grid = this_grid();
     grid.sync(); // 网格级同步

     此特性需GPU架构支持（如Compute Capability 6.0+），适用于大规模并行任务调度‌。