CUDA Cooperative Groups是CUDA编程模型中引入的一组高级特性，提供了更灵活的线程组织和同步机制
通过Cooperative Groups，开发者可以在不同层次上组织线程，并执行更高效的并行操作
grid_group.sync 可用于整个grid同步

一.测试描述及小结

1.任务描述

• 一个thread block只有2个线程,4个thread block
• 用cooperative_groups的grid_group做所有线程的同步
• 因为grid_group没有广播功能,于是采用tid=0 的sm时钟做全局时钟
• 在Kernel中记录当前当前线程对应的smid、全局时钟、当前时钟

2.输出

tid:00 smid:00 local_ts:477113991510614 global_ts:477113991321194
tid:01 smid:00 local_ts:477113991510614 global_ts:477113991321194
tid:06 smid:06 local_ts:477113991510702 global_ts:477113991321194
tid:07 smid:06 local_ts:477113991510702 global_ts:477113991321194
tid:02 smid:02 local_ts:477136243949393 global_ts:477113991321194
tid:03 smid:02 local_ts:477136243949393 global_ts:477113991321194
tid:04 smid:04 local_ts:477161370613356 global_ts:477113991321194
tid:05 smid:04 local_ts:477161370613356 global_ts:477113991321194

3.小结

• 通过cooperative_groups的grid_group可以做所有线程块的同步,而__syncthreads()只能实现线程块内同步
• clock64()读取的是每个SM上的时钟计数器,该计数器从设备启动时开始计数,但不同SM之间并不保证同步
• 使用cooperative_groups的grid_group进行全网格同步(grid.sync())可以确保所有线程在同步点之前的操作都已完成
  但无法保证同步点之后的指令在所有SM上同时开始执行。由于硬件调度和指令级并行的存在,不同SM上的线程在同步点之后可能仍会有微小的执行时间差异。
  即使线程在同步后执行完全相同的指令序列,GPU的指令调度器可能会因各种原因导致不同SM上的指令开始执行的时刻略有差异,如：
  • 指令缓存命中率：不同 SM 的指令缓存状态可能不同，导致指令取指时间不同。
  • 资源竞争：SM 上的共享资源（如内存带宽）可能受到其他线程块的影响。
  • 硬件层面的不可控因素：GPU 硬件内部的微架构特性可能引入额外的延迟。
• 查看PTX和SASS指令,该功能是通过循环读取dram中的变量并判断实现的
• 测试的架构每个GPC有二个SM,从调度的顺序可见(4个thread_block采用的smid分别是0 2 4 6).用到了4个GPC,每个GPC出一个SM,而不是2个GPC

二.复现步骤

tee cooperative_groups.cu<<-'EOF'
#include <iostream>
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <stdio.h>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <algorithm>

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

#define CHECK_CUDA(call)                      \
  do {                              \
    cudaError_t err = call;                  \
    if (err != cudaSuccess) {                 \
      std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__; \
      std::cerr << " code=" << err << " (" << cudaGetErrorString(err) << ")" << std::endl; \
      exit(EXIT_FAILURE);                  \
    }                             \
  } while (0)

__device__ unsigned long long global_clock = 0;

struct node_data
{
  unsigned long long local_ts;
  unsigned long long global_ts;
  unsigned int smid;
};

__global__ void kernel_grid_sync(node_data *pdata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  unsigned int smid;
  asm volatile("mov.u32 %0, %smid;" : "=r"(smid));  
  cg::grid_group grid = cg::this_grid();
  __prof_trigger(0);//仅用于标记代码
  grid.sync();
  __prof_trigger(1);
  pdata[tid].smid=smid;
}

