CN113284038B - 用于执行计算的方法、计算设备、计算系统和存储介质 - Google Patents

用于执行计算的方法、计算设备、计算系统和存储介质 Download PDF

Info

Publication number
CN113284038B
CN113284038B CN202110244679.9A CN202110244679A CN113284038B CN 113284038 B CN113284038 B CN 113284038B CN 202110244679 A CN202110244679 A CN 202110244679A CN 113284038 B CN113284038 B CN 113284038B
Authority
CN
China
Prior art keywords
thread
corresponding program
program segment
compiled
thread group
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Active
Application number
CN202110244679.9A
Other languages
English (en)
Other versions
CN113284038A (zh
Inventor
不公告发明人
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Shanghai Bi Ren Technology Co ltd
Original Assignee
Shanghai Biren Intelligent Technology Co Ltd
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Shanghai Biren Intelligent Technology Co Ltd filed Critical Shanghai Biren Intelligent Technology Co Ltd
Priority to CN202110244679.9A priority Critical patent/CN113284038B/zh
Publication of CN113284038A publication Critical patent/CN113284038A/zh
Priority to US17/686,413 priority patent/US11886846B2/en
Application granted granted Critical
Publication of CN113284038B publication Critical patent/CN113284038B/zh
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06T—IMAGE DATA PROCESSING OR GENERATION, IN GENERAL
    • G06T1/00—General purpose image data processing
    • G06T1/20—Processor architectures; Processor configuration, e.g. pipelining
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G06F8/45—Exploiting coarse grain parallelism in compilation, i.e. parallelism between groups of instructions
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/30003—Arrangements for executing specific machine instructions
    • G06F9/3005—Arrangements for executing specific machine instructions to perform operations for flow control
    • G06F9/30054—Unconditional branch instructions
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/30003—Arrangements for executing specific machine instructions
    • G06F9/30076—Arrangements for executing specific machine instructions to perform miscellaneous control operations, e.g. NOP
    • G06F9/30087—Synchronisation or serialisation instructions
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/30003—Arrangements for executing specific machine instructions
    • G06F9/30076—Arrangements for executing specific machine instructions to perform miscellaneous control operations, e.g. NOP
    • G06F9/3009—Thread control instructions
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/30145—Instruction analysis, e.g. decoding, instruction word fields
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38—Concurrent instruction execution, e.g. pipeline or look ahead
    • G06F9/3836—Instruction issuing, e.g. dynamic instruction scheduling or out of order instruction execution
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38—Concurrent instruction execution, e.g. pipeline or look ahead
    • G06F9/3836—Instruction issuing, e.g. dynamic instruction scheduling or out of order instruction execution
    • G06F9/3851—Instruction issuing, e.g. dynamic instruction scheduling or out of order instruction execution from multiple instruction streams, e.g. multistreaming
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/46—Multiprogramming arrangements
    • G06F9/48—Program initiating; Program switching, e.g. by interrupt
    • G06F9/4806—Task transfer initiation or dispatching
    • G06F9/4843—Task transfer initiation or dispatching by program, e.g. task dispatcher, supervisor, operating system
    • G06F9/4881—Scheduling strategies for dispatcher, e.g. round robin, multi-level priority queues

Landscapes

  • Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • Software Systems (AREA)
  • General Engineering & Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • Multimedia (AREA)
  • Devices For Executing Special Programs (AREA)

Abstract

本公开涉及一种用于执行计算的方法、计算设备、计算系统和存储介质。该方法包括：经由编译器，确认待编译核函数中是否存在关于修改线程组请求的调用指令；响应于确认待编译核函数中存在关于修改线程组请求的调用指令，确定与调用指令相关联的对应程序段；为对应程序段配置所需的线程组和线程私有寄存器；以及在对应程序段插入控制指令，以用于使得为对应程序段所配置的线程组执行对应程序段的相关计算，并且未被配置的线程组不执行相关计算。本公开在针对以上需要改变线程组的拓扑结构的计算任务时能够提高计算设备整体性能和使得编写与维护容易，降低代码的错误率。

Description

用于执行计算的方法、计算设备、计算系统和存储介质
技术领域
本公开的实施例总体上涉及信息处理领域，更具体地涉及一种用于执行计算的方法、计算设备、计算系统和存储介质。
背景技术
在传统的计算设备中，例如图形处理器通用计算设备GPGPU(General-purposecomputing on graphics processing units)计算设备中，是让CPU端(Host端)启动GPU设备端(Device端)的代码进行运行。设备端代码通常也称为核函数(Kernel Function)。由于GPU执行大规模高并发的计算任务时会同时启动成千上万的线程(thread)来进行执行计算。为了有效编写与执行并发任务，启动核函数时，并发任务所对应的大规模线程组会被程序代码按照一定的拓扑结构进行组织。例如，现有的Nvidia CUDA使用网格(Grid)、线程组(Thread Block)来组织线程；OpenCL编程里面使用NDRange和WorkGroup进行组织。
在传统的用于执行计算的方案中，如果核函数被启动运行，其线程组的拓扑结构(例如，线程组大小或者三维尺寸)通常不被改变。在真实应用中，一个完整的GPU计算任务经常需要不同的线程组拓扑结构来进行计算。例如，针对一个AI计算任务，第一步中一个线程组(thread block)使用8x8x8做卷积计算，此时需要512个线程参与执行。第二步一个线程组(thread block)使用32x32做全连接的矩阵运算，此时需要 1024个线程参与执行。由于在传统的用于执行计算的方案中，启动核函数后通常不改变线程组的拓扑结构，因此，针对以上需要改变线程组的拓扑结构的计算任务时，通常采用以下两种方法。第一种方法例如包括：把该计算任务拆解成多个核函数。例如，第一核函数用于启动8x8x8的卷积计算，第二核函数用于启动32x32全连接的矩阵运算。两个核函数依次执行，使得512个线程和1024个线程分别在第一核函数和第二核函数被启动运行时参与执行。第二种方法例如包括：不使用GPGPU提供的线程组的拓扑结构，而是直接使用线程(thread)物理索引(index) 进行编程。而上述第一种方法容易造成多次启动核函数开销大，并且消耗内存存储中间结果，浪费内存带宽，增加数据延迟。上述第二种方法中，线程的编号不再具备任务分配拓扑的属性。用户需要自行计算线程的编号到任务拓扑的映射，因而增加了编程复杂性。同时，当一个线程有可能处理多个数据点时，任务分配复杂，且不一定均衡，有的线程可能处理比较多的数据点，有的可能处理比较少的数据点，进而造成代码编写困难、工作量大以及容易产生错误。
综上，传统的用于执行计算的方案，存在的不利之处在于，针对需要改变线程组的拓扑结构的计算任务时，容易因多次启动核函数开销大、浪费内存带宽、增加数据延迟等原因使得计算性能变差，或者代码编写困难且容易出错。
发明内容
本公开提供了一种用于执行计算的方法、计算设备、计算系统和计算机可读存储介质，能够针对以上需要改变线程组的拓扑结构的计算任务时能够提高计算设备整体性能，使得代码编写与维护容易，以及降低代码的错误率。
根据本公开的第一方面，提供了一种用于执行计算的方法。该方法包括：经由编译器，确认待编译核函数中是否存在关于修改线程组请求的调用指令；响应于确认待编译核函数中存在关于修改线程组请求的调用指令，确定与调用指令相关联的对应程序段；为对应程序段配置所需的线程资源；以及在对应程序段插入控制指令，以用于使得为对应程序段所配置的线程资源执行对应程序段的相关计算，并且未被配置的线程资源不执行相关计算。
根据本发明的第二方面，还提供了一种计算设备。该计算设备包括：至少一个处理器；以及与至少一个处理器通信连接的存储器；其中存储器存储有可被至少一个处理器执行的指令，指令被至少一个处理器执行，以使计算设备能够本公开的第一方面的方法。
根据本公开的第三方面，还提供了一种计算机可读存储介质。该计算机可读存储介质上存储有计算机程序，计算机程序被机器执行时执行本公开的第一方面的方法。
根据本公开的第四方面，还提供了一种计算系统。该计算系统包括：编译器模块，被配置为识别所接收的待编译核函数中包括的关于修改线程组请求的调用指令，以及将调用指令的待编译核函数进行转换，以便生成经编译核函数，经编译核函数所包括的对应程序段配置有控制指令；驱动程序模块，被配置为基于经编译核函数确定经编译核函数所需启动的最大线程数量；以及执行单元计算核模块，被配置为响应于经编译核函数中的对应程序段的控制指令被执行，修改线程组的函数调用。
在一些实施例中，确定与调用指令相关联的对应程序段包括：基于调用指令的数量，将待编译核函数分成多个对应程序段。
在一些实施例中，为对应程序段配置所需的线程资源包括：为多个对应程序段中的每一个对应程序段配置所需的线程组和线程私有寄存器。
在一些实施例中，用于执行计算的方法还包括：响应于确认经编译核函数被启动，驱动程序确认经编译核函数所需的最大线程数量，经编译核函数包括被分别插入控制指令的多个对应程序段；以及执行单元计算核模块启动最大线程数量。
在一些实施例中，用于执行计算的方法还包括：响应于对应程序段的控制指令被调用，执行以下各项：使得为对应程序段所配置的线程组的所有线程同步；为所配置的线程组的每一个线程分配线程私有寄存器；重新配置每一个线程的编号，以便将所重新配置的编号写入所分配的线程私有寄存器。
在一些实施例中，用于执行计算的方法还包括：经由驱动程序，启动经编译核函数；经由编译器或者驱动程序，确认经编译核函数所需的最大线程数量；以及经由驱动程序，向图形处理器请求最大线程数量。
在一些实施例中，使得为对应程序段所配置的线程资源执行对应程序段的相关计算包括：通过控制流代码控制执行流封装对应程序段，使得如果确定当前线程的线程标识属于为对应程序段所配置的线程组标识集合，以当前线程执行对应程序段的相关计算；以及如果确定当前线程的线程标识不属于为对应程序段所配置的线程组标识集合，使得当前线程被跳过。
在一些实施例中，确认待编译核函数中是否存在关于修改线程组请求的调用指令包括：经由编译器，扫描待编译核函数，以确认待编译核函数中是否包括关于修改线程组请求的调用指令，关于修改线程组请求的调用指令指示对应程序段所需的线程配置。
在一些实施例中，驱动程序模块被配置在中央处理器，执行单元计算核模块被配置在图形处理器，中央处理器和图形处理器由计算系统包括。
在一些实施例中，编译器模块还被配置为：如果确认待编译核函数中包括关于修改线程组请求的调用指令，为与调用指令相关联的对应程序段配置所需的线程资源；以及在对应程序段插入控制指令，以用于使得为对应程序段所配置的线程资源执行对应程序段的相关计算，并且未被配置的线程资源不执行相关计算。
在一些实施例中，修改线程组的函数调用包括：重新恢复或者暂停线程、为每一个被配置的线程分配线程私有寄存器。
应当理解，本部分所描述的内容并非旨在标识本公开的实施例的关键或重要特征，也不用于限制本公开的范围。本公开的其它特征将通过以下的说明书而变得容易理解。
附图说明
图用于更好地理解本方案，不构成对本申请的限定。
图1示意性示出了传统的计算设备的示意图。
图2示出了用于实施根据本公开的实施例的用于执行计算的方法的计算设备的示意图。
图3示出了根据本公开的一些实施例的用于执行计算的方法的流程图。
图4示出了根据本公开的一些实施例的核函数的示意图。
图5示出了根据本公开的一些实施例的用于执行控制指令的方法的流程图。
图6示出了用于实施根据本公开的实施例的用于执行计算的方法的计算系统的示意图。
图7示出了用于实施根据本公开的实施例的图形处理器的示意图。
在各个附图中，相同或对应的标号表示相同或对应的部分。
具体实施方式
下面将参照附图更详细地描述本公开的优选实施例。虽然附图中显示了本公开的优选实施例，然而应该理解，可以以各种形式实现本公开而不应被这里阐述的实施例所限制。相反，提供这些实施例是为了使本公开更加透彻和完整，并且能够将本公开的范围完整地传达给本领域的技术人员。
在本文中使用的术语“包括”及其变形表示开放性包括，即“包括但不限于”。除非特别申明，术语“或”表示“和/或”。术语“基于”表示“至少部分地基于”。术语“一个示例实施例”和“一个实施例”表示“至少一个示例实施例”。术语“另一实施例”表示“至少一个另外的实施例”。术语“第一”、“第二”等等可以指代不同的或相同的对象。
如前文所描述，在传统的用于执行计算的方案中，针对需要改变线程组的拓扑结构的计算任务时，通常通过把该计算任务所对应的核函数拆解成很多个核函数。图1示意性示出了传统的计算设备的示意图。如图1所示，计算设备100至少包括CPU端110(或称Host端)和GPU 端120(或称Device端)。假设一个CPU端110执行的原始核函数112 的名称例如为processImage。原始核函数112的第一步例如需要使用具有第一拓扑结构的网格122(Grid)所示的线程组(thread block)做卷积计算。原始核函数112的第二步例如需要使用具有第二拓扑结构的网格 124(Grid)所示的线程组(thread block)做全连接的矩阵运算。为了解决一个计算任务需要不同的线程组拓扑结构来进行计算的问题。传统的用于执行计算的方案例如将原始核函数112拆解为第一核函数112和第二核函数114。例如，通过代码“procesImage<<<dim3(2,2,1), dim3(10,10,3)>>>(…)”启动第一核函数112时，第一计算任务(例如前述卷积计算任务)在网格122(Grid)层次有2x2x1个线程组(Thread Block)，而每个线程组(Thread Block)例如有10x10x3个线程参与执行。然后，例如通过代码“procesImage<<<dim3(3,2,1), dim3(10,10,3)>>>(…)”启动第二核函数114时，第二计算任务(例如前述全连接的矩阵运算任务)在网格124(Grid)层次有3x2x1个线程组 (ThreadBlock)，而每个线程组(Thread Block)例如有10x10x3个线程参与执行。由此实现通过启动两个核函数分别满足不同网格大小的计算。然而，上述方法需要将一个核函数拆解成很多个核函数，进而造成多次启动核函数开销大；同时使得核函数启动有延迟，占用GPU控制单元(CP)的资源。并且，每个核函数的输出数据需要从GPU计算核 (Computation Core)写出到内存，下一个核函数启动后再读回到GPU 计算核心，消耗内存存储中间结果，浪费内存带宽，增加数据延迟，不利于计算设备整体性能的提升。
为了至少部分地解决上述问题以及其他潜在问题中的一个或者多个，本公开的示例实施例提出了一种用于执行计算的方法、计算设备、计算系统和计算机可读存储介质。在本公开方法中：通过编译器在确认待编译核函数中存在关于修改线程组请求的调用指令时，确定与调用指令相关联的对应程序段；并为对应程序段配置所需的线程资源；以及在对应程序段插入控制指令，以用于使得为对应程序段所配置的线程资源执行对应程序段的相关计算，并且未被配置的线程资源不执行相关计算。本公开能够在一个核函数内部动态修改线程组的大小(或称“线程组的数量”)以及拓扑结构，可以依然使用线程组内置的坐标信息进行数据访问，并且无需通过多次启动核函数来满足需要改变线程组的拓扑结构的计算任务的需要。因此，本公开减少了核函数启动的次数，代码更集中，性能更好，而且方便程序员或者编译器进行边界检查，降低代码的错误率。因而，本公开能够针对以上需要改变线程组的拓扑结构的计算任务时能够提高计算设备整体性能，使得代码编写与维护容易，以及降低代码的错误率。
图2示出了用于实施根据本公开的实施例的用于执行计算的方法的计算设备200的示意图。
如图2所示，计算设备200用于执行需改变线程组的拓扑结构的计算任务。计算设备200包括：调用指令确认单元202、对应程序段确认单元204、对应程序段线程资源配置206、控制指令插入单元208。计算设备200例如还包括：至少一个处理器(至少一个图形处理器和至少一个中央处理器，未示出)；以及与至少一个处理器通信连接的存储器(未示出)；其中存储器存储有可被至少一个处理器执行的指令，指令被至少一个处理器执行。计算设备200例如还包括编译器。
关于调用指令确认单元202，其用于经由编译器，确认待编译核函数中是否存在关于修改线程组请求的调用指令。
关于对应程序段确认单元204，其用于响应于确认待编译核函数中存在关于修改线程组请求的调用指令，确定与调用指令相关联的对应程序段。
关于对应程序段线程资源配置206，其用于为对应程序段配置所需的线程资源。所配置的线程资源例如：线程私有寄存器(Thread Local Register，简称为TLR)、执行组标量寄存器(Warp Scalar Register，简称为WSR)、常量寄存器(Constant Register，简称为CSR)。
关于控制指令插入单元208，其用于在对应程序段插入控制指令，以用于使得为对应程序段所配置的线程资源执行对应程序段的相关计算，并且未被配置的线程资源不执行相关计算。
以下结合图3和图4说明用于执行计算的方法300。图3示出了根据本公开的一些实施例的用于执行计算的方法300的流程图。图4示出了根据本公开的一些实施例的核函数的示意图。应当理解，方法300例如可以在图2所描述的计算设备200处执行。应当理解，方法300还可以包括未示出的附加组成部分、动作和/或可以省略所示出的组成部分、动作，本公开的范围在此方面不受限制。
在步骤302处，经由编译器，确认待编译核函数中是否存在关于修改线程组请求的调用指令。例如，经由编译器，扫描待编译核函数，以确认待编译核函数中是否包括关于修改线程组请求的调用指令，关于修改线程组请求的调用指令指示对应程序段所需的线程配置。如果确认不存在关于修改线程组请求的调用指令，跳转至步骤310处，针对核函数的代码进行编译。
例如，图4所示的待编译核函数402，即testCode()，其包括两部分代码，第一部分代码404需要8x8x8，即512个线程参与执行，第二部分代码406需要32x32，即1024个线程参与执行。编译器扫描整个待编译核函数402，确定是否存在关于修改线程组请求的调用指令408，调用指令408例如而不限于是change_block_shape()。调用指令408例如指示调整后的线程资源配置。
修改线程组例如包含但不限于修改线程组里面的线程数量和/或修改多个线程的组织拓扑结构。在步骤304处，响应于确认待编译核函数中存在关于修改线程组请求的调用指令，确定与调用指令相关联的对应程序段。
关于确定与调用指令相关联的对应程序段的方法，其例如包括：如果编译器确认待编译核函数中存在关于修改线程组请求的调用指令，基于调用指令的数量，将待编译核函数分成多个对应程序段。例如，如果编译器确认存在N次关于修改线程组请求的调用指令，则将待编译核函数代码分成N+1个对应的代码段。以及编译器确定每一个对应程序段所分配的线程组和线程私有寄存器。例如，如果编译器确认待编译核函数 402中包括关于修改线程组请求的调用指令408，则例如基于调用指令 408的数量，例如为1个，将核函数分成2个对应程序段，即第一对应程序段412和第二对应程序段414。
在步骤306处，为对应程序段配置所需的线程资源。例如，编译器为多个对应程序段中的每一个对应程序段配置所需的线程组和线程私有寄存器。例如，编译器确定第一对应程序段412需要8x8x8，即512个线程参与执行，第二对应程序段414需要32x32，即1024个线程参与执行。
在步骤308处，在对应程序段插入控制指令，以用于使得为对应程序段所配置的线程资源执行对应程序段的相关计算，并且未被配置的线程资源不执行相关计算。
关于在对应程序段插入控制指令的方式，其例如包括：通过控制流代码控制执行流封装对应程序段，使得如果确定当前线程的线程标识属于为对应程序段所配置的线程组标识集合，以当前线程执行对应程序段的相关计算；以及如果确定当前线程的线程标识不属于为对应程序段所配置的线程组标识集合，使得当前线程被跳过。关于控制流代码，其例如而不限于包括：“if…else”代码,“Switch”代码、go to代码等等。关于控制指令，其例如如图4中的第二对应程序段414所示，其包括：使得所分配的线程组的所有线程同步(对应的代码例如为“_sync_block__ threads()”)；如果确定当前线程标识属于(例如小于)为对应程序段(例如第二对应程序段414)所分配的线程标识集合(对应的代码例如为“if(thread_idx<Required_Thread_in_Section2)”)；为所分配的线程组的每一个线程分配线程私有寄存器(对应的代码例如为“set TLR for index”)；重新配置每一个线程(例如32x32线程)的编号，以便将所重新配置的编号存储在线程私有寄存器中(对应的代码例如为“Code in 32x32 thread block shape”)。经插入控制指令的对应程序段的代码例如如下，其中符号“…”示意其他未示出的代码。
“_sync_block__threads()
if(thread_idx<Required_Thread_in_Section2){
_change_block_shape(32,32,1)；//set Thread local register for index
…//Code in 32x32 thread block shape
}”。
在上述方案中，通过编译器在确认待编译核函数中存在关于修改线程组请求的调用指令时，确定与调用指令相关联的对应程序段；并为对应程序段配置所需的线程资源；以及在对应程序段插入控制指令，以用于使得为对应程序段所配置的线程资源执行对应程序段的相关计算，并且未被配置的线程资源不执行相关计算。本公开能够在一个核函数内部动态修改线程组的大小以及拓扑结构，可以依然使用线程组内置的坐标信息进行数据访问，并且无需通过多次启动核函数来满足需要改变线程组的拓扑结构的计算任务的需要。因此，本公开减少了核函数启动的次数，代码更集中，性能更好，而且方便程序员或者编译器进行边界检查，降低代码的错误率。因而，本公开能够针对以上需要改变线程组的拓扑结构的计算任务时能够提高计算设备整体性能，使得代码编写与维护容易，以及降低代码的错误率。
在一些实施例中，方法300还包括：经由驱动程序，启动经编译核函数；经由编译器或者驱动程序，确认经编译核函数所需的最大线程数量；以及经由驱动程序，向图形处理器请求最大线程数量。
在一些实施例中，方法300还包括：响应于确认经编译核函数被启动，驱动程序确认经编译核函数所需的最大线程数量，经编译核函数包括被分别插入控制指令的多个对应程序段；以及执行单元计算核模块启动最大线程数量。例如，当经编译核函数启动后，执行单元计算核模块 (硬件)会启动该经编译核函数所需的最大规模的线程数量(该最大规模的线程数量例如由编译器或者Driver确定)，最大规模的线程数量例如为1024个线程。在第一对应程序段被执行时，例如，只有512个线程会被使用。需要用到的512线程按照8x8x8进行编号组织，同时，执行单元计算核模块的其他线程资源，比如线程私有寄存器(ThreadLocal Register，TLR)都分配给这512个线程。剩余的512个线程直接挂起，以等待第二对应程序段被执行。当第一对应程序段被执行完，进行 _change_block_shape(32,32,1)时，所有线程都停下，然后，所有线程都被激活，然后按照新的拓扑编号，例如32x32，执行单元计算核模块将其他线程资源(如TLR)也对应分配给所有1024个线程。这样执行第二对应程序段时，线程组的线程拓扑就被动态修改了。修改线程组的线程拓扑例如是通过屏障寄存器，根据线程组的修改后的线程拓扑分配线程私有寄存器(Thread Local Register，TLR)。以及为线程编号和线程组尺寸配置TLR、WSR、CSR。关于屏障寄存器，下文将结合图7加以说明，在此，不再赘述。
在一些实施例中，方法300还包括对应程序段的控制指令被执行的方法500。图5示出了根据本公开的一些实施例的用于执行控制指令的方法500的流程图。应当理解，方法500例如可以在图2所描述的计算设备200处执行。应当理解，方法500还可以包括未示出的附加组成部分、动作和/或可以省略所示出的组成部分、动作，本公开的范围在此方面不受限制。
在步骤502，响应于对应程序段的控制指令被调用，使得为对应程序段所配置的线程组的所有线程同步。
由于对应程序段所配置的每一个线程调整线程编号，与分配线程私有寄存器资源等需要同时进行。因此，需要使得所有线程等待到该同步点后进行同步。
在步骤504，为所配置的线程组的每一个线程分配线程私有寄存器。针对对应程序段所配置的每一个线程进行同步之后，调整每一个线程的线程编号，以及为每一个线程分配线程私有寄存器。
关于为每一个线程分配线程私有寄存器的方法，其例如包括：按照现有GPU通常设计，一个线程组(Thread Block)总共的线程私有寄存器(TLR)的数量是固定的。启动的线程越多，每个线程所分配的线程私有寄存器的数量越少。例如，线程组一共有K个TLR，如果，启动N 个线程，则每个线程有K/N个TLR。如果启动2N个线程，则每个线程就只有K/2N个TLR。
关于线程私有寄存器的分割方法，其例如包括：使得每个线程都可以看到一段完整的线程私有寄存器，并且每个线程所能看到的完整的线程私有寄存器都从寄存器地址0开始编号的。
在步骤506，重新配置每一个线程的编号，以便将所重新配置的编号写入所分配的线程私有寄存器。关于重新配置每一个线程的编号的方法，其例如包括：在调整前，线程编号例如是是从(0,0,0),(0,0,1)到(7,7,7)，这些是记录在线程私有寄存器中的。修改线程私有寄存器和线程数量后，需要重新设置每个线程编号，比如从(0,0),(0,1)到(31,32)。该步骤根据线程的绝对物理地址把每个线程的新坐标重新计算后，写入线程私有寄存器即可。
通过采用上述技术手段，本公开可以通过一次核函数启动即可实现线程组形状的动态调整，进而能够动态支持需要多种形状线程组的计算任务的执行。
图6示出了用于实施根据本公开的实施例的用于执行计算的方法的计算系统600的示意图。应当理解，计算系统600还可以包括未示出的附加组成部分，本公开的范围在此方面不受限制。
如图6所示，计算系统600至少包括：编译器模块610、驱动程序模块620、执行单元计算核模块630。计算系统600还包括：编译器612、中央处理器622和图形处理器632。例如，编译器模块610例如配置在编译器612处，驱动程序模块620例如被配置在中央处理器622处，执行单元计算核模块630例如被配置在图形处理器632处。
关于编译器模块610，其被配置为识别所接收的待编译核函数614 中包括的关于修改线程组请求的调用指令，以及将调用指令的待编译核函数进行转换，以便生成经编译核函数616，经编译核函数所包括的对应程序段配置有控制指令。
关于将包括关于修改线程组请求的调用指令的核函数进行转换，编译器模块610被配置为：如果确认待编译核函数中包括关于修改线程组请求的调用指令，为与调用指令相关联的对应程序段配置所需的线程资源；以及在对应程序段插入控制指令，以用于使得为对应程序段所配置的线程资源执行对应程序段的相关计算，并且未被配置的线程资源不执行相关计算。
关于驱动程序模块620，其用于基于经编译核函数616确定经编译核函数所需启动的最大线程数量。例如，驱动程序模块620需要检查整个经编译核函数616中最大需要的线程数量。例如，整个经编译核函数 616包括两段，第一对应程序段需要512个线程(8x8x8)，第二对应程序段需要1024和线程(32x32)，那么，在启动经编译核函数616时，驱动程序模块620需要按照1024个线程数量请求GPU线程数量。需要说明的是，在经编译核函数616的第一对应程序段，驱动程序模块620 还是只配置第一对应程序段所需要的512个线程的坐标和线程私有寄存器以参与相关计算的执行。第二对应程序段之后的配置在通过控制指令被启动时在执行单元计算核模块动态修改线程组的函数调用。关于修改线程组，其例如包含但不限于修改线程组里面的线程数量和/或修改多个线程的组织拓扑结构。
执行单元计算核模块630，被配置响应于经编译核函数616中的对应程序段的控制指令被执行，修改线程组的函数调用。修改线程组的函数调用包括：重新恢复或者暂停线程、为每一个线程分配线程私有寄存器。
通过采用上述手段，本公开能够在一个核函数内部可以动态修改线程组的大小以及拓扑结构，从而减少了核函数启动的次数，代码更集中，性能更好，而且方便程序员或者编译器进行边界检查，降低代码的错误率。因而，本公开能够针对以上需要改变线程组的拓扑结构的计算任务时能够提高计算设备整体性能，使得代码编写与维护容易，以及降低代码的错误率。
以下结合图7说明图形处理器700。图7示出了根据本公开的实施例的图形处理器700的示意图。图形处理器700包括但不限于：线程组指令调度器702(Warp InstructionScheduler)、指令缓存704(Instruction Cache)、屏障寄存器706(Barriers Register)、各线程组资源寄存器708 Resource-per-warp Register)、算术逻辑单元710(ArithmeticLogical Unit， ALU)和通用寄存器712(General-Purpose Registers，GPRs)，该通用寄存器712可以为前文提及的TLR。
指令缓存704用于存储核函数的多条指令。
线程组指令调度器702用于为每个线程组提取一系列指令并存储至指令缓存704，以及依据程序计数器为每个线程组从指令缓存获取待执行的指令。每个线程组例如拥有独立的程序计数器(Program Counter， PC)寄存器，用于记录现在正执行的指令位置(即，指令地址)。每当为一个线程组从指令缓存提取一条指令后，相应的程序计数器寄存器累加一。线程组指令调度器702可以将指令送到算术逻辑单元710来执行，这些指令在指令集架构(Instruction Set Architecture，ISA)中定义。线程组指令调度器702可依据核函数中的相关指令让每个线程组负责处理同一个核函数中指定段的多个指令。
算术逻辑单元710可执行各种操作，算术逻辑单元710在执行操作的过程中，可从通用寄存器712的指定位置(又可称为来源地址)读取数据，和回写执行结果到通用寄存器的指定位置(又可称为目的地地址)。
屏障寄存器706例如可用于让软件同步不同线程组之间的执行。
各线程组资源寄存器708可用于让软件在执行时动态配置每个线程资源时能够使用的通用寄存器的空间范围。各线程组资源寄存器708例如存储每个线程组的基础位置的信息，每个基础位置指向通用寄存器712的特定地址。为了让不同线程组可以存取通用寄存器712中没有重叠的存储空间，软件可动态改变各线程组资源寄存器708的内容，用于设定每个线程组的基础位置。例如，软件可为多个线程组将通用寄存器 712分为对应的多块，以分别对应多个线程组。软件可在核函数的执行之前或执行的一开始设定各线程组资源寄存器708的内容，用于指向通用寄存器712中关联于每个线程组的基础地址。线程组指令调度器702 为第i个线程组从指令缓存702提取指令后，可依据第i个线程组资源寄存器的内容调整指令的来源地址和目的地地址，用于映射到通用寄存器712中动态配置给第i个线程组的存储空间。
上文所描述的各个过程和处理，例如方法300、500，可由计算设备处执行。该计算设备例如包括：至少一个处理器(至少一个图形处理器和至少一个中央处理器)；以及与至少一个处理器通信连接的存储器；其中存储器存储有可被至少一个处理器执行的指令，指令被至少一个处理器执行。在一些实施例中，300、500，可被实现为计算机软件程序，其被有形地包含于机器可读介质。在一些实施例中，计算机程序的部分或者全部可以经由ROM和/或通信单元而被载入和/或安装到计算设备上。当计算机程序被加载到RAM并由GPU和CPU执行时，可以执行上文描述的方法300、500，的一个或多个动作。
本公开可以是方法、装置、系统和/或计算机程序产品。计算机程序产品可以包括计算机可读存储介质，其上载有用于执行本公开的各个方面的计算机可读程序指令。计算机可读存储介质可以是可以保持和存储由指令执行设备使用的指令的有形设备。计算机可读存储介质例如可以是但不限于电存储设备、磁存储设备、光存储设备、电磁存储设备、半导体存储设备或者上述的任意合适的组合。
这里所描述的计算机可读程序指令可以从计算机可读存储介质下载到各个计算/处理设备，或者通过网络、例如因特网、局域网、广域网和 /或无线网下载到外部计算机或外部存储设备。这里参照根据本公开实施例的方法、装置(系统)和计算机程序产品的流程图和/或框图描述了本公开的各个方面。应当理解，流程图和/或框图的每个方框以及流程图和 /或框图中各方框的组合，都可以由计算机可读程序指令实现。
这些计算机可读程序指令可以提供给通用计算机、专用计算机或其它可编程数据处理装置的中央处理单元，从而生产出一种机器，使得这些指令在通过计算机或其它可编程数据处理装置的中央处理单元执行时，产生了实现流程图和/或框图中的一个或多个方框中规定的功能/动作的装置。也可以把这些计算机可读程序指令存储在计算机可读存储介质中，这些指令使得计算机、可编程数据处理装置和/或其他设备以特定方式工作，从而，存储有指令的计算机可读介质则包括一个制造品，其包括实现流程图和/或框图中的一个或多个方框中规定的功能/动作的各个方面的指令。
附图中的流程图和框图显示了根据本公开的多个实施例的系统、方法和计算机程序产品的可能实现的体系架构、功能和操作。在这点上，流程图或框图中的每个方框可以代表一个模块、程序段或指令的一部分，模块、程序段或指令的一部分包含一个或多个用于实现规定的逻辑功能的可执行指令。在有些作为替换的实现中，方框中所标注的功能也可以以不同于附图中所标注的顺序发生。例如，两个连续的方框实际上可以基本并行地执行，它们有时也可以按相反的顺序执行，这依所涉及的功能而定。也要注意的是，框图和/或流程图中的每个方框、以及框图和/ 或流程图中的方框的组合，可以用执行规定的功能或动作的专用的基于硬件的系统来实现，或者可以用专用硬件与计算机指令的组合来实现。
应该理解，可以使用上面所示的各种形式的流程，重新排序、增加或删除步骤。例如，本申请中记载的各步骤可以并行地执行也可以顺序地执行也可以不同的次序执行，只要能够实现本申请公开的技术方案所期望的结果，本文在此不进行限制。
上述具体实施方式，并不构成对本申请保护范围的限制。本领域技术人员应该明白的是，根据设计要求和其他因素，可以进行各种修改、组合、子组合和替代。

