本文测试了cuda block之间的同步行为

一.小结

• 1.cuda没有提供block之间的同步机制
• 2.本文通过一个计数器,实现同步(while判断全局变量计数是否等于总的线程数)
• 3.当GPU可以容纳所有的线程时
  • A.while循环中为空,kernel无法退出(从SASS指令看,没有退出条件,会死循环)
  • B.while循环中加入printf,kernel可以退出(从SASS指令看,有退出机制),结果符合预期
  • C.如何让编译器生成符合预期的指令呢?(TODO)
• 4.当GPU不能容纳所有的线程时,kernel阻塞
  • 因为,第一批ThreadBlock无法退出(释放资源),导致每二批TheadBlock无法调度
  • 因此,从设计上,cuda不能实现,同一个kernel的block之间同步.只能拆分成二个kernel
• 5.atom操作只是保证多线程并发对一个地址的操作是串行的(让每个线程的操作都能得到正确,顺序不一定)
  • atom单元是异步执行的,指令发射后,并不会等待执行完成,想要获得最新的值,要靠同步指令
• 6.__threadfence() 只是保证该函数调用点之前的所有全局内存和共享内存的写入
• 7.__threadfence_block() 只是确保在调用该函数之前线程块中的所有共享内存和全局内存的写入

二.复现过程

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

#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 int add_count = 0;
__device__ unsigned int reach_point = 0;

__global__ void kernel_atom_add(int *addr,unsigned int thread_size)
{
    unsigned int tid  = threadIdx.x + blockIdx.x * blockDim.x;
    atomicAdd_system(&addr[0],1); //每个thread加1
    atomicAdd_system(&reach_point,1); //每个thread加1
    __prof_trigger(0);
    while(thread_size!=reach_point)
    {
        printf("tid:%d\n",tid);
    }
    __prof_trigger(1);
    //才能让每个线程的结果都一样
    if(threadIdx.x==0) printf("blkid:%04d tid:%08d value:%08d\n",blockIdx.x,tid,addr[0]);
}

__global__ void kernel_atom_add_v1(int *addr,unsigned int thread_size)
{
    unsigned int tid  = threadIdx.x + blockIdx.x * blockDim.x;
    atomicAdd_system(&addr[0],1);
    //__syncthreads();//如果不加同步,下面每个线程不能保证一样,说明atom指令不能达到同步的效果,是异步执行的
    //__threadfence();//也能实现同步的效果
    if(threadIdx.x%32==0) printf("blkid:%04d tid:%08d value:%08d\n",blockIdx.x,tid,addr[0]);
}

int main(int argc,char *argv[])
{
    int deviceid=0;cudaSetDevice(deviceid);  
    //测试1 
    if(1)
    {
        int block_count=28;int block_size=32*4*8;
        int thread_size=block_count*block_size;
        {
            int *addr;CHECK_CUDA(cudaMalloc(&addr, thread_size*4));
            kernel_atom_add<<<block_count, block_size>>>(addr,thread_size);CHECK_CUDA(cudaDeviceSynchronize());
        }    
    }
    //测试2
    if(0)
    {
        int block_count=1;int block_size=32*4*8;
        int thread_size=block_count*block_size;
        {
            int *addr;CHECK_CUDA(cudaMalloc(&addr, thread_size*4));
            kernel_atom_add_v1<<<block_count, block_size>>>(addr,thread_size);CHECK_CUDA(cudaDeviceSynchronize());
        }
    }
    return 0;
}
EOF
/usr/local/cuda/bin/nvcc -std=c++17 -arch=sm_86 -lineinfo  -o atom_test atom_test.cu -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcuda
./atom_test
/usr/local/cuda/bin/cuobjdump --dump-sass ./atom_test

三.输出

• while中为空(死循环)

/*0100*/                   PMTRIG 0x1 ;                                  /* 0x0000000100007801 */
                                                                         /* 0x000fc60003800000 */
/*0110*/                   LDG.E R4, [R4.64] ;                           /* 0x0000000404047981 */
                                                                         /* 0x001ea4000c1e1900 */
/*0120*/                   ISETP.NE.AND P0, PT, R4, c[0x0][0x168], PT ;  /* 0x00005a0004007a0c */
                                                                         /* 0x004fda0003f05270 */
/*0130*/              @!P0 BRA 0x150 ;                                   /* 0x0000001000008947 */
                                                                         /* 0x000fea0003800000 */