__global__ void kernel(node_data *pdata)
{
  unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
  unsigned int smid;
  asm volatile("mov.u32 %0, %smid;" : "=r"(smid));
  
  cg::grid_group grid = cg::this_grid();
  cg::thread_block block = cg::this_thread_block();
  
  __nanosleep(blockIdx.x*1000000000);
  block.sync();
  
  unsigned long long local_ts = 0;
  asm volatile ("mov.u64 %0, %clock64;" : "=l"(local_ts) :: "memory");
  if(tid==0)
  {
    global_clock=local_ts; //生成全局时钟
  }
  grid.sync();//全网格同步
  asm volatile ("mov.u64 %0, %clock64;" : "=l"(local_ts) :: "memory");
  
  pdata[tid].local_ts=local_ts;
  pdata[tid].global_ts=global_clock;
  pdata[tid].smid=smid;
}

int main(int argc,char *argv[])
{
  int deviceid=0;cudaSetDevice(deviceid); 
  int block_count=4;int block_size=2;
  int thread_size=block_count*block_size;
  node_data *pdata;
  CHECK_CUDA(cudaHostAlloc(&pdata,thread_size*sizeof(node_data),cudaHostAllocDefault));
  void *kernelArgs[] = {&pdata};
  cudaLaunchCooperativeKernel((void*)kernel_grid_sync, block_count, block_size, kernelArgs);
  cudaLaunchCooperativeKernel((void*)kernel, block_count, block_size, kernelArgs);
  CHECK_CUDA(cudaDeviceSynchronize());

  std::vector<int> indices(thread_size);
  for (int i = 0; i < thread_size; ++i) {
    indices[i] = i;
  }
  //按本地时钟大小排序(其实没有意义,因为不同SM的时钟没有可比性)
  std::sort(indices.begin(), indices.end(), [&pdata](int a, int b) {
    return pdata[a].local_ts < pdata[b].local_ts;
  });  
  for(int i=0;i<thread_size;i++)
  {
    int idx=indices[i];
    printf("tid:%02d smid:%02d local_ts:%lld global_ts:%lld\n",
        idx,pdata[idx].smid,
        pdata[idx].local_ts,
        pdata[idx].global_ts);
  }
  CHECK_CUDA(cudaFreeHost(pdata));
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -lineinfo -o cooperative_groups cooperative_groups.cu \
 -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
./cooperative_groups

# 用NCU查看CUDA C/PTX/SASS的对应关系
/usr/local/NVIDIA-Nsight-Compute/ncu --set full --target-processes all \
        --export ncu_report_cooperative_groups -f ./cooperative_groups

三.grid_group.sync 代码对照

1.CUDA C

  __prof_trigger(0);
  grid.sync();
  __prof_trigger(1);

2.PTX

  mov.u32 %rd6, %envreg2;   # 特殊寄存器 %envreg<32> 是PTX的32个预定义的只读寄存器集合,在内核启动之前由驱动程序初始化。
  pmevent 0;
  setp.ne.s64 %p1, %rd1, 0; # 使用关系运算符比较两个数值，然后（可选地）通过应用布尔运算符将这个结果与谓词值结合起来。
  @%p1 bra $L__BB0_2;       # 在目标处继续执行。条件分支通过使用保护谓词来指定。分支目标必须是标签。
  trap;                     # 中止执行并生成一个中断到主机CPU。
$L__BB0_2:
  mov.u32 %r2, %ctaid.x;
  mov.u32 %r3, %tid.x;
  mov.u32 %r8, %tid.y;
  add.s32 %r9, %r3, %r8;
  mov.u32 %r10, %tid.z;
  neg.s32 %r11, %r10;
  setp.ne.s32 %p2, %r9, %r11;
  barrier.sync 0;           # 在CTA内同步,0指定一个逻辑屏障资源，该资源可以是立即常量或寄存器，其值为0到15。
  @%p2 bra $L__BB0_5;
  add.s64 %rd6, %rd1, 4;
  mov.u32 %r14, %ctaid.z;
  neg.s32 %r15, %r14;
  mov.u32 %r16, %ctaid.y;
  add.s32 %r17, %r2, %r16;
  setp.eq.s32 %p3, %r17, %r15;
  mov.u32 %r18, %nctaid.z;
  mov.u32 %r19, %nctaid.x;
  mov.u32 %r20, %nctaid.y;
  mul.lo.s32 %r21, %r19, %r20;
  mul.lo.s32 %r22, %r21, %r18;
  mov.u32 %r23, -2147483647;
  sub.s32 %r24, %r23, %r22;
  selp.b32 %r13, %r24, 1, %p3;
  atom.add.release.gpu.u32 %r12,[%rd6],%r13;
$L__BB0_4:
  ld.acquire.gpu.u32 %r25,[%rd6];
  xor.b32  %r26, %r25, %r12;
  setp.gt.s32 %p4, %r26, -1;
  @%p4 bra $L__BB0_4;
$L__BB0_5:
  barrier.sync 0;
  pmevent 1;

3.SASS

 PMTRIG 0x1 
 ISETP.NE.U32.AND P0, PT, RZ, c[0x0][0x90], PT 
 ISETP.NE.AND.EX P0, PT, RZ, c[0x0][0x8c], PT, P0 
@P0  BRA 0x7f13ef054d70 
 BPT.TRAP 0x1 
 S2R R2, SR_TID.Z 
 ULDC.64 UR6, c[0x0][0x118] 
 BSSY B0, 0x7f13ef055040 
 S2R R9, SR_TID.X 
 S2R R0, SR_TID.Y 
 S2R R6, SR_CTAID.X 
 BAR.SYNC 0x0 
 IMAD.MOV R3, RZ, RZ, -R2 
 IADD3 R0, R9, R0, RZ 
 ISETP.NE.AND P0, PT, R0, R3, PT 
@P0  BRA 0x7f13ef055030 
 S2UR UR4, SR_CTAID.Z 
 S2R R3, SR_LANEID 
 IMAD.MOV.U32 R0, RZ, RZ, c[0x0][0xc] 
 S2UR UR5, SR_CTAID.Y 
 UIADD3 UR4, -UR4, URZ, URZ 
 IADD3 R2, R6, UR5, RZ 
 ISETP.NE.AND P0, PT, R2, UR4, PT 
 MEMBAR.ALL.GPU 
 VOTEU.ANY UR4, UPT, PT 
 IMAD.MOV R0, RZ, RZ, -R0 
 FLO.U32 R4, UR4 
 MOV R5, c[0x0][0x14] 
 UPOPC UR5, UR4 
 IMAD R0, R0, c[0x0][0x10], RZ 
 ERRBAR
 IMAD.MOV.U32 R2, RZ, RZ, c[0x0][0x90] 
 IMAD R0, R0, R5, -0x7fffffff 
 SEL R0, R0, 0x1, !P0 
 ISETP.EQ.U32.AND P1, PT, R4, R3, PT 
 IMAD R5, R0, UR5, RZ 
 MOV R3, c[0x0][0x8c] 
@P1  ATOM.E.ADD.STRONG.GPU PT, R5, [R2.64+0x4], R5 
 S2R R8, SR_LTMASK 
 LOP3.LUT R8, R8, UR4, RZ, 0xc0, !PT 
 POPC R8, R8 
 SHFL.IDX PT, R11, R5, R4, 0x1f 
 IMAD R0, R0, R8, R11 
 LD.E.STRONG.GPU R5, [R2.64+0x4] 
 YIELD 
 LOP3.LUT R4, R5, R0, RZ, 0x3c, !PT 
 CCTL.IVALL 
 ISETP.GT.AND P0, PT, R4, -0x1, PT 
@P0  BRA 0x7f13ef054fd0 
 BSYNC B0 
 BRA.CONV ~URZ, 0x7f13ef055080 
 MOV R2, 0x370 
 CALL.REL.NOINC 0x7f13ef0550f0 
 BRA 0x7f13ef055090 
 BAR.SYNC 0x0 
 PMTRIG 0x2