Claims (11)

1.一种用于执行计算的方法，包括：
经由编译器，确认待编译核函数中是否存在关于修改线程组请求的调用指令，关于修改线程组请求包括：关于修改线程组里面的线程数量和/或修改多个线程的组织拓扑结构的请求；
响应于确认待编译核函数中存在关于修改线程组请求的调用指令，确定与所述调用指令相关联的对应程序段；
为所述对应程序段配置所需的线程资源；以及
在所述对应程序段插入控制指令，以用于使得为所述对应程序段所配置的线程资源执行所述对应程序段的相关计算，并且未被配置的线程资源不执行相关计算；
其中，确定与所述调用指令相关联的对应程序段包括：
基于调用指令的数量，将待编译核函数分成多个对应程序段，所述对应程序段的数量比调用指令的数量大1；
并且，为所述对应程序段配置所需的线程资源包括：
为所述多个对应程序段中的每一个对应程序段配置所需的线程组和线程私有寄存器。
2.根据权利要求1所述的方法，还包括：
响应于确认经编译核函数被启动，基于经编译所述核函数所需的最大线程数量配置线程组的编号以及分配线程私有寄存器，所述经编译核函数包括被分别插入控制指令的多个对应程序段。
3.根据权利要求1所述的方法，还包括：
响应于所述对应程序段的所述控制指令被调用，执行以下各项：
使得为所述对应程序段所配置的线程组的所有线程同步；
为所配置的线程组的每一个线程分配线程私有寄存器；
重新配置所述每一个线程的编号，以便将所重新配置的编号写入所分配的线程私有寄存器。
4.根据权利要求1所述的方法，还包括：
经由驱动程序，启动经编译核函数；
经由所述编译器或者驱动程序，确认经编译核函数所需的最大线程数量；以及
经由驱动程序，向图形处理器请求所述最大线程数量。
5.根据权利要求1所述的方法，其中使得为所述对应程序段所配置的线程资源执行所述对应程序段的相关计算包括：
通过控制流代码控制执行流封装所述对应程序段，使得如果确定当前线程的线程标识属于为所述对应程序段所配置的线程组标识集合，以所述当前线程执行所述对应程序段的相关计算；以及
如果确定当前线程的线程标识不属于为所述对应程序段所配置的线程组标识集合，使得所述当前线程被跳过。
6.根据权利要求1所述的方法，其中确认待编译核函数中是否存在关于修改线程组请求的调用指令包括：
经由编译器，扫描待编译核函数，以确认待编译核函数中是否包括关于修改线程组请求的调用指令，所述关于修改线程组请求的调用指令指示对应程序段所需的线程配置。
7.一种计算设备，包括：
至少一个处理器；以及
与所述至少一个处理器通信连接的存储器；其中
所述存储器存储有可被所述至少一个处理器执行的指令，所述指令被所述至少一个处理器执行，以使所述至少一个处理器能够执行权利要求1-6中任一项所述的方法。
8.一种计算机可读存储介质，所述计算机可读存储介质上存储有计算机程序，所述计算机程序被机器执行时执行根据权利要求1-6中任一项所述的方法。
9.一种计算系统，包括：
编译器模块，被配置为识别所接收的待编译核函数中包括的关于修改线程组请求的调用指令，以及将所述调用指令的待编译核函数进行转换，以便生成经编译核函数，所述经编译核函数所包括的对应程序段配置有控制指令，关于修改线程组请求包括：关于修改线程组里面的线程数量和/或修改多个线程的组织拓扑结构的请求；
驱动程序模块，被配置为基于所述经编译核函数确定所述经编译核函数所需启动的最大线程数量；以及
执行单元计算核模块，被配置为响应于所述经编译核函数中的对应程序段的控制指令被执行，修改线程组的函数调用；
其中所述编译器模块还被配置为：如果确认待编译核函数中包括关于修改线程组请求的调用指令，为与所述调用指令相关联的对应程序段配置所需的线程资源；以及在所述对应程序段插入所述控制指令，以用于使得为所述对应程序段所配置的线程资源执行所述对应程序段的相关计算，并且未被配置的线程资源不执行相关计算；
其中，与所述调用指令相关联的对应程序段是通过以下确定的：基于调用指令的数量，将待编译核函数分成多个对应程序段，所述对应程序段的数量比调用指令的数量大1，并且为与所述调用指令相关联的对应程序段配置所需的线程资源包括：为所述多个对应程序段中的每一个对应程序段配置所需的线程组和线程私有寄存器源。
10.根据权利要求9所述的计算系统，其中所述驱动程序模块被配置在中央处理器，所述执行单元计算核模块被配置在图形处理器，所述中央处理器和图形处理器由所述计算系统包括。
11.根据权利要求9所述的计算系统，其中修改线程组的函数调用包括：重新恢复或者暂停线程、为每一个被配置的线程分配线程私有寄存器。
CN202110244679.9A 2021-03-05 2021-03-05 用于执行计算的方法、计算设备、计算系统和存储介质 Active CN113284038B (zh)

