CUDA cooperative_groups grid

CUDA cooperative_groups grid_group测试

一.测试描述及小结
- 1.任务描述
- 2.输出
- 3.小结
二.复现步骤
三.grid_group.sync 代码对照
- 1.CUDA C
- 2.PTX
- 3.SASS

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 代码对照

CUDA_C_153">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 ERRBARIMAD.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