/*0140*/                   BRA 0x140 ;                                   /* 0xfffffff000007947 */
                                                                         /* 0x000fea000383ffff */
/*0150*/                   S2R R0, SR_TID.X ;                            /* 0x0000000000007919 */
                                                                         /* 0x000e220000002100 */
/*0160*/                   PMTRIG 0x2 ;                                  /* 0x0000000200007801 */

• while中加了printf(有退出机制)

/*0130*/                   PMTRIG 0x1 ;                                  /* 0x0000000100007801 */
                                                                         /* 0x000fc60003800000 */
/*0140*/                   LDG.E R4, [R4.64] ;                           /* 0x0000002404047981 */
                                                                         /* 0x001ea2000c1e1900 */
/*0150*/                   IADD3 R18, P1, R1, c[0x0][0x20], RZ ;         /* 0x0000080001127a10 */
                                                                         /* 0x000fca0007f3e0ff */
/*0160*/                   IMAD.X R2, RZ, RZ, c[0x0][0x24], P1 ;         /* 0x00000900ff027624 */
                                                                         /* 0x000fe200008e06ff */
/*0170*/                   ISETP.NE.AND P0, PT, R4, c[0x0][0x168], PT ;  /* 0x00005a0004007a0c */
                                                                         /* 0x004fda0003f05270 */
/*0180*/              @!P0 BRA 0x2a0 ;                                   /* 0x0000011000008947 */
                                                                         /* 0x000fea0003800000 */
/*0190*/                   UMOV UR38, 0x0 ;                              /* 0x0000000000267882 */
                                                                         /* 0x000fe40000000000 */
/*01a0*/                   UMOV UR39, 0x0 ;                              /* 0x0000000000277882 */
                                                                         /* 0x000fc40000000000 */
/*01b0*/                   STL [R1], R17 ;                               /* 0x0000001101007387 */
                                                                         /* 0x0001e20000100800 */
/*01c0*/                   MOV R4, UR38 ;                                /* 0x0000002600047c02 */
                                                                         /* 0x000fe20008000f00 */
/*01d0*/                   IMAD.U32 R5, RZ, RZ, UR39 ;                   /* 0x00000027ff057e24 */
                                                                         /* 0x000fe2000f8e00ff */
/*01e0*/                   MOV R7, R2 ;                                  /* 0x0000000200077202 */
                                                                         /* 0x000fe20000000f00 */
/*01f0*/                   IMAD.MOV.U32 R6, RZ, RZ, R18 ;                /* 0x000000ffff067224 */
                                                                         /* 0x000fc400078e0012 */
/*0200*/                   MOV R20, 0x0 ;                                /* 0x0000000000147802 */
                                                                         /* 0x000fe40000000f00 */
/*0210*/                   MOV R21, 0x0 ;                                /* 0x0000000000157802 */
                                                                         /* 0x000fc80000000f00 */
/*0220*/                   CALL.ABS.NOINC 0x0 ;                          /* 0x0000000000007943 */
                                                                         /* 0x001fea0003c00000 */
/*0230*/                   UMOV UR4, 0x0 ;                               /* 0x0000000000047882 */
                                                                         /* 0x000fe40000000000 */
/*0240*/                   UMOV UR5, 0x0 ;                               /* 0x0000000000057882 */
                                                                         /* 0x000fe20000000000 */
/*0250*/                   IMAD.U32 R4, RZ, RZ, UR4 ;                    /* 0x00000004ff047e24 */
                                                                         /* 0x000fc4000f8e00ff */
/*0260*/                   IMAD.U32 R5, RZ, RZ, UR5 ;                    /* 0x00000005ff057e24 */
                                                                         /* 0x000fca000f8e00ff */
/*0270*/                   LDG.E R4, [R4.64] ;                           /* 0x0000002404047981 */
                                                                         /* 0x000ea4000c1e1900 */
/*0280*/                   ISETP.NE.AND P0, PT, R4, c[0x0][0x168], PT ;  /* 0x00005a0004007a0c */
                                                                         /* 0x004fda0003f05270 */
/*0290*/               @P0 BRA 0x1b0 ;                                   /* 0xffffff1000000947 */
                                                                         /* 0x000fea000383ffff */
/*02a0*/                   PMTRIG 0x2 ;                                  /* 0x0000000200007801 */