Priority Applications (2)

Application Number Priority Date Filing Date Title
CN202110244679.9A CN113284038B (zh) 2021-03-05 2021-03-05 用于执行计算的方法、计算设备、计算系统和存储介质
US17/686,413 US11886846B2 (en) 2021-03-05 2022-03-04 Method for executing computation, computing device, computing system, and storage medium

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN202110244679.9A CN113284038B (zh) 2021-03-05 2021-03-05 用于执行计算的方法、计算设备、计算系统和存储介质

Publications (2)

Publication Number Publication Date
CN113284038A CN113284038A (zh) 2021-08-20
CN113284038B true CN113284038B (zh) 2022-10-18

Family

ID=77276138

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202110244679.9A Active CN113284038B (zh) 2021-03-05 2021-03-05 用于执行计算的方法、计算设备、计算系统和存储介质

Country Status (2)

Country Link
US (1) US11886846B2 (zh)
CN (1) CN113284038B (zh)

Families Citing this family (5)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN114546329B (zh) * 2022-03-01 2023-07-18 上海壁仞智能科技有限公司 用于实现数据奇偶重排的方法、设备和介质
CN114968361B (zh) * 2022-06-02 2024-04-30 上海壁仞科技股份有限公司 存储有程序的机器可读介质、计算机系统和一种操作方法
CN115016953B (zh) * 2022-06-02 2024-03-22 上海壁仞智能科技有限公司 存储有程序的机器可读介质、计算机系统和一种操作方法
CN116302504A (zh) * 2023-02-23 2023-06-23 海光信息技术股份有限公司 线程块的处理系统、方法及相关设备
CN117591776B (zh) * 2024-01-18 2024-05-03 北京壁仞科技开发有限公司 用于计算的方法、计算装置、介质和程序产品

