首页 > 其他分享 >cuda block之间的同步测试

cuda block之间的同步测试

时间:2024-09-09 21:53:17浏览次数:14  
标签:同步 addr int 0x0 RZ cuda block size

cuda block之间的同步测试

本文测试了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 */

标签:同步,addr,int,0x0,RZ,cuda,block,size
From: https://blog.csdn.net/m0_61864577/article/details/142055128

相关文章

  • cuda下载
    参考文章:https://blog.csdn.net/mbdong/article/details/121769951CUDAdownload:https://developer.nvidia.com/cuda-12-2-0-download-archiveCUDNNdownload:https://developer.nvidia.com/rdp/cudnn-archive下载的cuDNN是一个压缩文件,将它解压并把所有的目录复制到CUDA安装......
  • 数据同步方式何来“高级”与“低级”之说?场景匹配才是真理!
    导读:数据同步方式的重要性对于数据集成领域的兴从业者不言而喻,选择正确的数据同步方式能让数据同步工作的成果事半功倍。目市面上的数据同步工具很多,提供的数据同步方式也有多种,不同的数据同步方式有什么区别?如何选择适合自己业务需求的数据同步方式呢?本文将对此进行深入分析,并深......
  • 数据同步方式何来“高级”与“低级”之说?场景匹配才是真理!
    导读:数据同步方式的重要性对于数据集成领域的兴从业者不言而喻,选择正确的数据同步方式能让数据同步工作的成果事半功倍。目市面上的数据同步工具很多,提供的数据同步方式也有多种,不同的数据同步方式有什么区别?如何选择适合自己业务需求的数据同步方式呢?本文将对此进行深入分析,并......
  • 多线程篇(阻塞队列- PriorityBlockingQueue)(持续更新迭代)
    目录一、简介二、类图三、源码解析1.字段讲解2.构造方法3.入队方法put浮调整比较器方法的实现入队图解4.出队方法takedequeue下沉调整比较器方法的实现出队图解四、总结一、简介PriorityBlockingQueue队列是JDK1.5的时候出来的一个阻塞队列。但是该队......
  • 数据同步和数据备份
       日常使用的移动手机或者是电脑等其它电子产品都是每天在产生不同的数据。数据安全性的 保证需要有很多的计算机程序设计的运行程序进行有效保证。电子产品是硬件设备,硬件设 备就像机器机械一样是可以看得见摸得着的具体物件设施。物质和能量,源头物料可以通过 不同的设......
  • Java 同步锁性能分析
    同步锁在多线程编程中是保证线程安全的重要工具,其性能开销一直是不可忽视的存在。(一)性能验证说明为了直观说明我们可以直接先准备两个Java代码用例,我们通过高并发环境下的计数器递增操作来对比使用同步锁和不使用同步锁的性能差异。1.使用同步锁的代码示例使用ReentrantLock来保护......
  • Java的并发编程模型同步器
    在Java的并发编程中,同步器(Synchronizer)是一个非常重要的概念,它用于管理多个线程之间的协作,以确保线程间的正确交互和数据的一致性。Java并发包java.util.concurrent中提供了多种同步器,这些同步器主要用于实现锁(Locks)和其他并发原语(ConcurrencyPrimitives)。主要的同步器包括:......
  • [Java基础]IO的同步和阻塞
    同步与异步什么是同步与异步呢?百度百科是这样定义的:同步指两个或两个以上随时间变化的量在变化过程中保持一定的相对关系。异步与同步相对(这解释让我无言相对)所以,我们需要明确的是同步与异步针对的是两个或者两个以上的事物。对于同步而言,一个任务(调用者)的完成需要依赖另一个......
  • 多线程篇(阻塞队列- BlockingQueue)(持续更新迭代)
    目录一、了解什么是阻塞队列之前,需要先知道队列1.Queue(接口)二、阻塞队列1.前言2.什么是阻塞队列3.Java里面常见的阻塞队列三、BlockingQueue(接口)1.前言2.简介3.特性3.1.队列类型3.2.队列数据结构2.简介4.核心功能入队(放入数据)出队(取出数据)总结四......
  • 时间同步服务
    多主机协作工作时,各个主机的时间同步很重要,时间不一致会造成很多重要应用的故障,如:加密协议,日志,集群等。利用NTP(NetworkTimeProtocol)协议使网络中的各个计算机时间达到同步。目前NTP协议属于运维基础架构中必备的基本服务之一。主流的时间同步有两种实现方案ntpntp是......