Family Cites Families (13)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US5692193A (en) * 1994-03-31 1997-11-25 Nec Research Institute, Inc. Software architecture for control of highly parallel computer systems
CN1278227C (zh) * 2004-06-25 2006-10-04 中国科学院计算技术研究所 一种基于mips指令集的处理器的多线程方法和装置
GB0613289D0 (en) * 2006-07-04 2006-08-16 Imagination Tech Ltd Synchronisation of execution threads on a multi-threaded processor
US8392669B1 (en) * 2008-03-24 2013-03-05 Nvidia Corporation Systems and methods for coalescing memory accesses of parallel threads
US8650554B2 (en) * 2010-04-27 2014-02-11 International Business Machines Corporation Single thread performance in an in-order multi-threaded processor
JPWO2012131933A1 (ja) * 2011-03-30 2014-07-24 富士通株式会社 情報処理装置、プログラム、および解析方法
US10599404B1 (en) * 2012-06-01 2020-03-24 Altera Corporation M/A for compiling parallel program having barrier synchronization for programmable hardware
CN104063507B (zh) * 2014-07-09 2017-10-17 时趣互动（北京）科技有限公司 一种图计算方法及系统
CN106484519B (zh) * 2016-10-11 2019-11-08 东南大学苏州研究院 异步线程重组方法及基于该方法的simt处理器
US10437593B2 (en) * 2017-04-27 2019-10-08 Nvidia Corporation Techniques for comprehensively synchronizing execution threads
US10877757B2 (en) * 2017-11-14 2020-12-29 Nvidia Corporation Binding constants at runtime for improved resource utilization
US10606595B2 (en) * 2018-03-23 2020-03-31 Arm Limited Data processing systems
US10831500B2 (en) * 2018-06-10 2020-11-10 International Business Machines Corporation Adaptive locking in elastic threading systems

Also Published As

Publication number Publication date
US20220283790A1 (en) 2022-09-08
CN113284038A (zh) 2021-08-20
US11886846B2 (en) 2024-01-30

Similar Documents

Publication Publication Date Title
CN113284038B (zh) 用于执行计算的方法、计算设备、计算系统和存储介质
KR102228501B1 (ko) 컴파일러 방법
KR101759266B1 (ko) 프로세서들에 걸쳐 데이터-병렬 쓰레드들을 지닌 프로세싱 로직을 매핑하는 방법
CN111831287B (zh) 用于确定执行代码段所需的资源的方法、设备和程序产品
CN110032395B (zh) 用于提高资源利用率的统一寄存器文件
US7937567B1 (en) Methods for scalably exploiting parallelism in a parallel processing system
JP2020518881A (ja) コンピュータに実装する方法、コンピュータ可読媒体および異種計算システム
CN111507476A (zh) 部署机器学习模型的方法、设备和计算机程序产品
WO2007084700A2 (en) System and method for thread handling in multithreaded parallel computing of nested threads
CN102402419A (zh) 用所选执行运行时执行的用户代码的运行时不可知表示
US11385931B2 (en) Method, electronic device, and computer program product for processing computing job
KR20190044572A (ko) 명령 세트
CN111259205B (zh) 一种图数据库遍历方法、装置、设备及存储介质
US20210286654A1 (en) Method, device and computer program product for processing computing job
US8595726B2 (en) Apparatus and method for parallel processing
WO2015099562A1 (en) Methods and apparatus for data-parallel execution of operations on segmented arrays
CN110866610A (zh) 一种深度学习模型分布式运算的方法及装置
CN101980147B (zh) 多线程处理器及其指令执行与同步方法
JP6551751B2 (ja) マルチプロセッサ装置
KR102228502B1 (ko) 컴퓨터 프로세싱의 타이밍 제어
CN113204412A (zh) 用于任务调度的方法、电子设备和计算机存储介质
CN111240745A (zh) 交叉执行的增强型标量向量双管线架构
CN110990151A (zh) 一种基于异构计算平台的业务处理方法
US20020156611A1 (en) Performance simulation process, and multiprocessor application production process, and devices for implementing said processes
JP6817827B2 (ja) アクセラレータ処理管理装置、ホスト装置、アクセラレータ処理実行システム、方法およびプログラム

Legal Events

Date Code Title Description
PB01 Publication
PB01 Publication
SE01 Entry into force of request for substantive examination
SE01 Entry into force of request for substantive examination
GR01 Patent grant
GR01 Patent grant
CP03 Change of name, title or address
CP03 Change of name, title or address

Address after: 201114 room 1302, 13 / F, building 16, 2388 Chenhang Road, Minhang District, Shanghai

Patentee after: Shanghai Bi Ren Technology Co.,Ltd.

Country or region after: China

Address before: 201114 room 1302, 13 / F, building 16, 2388 Chenhang Road, Minhang District, Shanghai

Patentee before: Shanghai Bilin Intelligent Technology Co.,Ltd.

Country or region before: China