CN116405155A - 利用gpu解码vdif格式数据的方法 - Google Patents

利用gpu解码vdif格式数据的方法 Download PDF

Info

Publication number
CN116405155A
CN116405155A CN202310190655.9A CN202310190655A CN116405155A CN 116405155 A CN116405155 A CN 116405155A CN 202310190655 A CN202310190655 A CN 202310190655A CN 116405155 A CN116405155 A CN 116405155A
Authority
CN
China
Prior art keywords
frame
data
gpu
block
decoding
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.)
Pending
Application number
CN202310190655.9A
Other languages
English (en)
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 Astronomical Observatory of CAS
Original Assignee
Shanghai Astronomical Observatory of CAS
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 Astronomical Observatory of CAS filed Critical Shanghai Astronomical Observatory of CAS
Priority to CN202310190655.9A priority Critical patent/CN116405155A/zh
Publication of CN116405155A publication Critical patent/CN116405155A/zh
Pending legal-status Critical Current

Links

Images

Classifications

    • H—ELECTRICITY
    • H04—ELECTRIC COMMUNICATION TECHNIQUE
    • H04L—TRANSMISSION OF DIGITAL INFORMATION, e.g. TELEGRAPHIC COMMUNICATION
    • H04L1/00—Arrangements for detecting or preventing errors in the information received
    • H04L1/0001—Systems modifying transmission characteristics according to link quality, e.g. power backoff
    • H04L1/0036—Systems modifying transmission characteristics according to link quality, e.g. power backoff arrangements specific to the receiver
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F17/00—Digital computing or data processing equipment or methods, specially adapted for specific functions
    • G06F17/10—Complex mathematical operations
    • G06F17/14—Fourier, Walsh or analogous domain transformations, e.g. Laplace, Hilbert, Karhunen-Loeve, transforms
    • G06F17/141—Discrete Fourier transforms
    • G06F17/142—Fast Fourier transforms, e.g. using a Cooley-Tukey type algorithm
    • 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
    • G06T—IMAGE DATA PROCESSING OR GENERATION, IN GENERAL
    • G06T1/00—General purpose image data processing
    • G06T1/60—Memory management
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06T—IMAGE DATA PROCESSING OR GENERATION, IN GENERAL
    • G06T2200/00—Indexing scheme for image data processing or generation, in general
    • G06T2200/28—Indexing scheme for image data processing or generation, in general involving image processing hardware
    • Y—GENERAL TAGGING OF NEW TECHNOLOGICAL DEVELOPMENTS; GENERAL TAGGING OF CROSS-SECTIONAL TECHNOLOGIES SPANNING OVER SEVERAL SECTIONS OF THE IPC; TECHNICAL SUBJECTS COVERED BY FORMER USPC CROSS-REFERENCE ART COLLECTIONS [XRACs] AND DIGESTS
    • Y02—TECHNOLOGIES OR APPLICATIONS FOR MITIGATION OR ADAPTATION AGAINST CLIMATE CHANGE
    • Y02D—CLIMATE CHANGE MITIGATION TECHNOLOGIES IN INFORMATION AND COMMUNICATION TECHNOLOGIES [ICT], I.E. INFORMATION AND COMMUNICATION TECHNOLOGIES AIMING AT THE REDUCTION OF THEIR OWN ENERGY USE
    • Y02D10/00—Energy efficient computing, e.g. low power processors, power management or thermal management

Landscapes

  • Engineering & Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • Mathematical Physics (AREA)
  • Theoretical Computer Science (AREA)
  • Pure & Applied Mathematics (AREA)
  • Mathematical Analysis (AREA)
  • Data Mining & Analysis (AREA)
  • Mathematical Optimization (AREA)
  • Computational Mathematics (AREA)
  • Quality & Reliability (AREA)
  • Discrete Mathematics (AREA)
  • Signal Processing (AREA)
  • Algebra (AREA)
  • Computer Networks & Wireless Communication (AREA)
  • Databases & Information Systems (AREA)
  • Software Systems (AREA)
  • General Engineering & Computer Science (AREA)
  • Communication Control (AREA)

Abstract

本发明涉及一种利用GPU解码VDIF格式数据的方法，包括：初始化第一并发队列、第二并发队列、环形缓冲区、CPU输入缓冲区、GPU输入缓冲区、CPU输出缓冲区和GPU输出缓冲区；读取VDIF格式数据，并将其写入第一并发队列；将第一并发队列中的VDIF格式数据按时间顺序排序，并将排序后的VDIF格式数据写入第二并发队列；将排序后的数据依次写入CPU输入缓冲区，然后拷贝至GPU输入缓冲区中，GPU对数据进行解码，并将解码后的数据存储在GPU输出缓冲区，然后将其拷贝到CPU输出缓冲区。本发明的利用GPU解码VDIF格式数据的方法，帧读取、帧重排和帧解码可以同时进行，大大减少处理时间，提高了解码速度。

Description

利用GPU解码VDIF格式数据的方法
技术领域
本发明涉及射电天文的基带信号处理领域，更具体地涉及一种利用GPU解码VDIF格式数据的方法。
背景技术
在射电天文的基带信号处理领域中，常用的基带数据格式有MarkIV、Mark5B和VDIF(VLBI Data Interface Format，VLBI数据接口格式)，其中VDIF是目前国际上VLBI领域使用最普遍的格式。VDIF格式的基带数据遵循一定的规则，每个数据帧包含帧头和帧体，其中帧头通常为固定的32字节，帧体长度则在帧头中定义。
在观测脉冲星时，主要有Searching(搜索)、Timing(计时)两种观测模式。在搜索模式时，通常先消除脉冲星信号发生的色散(DM)效应，然后进行傅里叶变换(FFT)获取周期，再按照周期进行折叠，最后验证候选体，验证成功即发现新的脉冲星。
在搜索过程中，会产生大量VDIF格式的基带数据，为了更好地搜索脉冲星，需要对产生的VDIF格式的基带数据进行实时解码。
现有技术中，采用CPU对VDIF格式的基带数据进行解码，其速度较慢，无法满足实时解码的要求。
发明内容
本发明的目的在于提供一种利用GPU解码VDIF格式数据的方法，由于GPU计算核心远大于CPU计算核心，因此可以快速对VDIF格式数据进行解码，满足实时解码的要求。
基于上述目的，本发明的目的在于提供一种利用GPU解码VDIF格式数据的方法，包括步骤：
S100：初始化第一并发队列、第二并发队列、环形缓冲区、CPU输入缓冲区、GPU输入缓冲区、CPU输出缓冲区和GPU输出缓冲区；
S200：从采集设备的输出端口读取VDIF格式数据，并将VDIF格式数据写入所述第一并发队列中；
S300：将所述第一并发队列中的VDIF格式数据按时间顺序排序，并将排序后的VDIF格式数据装填到所述第二并发队列中；
S400：将所述第二并发队列中的数据依次写入所述CPU输入缓冲区，然后将所述CPU输入缓冲区的数据拷贝至所述GPU输入缓冲区中，GPU对所述GPU输入缓冲区中的数据进行解码，并将解码后的数据存储在所述GPU输出缓冲区，然后将所述GPU输出缓冲区中的解码后的数据拷贝到所述CPU输出缓冲区。
进一步地，所述环形缓冲区被分为cqs个块，步骤S300进一步包括：
S310：从第一并发队列中读出当前帧的帧头，根据当前帧的帧头得到当前帧的帧体长度、帧秒和秒内帧序号；然后根据当前帧的帧体长度从第一并发队列中读取当前帧的帧体；
S320：判断当前帧是否为第一帧；
S330：若当前帧是第一帧，则将当前帧的帧头和帧体写入环形缓冲区的其中一块的首位，并将存储第一帧的环形缓冲区的块作为读取块，将读取块的首位存储的帧作为基准帧，然后回到步骤S310，以从第一并发队列中读取下一帧数据，并将下一帧作为当前帧；
S340：若当前帧不是第一帧，则根据当前帧的帧秒和秒内帧序号、基准帧的帧秒和秒内帧序号以及环形缓冲区的块数和块内可容纳帧数获得当前帧的相对块号和块内序号，判断相对块号是否小于0或大于cqs-1；
S350：若相对块号小于0或大于cqs-1，则回到步骤S310；否则，进一步判断相对块号是否等于cqs-1；
S360：若否，说明相对块号大于等于0且小于cqs-1，则根据相对块号和块内序号将当前帧的帧头和帧体写入环形缓冲区的对应位置，然后回到步骤S310；
若是，则根据相对块号和块内序号将当前帧的帧头和帧体写入环形缓冲区的对应位置；然后通过扫描读取块获得非法帧，并将读取块中的所有帧的帧体写入第二并发队列，在写入时，若遇到非法帧，则用非法帧的上一个合法帧的帧体替代非法帧的帧体写入第二并发队列中；写入完成后，将读取块的下一个块作为新的读取块，然后回到步骤S310。
进一步地，根据当前帧的帧秒和秒内帧序号、基准帧的帧秒和秒内帧序号以及环形缓冲区的块数和块内可容纳帧数获得当前帧的相对块号和块内序号，具体包括：
根据当前帧的帧秒和秒内帧序号、基准帧的帧秒和秒内帧序号以及环形缓冲区的块数和块内可容纳帧数获得当前帧与基准帧的帧差，将所述帧差除以块内可容纳帧数，并向下取整，得到当前帧的相对块号，将所述帧差对块内可容纳帧数取余，得到当前帧的块内序号。
进一步地，步骤S350还包括：在遇到非法帧时，将非法帧的帧秒和秒内帧序号写入丢帧文件中。
进一步地，cqs为3。
步骤S400进一步包括：
S410：创建输入流、解码流、输出流、输入事件和解码事件；
S420：从第二并发队列中读取第一预设长度的数据到CPU输入缓冲区；
S430：CPU异步调用输入流，以通过输入流从CPU输入缓冲区中拷贝第二预设长度的数据至GPU输入缓冲区，拷贝完成后由输入流触发一次输入事件，其中第一预设长度为第二预设长度的整数倍；
S440：CPU异步调用解码流等待输入事件，在输入事件触发后通过解码流执行解码核函数操作，对GPU输入缓冲区中的第二预设长度的数据进行解码，并将解码后的数据存储在GPU输出缓冲区中，解码完成后解码流触发一次解码事件；
S450：CPU异步调用输出流等待解码事件，在解码事件触发后将解码后的数据从GPU输出缓冲区拷贝到CPU输出缓冲区；
S460：重新执行步骤S440-步骤S450，直至CPU输入缓冲区中的数据全部被解码，然后回到步骤S420。
步骤S440进一步包括：
S441：将GPU输入缓冲区的数据与解码核函数的线程一一对应；
S442：解码核函数的各线程根据预设的量化方法对各线程对应的数据进行解码，并将解码后的数据存储至GPU输出缓冲区。
进一步地，预设的量化方法为通过预设的索引数组与VDIF格式数据进行映射或将二进制的VDIF格式数据转换为十进制。
进一步地，所述输入流、所述解码流和所述输出流设置为并行处理。
进一步地，所述第二并发队列的存储空间大于所述第一并发队列。
本发明的利用GPU解码VDIF格式数据的方法具有以下技术效果：
首先，充分考虑到脉冲星的VDIF格式数据的时间连续性，包括：1)通过将上一个合法帧替代非法帧的方式保证时间的连续性；2)通过丢帧文件记录非法帧的位置；传统的补帧方式是通过计算历史数据的平均值、标准差等值进行模拟，相比而言本方法的速度更快；由于丢帧是局部的情况，而脉冲星Searching模式并不要求对精细结构有很大需求，所以只要最大程度保证时间连续，中间小部分数据丢失并不会影响搜索；此外，在不知道脉冲星参数时，模拟丢帧也并不能起到很好的效果，而且因需要邻近数据的统计值，将花费很多时间，会导致丢帧的增加，本发明的方法可以避免这些问题。
其次，各数据缓冲区的大小都为固定的，易于方法的实施，即设备的采购；而且，避免了缓冲区不设上限而导致的问题，如缓冲区动态扩展的空间过大而导致系统崩溃。
再次，CPU、GPU的输入输出缓冲区为锁页内存，读写速度快且稳定，避免了统一内存的问题，还防止了内存拓展导致的安全性问题。
最后，GPU的输入流、解码流和输出流并行处理，CPU的帧读取、帧重排和帧解码也并行处理，大大减少了处理时间，提高了解码速度。
附图说明
图1为根据本发明的利用GPU解码VDIF格式数据的方法的流程图。
具体实施方式
下面结合附图，给出本发明的较佳实施例，并予以详细描述。
如图1所示，本发明提供一种利用GPU解码VDIF格式数据的方法，其协同CPU(Central Processing Unit，中央处理器)和GPU(Graphics Processing Unit，图形处理器)快速解码VDIF格式数据，其主要分为三个并行的阶段(帧读取、帧重排和帧解码)来提高处理速度，首先，在帧读取阶段接收由采集设备发送的VDIF格式数据，帧重排阶段则对VDIF格式数据按时间顺序进行排列，帧解码阶段则通过GPU核函数对VDIF格式数据进行解码，得到适用于FFT批处理的三维数组格式(FFT次数-通道-FFT频点)。该方法包括以下步骤：
S100：初始化第一并发队列、第二并发队列、环形缓冲区、CPU输入缓冲区、GPU输入缓冲区、CPU输出缓冲区和GPU输出缓冲区。
第一并发队列、第二并发队列、环形缓冲区、CPU输入缓冲区、GPU输入缓冲区、CPU输出缓冲区和GPU输出缓冲区为计算设备上的7个固定存储空间，作为本发明的方法使用的数据缓冲区。
第一并发队列作为帧读取阶段的输出空间，用于存储接收到的VDIF格式数据，其中VDIF格式数据的单位为帧，每一帧的VDIF格式数据均包括帧头和帧体；第二并发队列作为帧重排阶段的输出空间，用于存储有序的每一帧的VDIF格式数据的帧体数据。第一并发队列和第二并发队列的数据单元大小均可设置为1字节(char类型)。由于帧读取的速度设置为快于帧重排的速度，通常第二并发队列的最大可用空间需要大于第一并发队列。第一并发队列的最大可用空间至少需要大于两个帧的字节数目，这样可以降低帧读取写入和帧重排读取出现并发资源争夺的概率，进而提高处理速度。而第二并发队列的最大可用空间，可以至少等于两倍的CPU输入缓冲区大小，这样可以降低帧重排写入和帧解码读取出现并发资源争夺的概率，进而提高处理速度。并发队列支持数据有序、并发，其实现方法为本领域公知，此处不再赘述。
环形缓冲区用于在帧重排阶段存储重新排列的VDIF帧，环形缓冲区可由cqs个块组成，在一些实施例中，cqs可以为3。为了将1秒内的数据切分成N块(N为正整数)来批处理，每个块的数据个数可设置为nframes(nframes也为正整数，表示1/N秒的帧数量)，数据单元大小为VHS+VBS(char数组类型)，其中VHS为帧头的字节大小，VBS为帧体的字节大小。在观测模式固定时，接收机的带宽固定，因此采集设备发送的帧大小、秒内帧数一般也是固定的，nframes的取值取决于脉冲星观测模式使用的带宽、VDIF比特位数、VDIF帧体的字节大小、切分数N。如果要提高处理速度，通过调节N的数据使环形缓冲区的总大小不超过CPU一级、二级和三级缓存总和。N满足关系式(1)，nframes满足关系式(2)，式1中
Figure BDA0004105380250000063
表示对a向下取整，即小于a的最大整数。
Figure BDA0004105380250000061
Figure BDA0004105380250000062
例如，在一个示例性的实施例中，脉冲星观测模式使用的带宽为8GHz、VDIF数据比特位nibit为2，VHS为32字节，VBS为8000字节，CPU缓存总和为64M字节，则根据式(1)可得到N≤62500，N可取62500，根据式(2)可得nframes＝256。nframes必须为整数，若不是整数，则需调整N的值，使其为整数。
CPU输入缓冲区用于在帧解码阶段中在CPU上存储待解码数据(从第二并发队列中输出的排好序的数据)，数据个数为nbytes，数据单元大小为1字节(char类型)。
GPU输入缓冲区用于在帧解码阶段中在GPU上存储待解码数据(从CPU输入缓冲区拷贝而来)，数据个数为nbytes，数据单元大小为1字节(char类型)。
GPU输出缓冲区用于在帧解码阶段中在GPU上存储解码后的数据(采用GPU对GPU输入缓冲区的数据解码后得到的数据)，数据个数为8*nbytes/nibit，nibit为VDIF帧的帧体数据的比特位数，数据单元大小为nobit比特(通常为32比特位浮点类型或64比特位双精度浮点类型)。
CPU输出缓冲区用于在帧解码阶段中在CPU上存储解码后的数据(从GPU输出缓冲区拷贝而来)，数据个数为8*nbytes/nibit，数据单元大小为nobit比特(通常为32比特位浮点类型或64比特位双精度浮点类型)。
S200：帧读取阶段：从采集设备的输出端口读取VDIF格式数据，并将VDIF格式数据写入第一并发队列中。
在一些实施例中，可通过用户数据报协议(UDP，User Datagram Protocol)从采集设备的输出端口读取VDIF格式数据。具体地，可绑定输出端口，然后监听UDP数据包(即VDIF格式数据)，若没有数据到达，则继续等待，若有数据到达，则调用第一并发队列的接口(需保证线程安全)，将数据写入第一并发队列中。
S300：帧重排阶段：将第一并发队列中的VDIF格式数据按时间顺序排序，并将排序后的VDIF格式数据装填到第二并发队列中。
VDIF格式数据都是按时间顺序产生的，但是在将读取到的VDIF格式数据写入第一并发队列中时，可能并不会按时间顺序排列，例如，1秒内可能会产生很多帧的VDIF格式数据，第1秒内的第n帧数据是先于第1秒内的第n+1帧数据而产生的，但是在写入第一并发队列中时，第1秒内的第n+1帧数据可能会先于第1秒内的第n帧数据写入第一并发队列中，其不能直接用于FFT计算，因此，需要重新对其进行排序，使其按VDFI数据产生的时间顺序排列。
步骤S300进一步包括：
S310：从第一并发队列中读出当前帧的帧头，根据当前帧的帧头得到当前帧的帧体长度、帧秒和秒内帧序号；然后根据当前帧的帧体长度从第一并发队列中读取当前帧的帧体。
在从第一并发队列中读取VDIF数据时，是以帧为单位进行读取，由于每一帧的帧体的字节大小是固定的32字节，因此，可先从第一并发队列中读取32字节的数据，即为当前帧的帧头。帧头中记录了当前帧的帧体长度、帧秒和秒内帧序号，其中，帧秒是绝对时间，指的是当前帧是哪一秒产生的，秒内帧序号表示的是当前帧在该秒内的帧序号，即当前帧属于该秒内的第几帧。例如，若帧秒为10000，那么表示当前帧是距基准时间第10000秒产生的，若秒内帧序号为20，那么表示当前帧是该秒内的第20帧。根据帧秒、秒内帧序号以及1秒内产生的帧数可以得到当前帧的具体产生时间，例如，若1秒内产生的帧数为1000，那么帧秒为10000，秒内帧序号为20的当前帧的产生时间为(10000-1)+20/1000秒。由于已得到当前帧的帧体长度，从第一并发队列中读取当前帧的帧体长度的数据，即为当前帧的帧体。
S320：判断当前帧是否为第一帧。
从第一并发队列中读取的第一个数据即为第一帧，通常认为初始化第一并发队列、第二并发队列、环形缓冲区、CPU输入缓冲区、GPU输入缓冲区、CPU输出缓冲区和GPU输出缓冲区后读取的第一帧数据即为第一帧，后面的数据则不属于第一帧。在一些实施例中，可以设置一个初始值firstVDIF，在初始化后，将firstVDIF赋值为1，在从第一并发队列读取第一帧数据后，由于firstVDIF为1，说明该数据为第一帧数据，然后将firstVDIF赋值为0，在后续从第一并发队列读取数据时，均判定firstVDIF的值，若为0，说明该数据不是第一帧。
S330：若当前帧是第一帧，则将当前帧的帧头和帧体写入环形缓冲区的其中一块的首位，并将存储第一帧的环形缓冲区的块作为读取块，将读取块的首位存储的帧作为基准帧，然后回到步骤S310，以从第一并发队列中读取下一帧数据，并将下一帧作为当前帧。
由于当前帧为第一帧，可以将其帧头和帧体直接写入环形缓冲区中，第一帧的绝对时间即为排序的基准，也是数据的起点。由于环形缓冲区被分成cqs个块，每个块均有各自的编号，例如分别为0,1..cqs-1，每个块中可存储nframes帧数据，每个帧数据的存储位置也具有各自的块内序号，例如分别为1，2...nframes，分别表示该块内的第1、第2...第nframes帧的存储位置，在向环形缓冲区读数据或写数据时，会先从块内序号1的位置开始，然后依次读取，最开始读或写数据的地址即为块的首位。例如，当读指针指向缓冲区的0块时，会先从第0块的首位(即1号)开始读数据，然后依次读至nframes，完成该块数据读取。
当前帧为第一帧时，存储第一帧的块即为读取块，当需要从环形缓冲区中取数据时，先从读取块开始读取，读取块的首位存储的帧可以作为基准帧，基准帧用于判断是否需要将环形缓冲区的数据写入第二并发队列以及用于定位后续读取的帧存储在环形缓冲区内的地址(即位置)，从而保证写入第二并发队列中的VDIF格式数据是按时间顺序排列的。
S340：若当前帧不是第一帧，则根据当前帧的帧秒和秒内帧序号、基准帧的帧秒和秒内帧序号以及环形缓冲区的块数和块内可容纳帧数获得当前帧的相对块号和块内序号，判断相对块号是否小于0或大于cqs-1。
S350：若相对块号小于0或大于cqs-1，则回到步骤S310；否则，进一步判断相对块号是否等于cqs-1。
由于环形缓冲区的存储空间有限，若当前帧与基准帧的帧差超过其存储空间，那么当前帧将无法存入环形缓冲区中，说明数据异常，当前帧无法被排序，需要废弃，然后回到步骤S310，从第一并发队列读取下一帧。具体地，可通过当前帧的帧秒和秒内帧序号以及基准帧的帧秒和秒内帧序号计算得到当前帧与基准帧的帧差，然后用帧差除以块内可容纳帧数nframes，将余数作为块内序号，将结果向下取整则可得到当前帧的相对块号。若相对块号小于0，说明当前帧是在基准帧之前产生的，若相对块号大于cqs-1，则说明当前帧是在基准帧之后产生的，且相差的帧数超过环形缓冲区的存储空间，无法存储在环形缓冲区中。
S360：若否，说明相对块号大于等于0且小于cqs-1，则根据相对块号和块内序号将当前帧的帧头和帧体写入环形缓冲区的对应位置，然后回到步骤S310；若是(即相对块号等于cqs-1)，则根据相对块号和块内序号将当前帧的帧头和帧体写入环形缓冲区的对应位置；然后通过扫描读取块获得非法帧，并将读取块中的所有帧的帧体写入第二并发队列，在写入时，若遇到非法帧，则用非法帧的上一个合法帧的帧体替代非法帧的帧体写入第二并发队列中；写入完成后，将读取块的下一个块作为新的读取块，然后回到步骤S310。
若相对块号大于等于0且小于cqs-1，说明当前帧与基准帧的帧差没有超过其存储空间，可以存储在环形缓冲区中。此时相对块号表示的是存储当前帧的块与存储基准帧的块之间相差几块，当前帧存储在环形缓冲区的对应位置可根据相对块号和块内序号得到。例如，若基准帧存储在第0块中，相对块号为1，那么存储当前帧的块与存储基准帧的块之间相差1个块，因此当前帧需存储在第1块中；若基准帧存储在第cqs-1块中，相对块号为2，那么存储当前帧的块与存储基准帧的块之间相差2个块，由于各块形成一个环形，第cqs-1块之后的2块依次为第0块和第1块，因此，当前帧需存储在第1块中；以此类推。块内序号是当前帧存储在块内的序号，若块内序号为2，那么当前帧需要存储在该块内的第2帧的位置。
若相对块号等于cqs-1，说明已有cqs-1个块存储有VDIF数据，可以将当前帧存储在环形缓冲区中，但是同时需要将读取块内的数据写入第二并发队列，从而将读取块释放，便于后续的数据存入。由于在从第一并发队列中读取数据时，可能会跳过一些不符合要求的帧(参见步骤330)，因此，读取块中的数据可能并不完整，即有些位置并没有写入数据。例如，假设N＝2，1秒内产生的数据为30帧，则nframes＝15，每个块内可以存储15帧数据，基准帧存储在读取块的首位，通常期望将基准帧的下一帧(秒内帧序号的下一位)存储在读取块的第2帧的位置，从而使其按时间顺序排列，但是基准帧的下一帧可能出现在基准帧之前，因此已经被废弃，那么在读取块的第2帧位置不会存储数据，此时基准帧的下一帧即为缺失帧或非法帧，而其他存储在读取块内的帧则为合法帧，通过扫描读取块，可得到所有非法帧，然后将非法帧的信息(包括帧秒和秒内帧序号)记录在丢帧文件中。由于第二并发队列中的数据是用于后续的解码，因此需要保证其数据是按时间顺序排列且是连续无缺失的，而在将读取块中的数据写入第二并发队列时，由于非法帧对应的存储位置没有数据，若直接按读取块中的数据写入，会造成第二并发队列中的数据缺失，导致时间不连续。因此在将读取块中的数据写入第二并发队列时，若遇到非法帧，可用非法帧的上一个合法帧替代非法帧写入至第二并发队列中，例如，若读取块中的第2帧位置、第10帧的位置为空，其余位置都存储有数据，那么在将读取块中的数据写入第二并发队列时，先从基准帧(即存储在读取块首位的帧)开始，将基准帧的帧体写入第二并发列，然后用基准帧的帧体替代读取块中的第2帧的位置对应的非法帧的帧体，将其写入第二并发队列，然后依次将第3帧、第4帧....第9帧的位置存储的帧的帧体写入第二并发队列，然后用第9帧的位置存储的帧的帧体替代第10帧的位置对应的非法帧的帧体，将其写入第二并发队列中，以此类推。最后在第二并发队列中的排列顺序为基准帧的帧体、基准帧的帧体、第3帧的帧体...第9帧的帧体、第9帧的帧体、第11帧的帧体...第N帧的帧体。这样，第二并发队列中的数据是严格按照时间顺序排列的，且没有缺失数据。
丢帧文件的记录格式可以设置为：一行一个丢帧记录，行中的帧秒、帧号用逗号隔开，写入方法通过操作系统的aio方式写入即可，其为公知方法，此处不再赘述。通过丢帧文件，可以明确丢帧的时间信息，便于在对解码后的数据进行处理(例如FFT)时知晓有哪些时间点的帧丢了。
读取块中的数据读取完后，将读取块的下一个块作为新的读取块，然后回到步骤S310，直至观测结束。
在一些实施例中，若环形缓冲区的块编号依次为0,1..cqs-1，则步骤S310-步骤S350可通过下述方法实现(方法中的“＝”号代表赋值)：
S310a：从第一并发队列中读出32字节的数据，即为当前帧的帧头，根据帧头得到当前帧的帧体长度、帧秒fsec和秒内帧序号fnum，然后再从第一并发队列中读取当前帧的帧体长度大小的数据，即为第一并发队列的帧体；
S320a：若当前帧为此次观测的第一帧，则设置fsec0＝fsec，fnum0＝fnum、blk＝0，并将当前帧的帧头和帧体写入环形缓冲区的第0块的首位，blk0用于标记当前状态下环形缓冲区的读取块和统计本次观测所重排的块数；
S330a：若当前帧不是此次观测的第一帧，则通过下式(3)和(4)计算得到当前帧的相对块号iblk、块内序号iframe，式中MOD表示取余运算符；
Figure BDA0004105380250000111
iframe＝帧差值MOD块内可容纳帧数＝[(fsec-fsec0)×nframes×N+fnum-fnum0]MOD[nframes] (4)
其中，fsec为当前帧的帧秒，fsec0为基准帧的帧秒，fnum为当前帧的秒内帧序号，fnum0为基准帧的秒内帧序号
若相对块号iblk小于0或大于cqs-1，则说明数据异常，跳过该帧，重新回到步骤S310a。
S340a：若相对块号iblk大于等于0且小于cqs-1，计算存储当前帧的块的索引位置blkn＝[blk0+iblk]MOD cqs，然后将当前帧的帧头和帧体写入环形缓冲区的第blkn块第iframes帧的位置，然后回到步骤S310a；
S350a：若相对块号iblk等于cqs-1，计算存储当前帧的块的索引位置blkn＝[blk0+iblk]MOD cqs，然后将当前帧的帧头和帧体写入环形缓冲区的第blkn块第iframes帧的位置，然后计算读取块的实际索引号(blk0+iblk)MOD cqs，根据实际索引号扫描读取块并获得非法帧，并将非法帧信息记录在丢帧文件中，然后将读取块中的所有帧的帧体写入第二并发队列，在写入时，若遇到非法帧，则用非法帧的上一个合法帧的帧体替代非法帧的帧体写入第二并发队列中；写入完成后，令blk0＝blk0+1，fsec0＝fsec0+(fnum0+nframes)/(nframes*N)，fnum0＝(fnum0+nframes)mod(nframes*N)，然后回到步骤S310a，如此循环往复直至本次观测结束。
在一个示例性的实施例中，1秒内总帧数为30，N＝2，cqs＝3，环形缓冲区的每块可容纳帧数nframes＝30/N＝15，环形缓冲区的块编号(即实际索引号)分别为0、1和2，每秒的秒内帧序号为0＝29，第一帧(基准帧)的帧秒fsec1＝0，帧序号fnum1＝0，则fsec0＝fsec1＝0，fnum0＝fnum1＝0，blk0＝0。按顺序到来了五个帧，分别是帧秒0且秒内帧序号9、帧秒1且秒内帧序号25、帧秒1且秒内帧序号11、帧秒1且秒内帧序号21、帧秒1且秒内帧序号22，其处理情况如下步骤：
1)当前帧的帧秒fsec＝0，秒内帧序号fnum＝9，当前帧与基准帧的帧差为[(0-0)×15×2+9-0]＝9，所以iframe＝9mod 15＝9，
Figure BDA0004105380250000121
由于0不等于cqs-1，因此不用释放，实际索引块号为(0+0)mod 3＝0，直接存到第0块，第9帧的位置；然后继续读取下一帧：帧秒1且秒内帧序号25。
2)当前帧的帧秒fsec＝1，秒内帧序号fnum＝25，则其与基准帧的帧差为[(1-0)×15×2+25-0]＝55，所以iframe＝55mod 10＝5，
Figure BDA0004105380250000122
由于3大于cqs-1，超出当前存储范围，为异常数据，跳过该帧，继续读取下一帧：帧秒1且秒内帧序号11。
3)当前帧的帧秒fsec＝1，秒内帧序号fnum＝11，则其与基准帧的帧差为[(1-0)×15×2+11-0]＝41，所以iframe＝41mod 15＝11，
Figure BDA0004105380250000123
Figure BDA0004105380250000124
由于2等于cqs-1，表示有两个块已满，需释放读取块，读取块的实际块编号为blk0 mod cqs＝0mod 3＝0，读取块中的数据按时间顺序写入第二并发队列中，然后清空第0块，并使blk0为0+1＝1，fesc0为0+(0+15)/(15*2)＝0，fnum0为0+15＝15；然后继续读取下一帧：帧秒1且秒内帧序号21。
4)当前帧的帧秒fsec＝1，秒内帧序号fnum＝21，该帧与基准帧的帧差为[(1-0)×15×2+21-15]＝36，所以iframe＝36mod 15＝6，
Figure BDA0004105380250000131
Figure BDA0004105380250000132
由于2等于cqs-1，表示有两个块已满，需释放读取块，读取块的的实际块号为blk0 mod cqs＝1mod 3＝1，读取块中的数据按时间顺序写入第二并发队列中，然后清空第0块，并使blk0为1+1＝2，fesc0为0+(15+15)/(15*2)＝1，fnum0为(15+15)mod(15*2)＝0；然后继续读取下一帧：帧秒1且秒内帧序号22。
5)当前帧的帧秒fsec＝1、帧序号fnum＝22，则该帧与第一帧的帧差为[(1-1)×15×2+22-0]＝22，所以iframe＝22mod 15＝7，
Figure BDA0004105380250000133
由于1不等于cqs-1，因此不用释放，实际索引块号为(2+1)mod 3＝0，直接将当前帧存到第0块，第7帧的位置；然后继续读取下一帧。
步骤S400：帧解码阶段：将第二并发队列中的数据依次写入CPU输入缓冲区，然后将CPU输入缓冲区的数据拷贝至GPU输入缓冲区中，GPU对GPU输入缓冲区中的数据进行解码，并将解码后的数据存储在GPU输出缓冲区，然后将GPU输出缓冲区中的解码后的数据拷贝到CPU输出缓冲区。
帧解码阶段主要是对第二并发队列中的连续的VDIF格式数据进行解码处理。由于GPU无法直接从第二并发队列中读取数据，因此先将第二并发队列中的数据依次写入CPU输入缓冲区，然后再拷贝至GPU输入缓冲区，以利用GPU解码GPU输入缓冲区中的数据，解码后的数据存储在GPU输出缓冲区，其可被拷贝到CPU输出缓冲区，以供后续使用。
步骤S400进一步包括：
S410：创建输入流、解码流、输出流、输入事件和解码事件。
流，也称为CUDA(Compute Unified Device Architecture，统一计算设备架构)流，是一种由NVIDIA推出的通用并行计算架构，该架构使GPU能够解决复杂的计算问题。流是GPU上的工作队列，队列里的工作将以特定的顺序执行，这些工作可以包括:内核函数的调用，cudaMemcpy系列传输，以及对CUDA事件的操作。各个流添加到队列的顺序将决定它们的执行顺序。每个CUDA流可以被视为单个任务，因此可以启动多个流来并行执行多个任务，在很多情况下，多个流中的工作可能同时执行。
在一些实施例中，输入流用于执行从CPU到GPU的内存拷贝操作的核函数，以将CPU输入缓冲区中的数据拷贝到GPU输入缓冲区；解码流用于执行解码相关的核函数，以对GPU输入缓冲区中的数据进行解码；输出流用于执行从GPU到CPU内存拷贝操作的核函数，以将GPU输出缓冲区的解码后的数据拷贝到CPU输出缓冲区；输入事件为输入完成后由输入流触发、由解码流等待的事件；解码事件为解码完成后由解码流触发、由输出流等待的事件。
S420：从第二并发队列中读取第一预设长度的数据到CPU输入缓冲区。
CPU输入缓冲区可储存数据的字节大小为nbytes，因此第一预设长度为nbytes字节。
S430：CPU异步调用输入流，以通过输入流从CPU输入缓冲区中拷贝第二预设长度的数据至GPU输入缓冲区，拷贝完成后由输入流触发一次输入事件，其中第一预设长度为第二预设长度的整数倍。
为了提高解码速度，可将CPU输入缓冲区中的数据分为m份，以实现并行处理，即第二预设长度可以为nbytes/m字节。在一些实施例中，nbytes和m需满足以下条件：1)nbytes可以被m整除；2)nbytes/m可以被nbin整除，其中nbin为FFT点数值，即后续对解码后的数据进行FFT时所使用的数据点数，通常nbin＝65536(或者2的其他次方数)，nfft＝nbytes/m/nbin表示并行处理中当行的数据量能够用于执行FFT的次数。
S440：CPU异步调用解码流等待输入事件，在输入事件触发后通过解码流执行解码核函数操作，对GPU输入缓冲区中的第二预设长度的数据进行解码，并将解码后的数据存储在GPU输出缓冲区中，解码完成后解码流触发一次解码事件。
异步调用是指CPU调用输入流后，不用等待输入流拷贝完成，而是直接调用解码流，因此，在输入流拷贝数据的过程中，解码流已经启动并在等待输入事件，输入事件被触发后，解码流将立即开始解码。
解码流通过预设的解码核函数来进行解码。在一些实施例中，解码核函数的线程块可以为CUDA提供的三维线程块、三维网格块结构。线程块格式如(blocksize.x，blocksize.y，blocksize.z)，具体数值由具体使用显卡型号决定(保证不超过最大支持量)，通常可取值为(1024，32，32)。GPU的网格块格式如(gridsize.x，gridsize.y，gridsize.z)，可设置其x为
Figure BDA0004105380250000152
Figure BDA0004105380250000153
(向上取整)、y为
Figure BDA0004105380250000151
z为/>
Figure BDA0004105380250000154
解码核函数的解码方式如下：
S441：将GPU输入缓冲区的数据与解码核函数的线程块一一对应。
具体地，可根据解码核函数的线程块格式和网格块格式计算当前线程所计算数据所属的通道号ichan、数据所属的FFT内的序号ibin以及数据点所属的FFT次数ifft，其计算公式如下：
ibin＝blockIdx.x*blockDim.x+threadIdx.x (5)
ifft＝blockIdx.y*blockDim.y+threadIdx.y (6)
ichan＝blockIdx.z*blockDim.z+threadIdx.z (7)
S442：解码核函数的各线程根据预设的量化方法对各线程对应的数据进行解码，并将解码后的数据存储至GPU输出缓冲区。
具体地，可根据ichan、ibin和ifft得到解码后的数据的索引位置idx1、待解码数据的存储的索引位置idx2，然后通过idx1和idx2计算得到中间值ivalue，idx1、idx2和ivalue的计算公式如下：
idx1＝ibin+nbin*ichan+nchan*nbin*ifft (8)
idx2＝ichan+nchan*ibin+nchan*nbin*ifft (9)
ibyte＝idx2＞＞nibit (10)
ibits＝nibit*(idx2&((1＜＜nibit)-1))) (11)
ivalue＝(nbytesBuf[ibyte]＞＞ibits)&((1＜＜nibit)-1)) (12)
其中，＞＞、＜＜和&均为位运算符，＞＞为右移，＜＜为左移，&为位与。
通过预设的量化方法可将ivlaue解码，然后通过idx1将解码后的数据映射到GPU输出缓冲区。量化方法是指通过既定的索引数组与VDIF格式数据进行映射，从而对VDIF格式数据进行解码。例如，VDIF格式数据通常是2比特，其为二进制，量化后的数值为0，1，2或3，其索引数组mask可以为[-3.3359，-1，1，3.3359]，那么在解码时，可将VDIF格式数据中的0解码为-3.3359，1则解码为-1，2则解码为1，3则解码为3.3359，从而得到解码后的数据。在一些实施例中，量化方式还可以直接将二进制的VDIF格式数据量化为十进制数值，作为解码后的数据，例如，若VDIF格式数据为01，那么解码后的数据为1，若为11，解码后的数据为3。
对于采用索引数组进行解码时，可通过下式得到解码后的数据并将其映射到GPU输出缓冲区：
Buffergo[idx1]＝mask[ivalue] (13)
Buffergo[idx1]是指将公式右边的值映射到GPU输出缓冲区的idx1的索引位置。
对于直接将二进制的VDIF格式数据量化为十进制进行解码时，可通过下式得到解码后的数据并将其映射到GPU输出缓冲区：
Buffergo[idx1]＝ivalue-(1＜＜(nibit-1)) (14)
S450：CPU异步调用输出流等待解码事件，在解码事件触发后将解码后的数据从GPU输出缓冲区拷贝到CPU输出缓冲区。
CPU调用解码流后，即调用输出流，输出流的工作可以与解码流和输入流同时进行。
S460：重新执行步骤S440-步骤S450，直至CPU输入缓冲区中的数据全部被解码，然后回到步骤S420。
由于在帧解码阶段，将CPU输入缓冲区中的数据分成整数份拷贝至GPU输入缓冲区，而输入流、解码流和输出流又可以并行处理，使输入、解码与输出同时进行，减少了计算时间，提高了计算速度。
在一些实施例中，帧读取阶段可由帧读取模块来实现，帧重排阶段可由帧重排模块来实现，帧解码阶段可由帧解码模块来实现。在一些实施例中，在利用本发明的方法解码VDIF格式数据时，可先依次启动帧解码模块、帧重排模块和帧读取模块，当帧解码模块获取不到CPU输入缓冲区中的数据时，则会等待并向帧读取模块发送信号，使帧重排模块将排列好的数据存入CPU输入缓冲区中，而当帧重排模块无法从第一并发队列获取数据时，则会等待并向帧读取模块发送信号，使帧读取模块从采集设备读取数据。
帧读取、帧重排和帧解码也可以实现并行处理，即在帧解码时，帧读取模块同时在读取新的数据，而帧重排模块也同时在对数据进行排列，从而让读取、重排和解码时间大量重叠，提高了解码效率。
本发明的利用GPU解码VDIF格式数据的方法具有以下技术效果：
首先，充分考虑到脉冲星的VDIF格式数据的时间连续性，包括：1)通过将上一个合法帧替代非法帧的方式保证时间的连续性；2)通过丢帧文件记录非法帧的位置；传统的补帧方式是通过计算历史数据的平均值、标准差等值进行模拟，相比而言本方法的速度更快；由于丢帧是局部的情况，而脉冲星Searching模式并不要求对精细结构有很大需求，所以只要最大程度保证时间连续，中间小部分数据丢失并不会影响搜索；此外，在不知道脉冲星参数时，模拟丢帧也并不能起到很好的效果，而且因需要邻近数据的统计值，将花费很多时间，会导致丢帧的增加，本发明的方法可以避免这些问题。
其次，各数据缓冲区的大小都为固定的，易于方法的实施，即设备的采购；而且，避免了缓冲区不设上限而导致的问题，如缓冲区动态扩展的空间过大而导致系统崩溃。
再次，CPU、GPU的输入输出缓冲区为锁页内存，读写速度快且稳定，避免了统一内存的问题，还防止了内存拓展导致的安全性问题。
最后，GPU的输入流、解码流和输出流并行处理，CPU的帧读取、帧重排和帧解码也并行处理，大大减少了处理时间，提高了解码速度。
以上所述的，仅为本发明的较佳实施例，并非用以限定本发明的范围，本发明的上述实施例还可以做出各种变化。即凡是依据本发明申请的权利要求书及说明书内容所作的简单、等效变化与修饰，皆落入本发明专利的权利要求保护范围。本发明未详尽描述的均为常规技术内容。

Claims (10)

1.一种利用GPU解码VDIF格式数据的方法，其特征在于，包括步骤：
S100：初始化第一并发队列、第二并发队列、环形缓冲区、CPU输入缓冲区、GPU输入缓冲区、CPU输出缓冲区和GPU输出缓冲区；
S200：从采集设备的输出端口读取VDIF格式数据，并将VDIF格式数据写入所述第一并发队列中；
S300：将所述第一并发队列中的VDIF格式数据按时间顺序排序，并将排序后的VDIF格式数据装填到所述第二并发队列中；
S400：将所述第二并发队列中的数据依次写入所述CPU输入缓冲区，然后将所述CPU输入缓冲区的数据拷贝至所述GPU输入缓冲区中，GPU对所述GPU输入缓冲区中的数据进行解码，并将解码后的数据存储在所述GPU输出缓冲区，然后将所述GPU输出缓冲区中的解码后的数据拷贝到所述CPU输出缓冲区。
2.根据权利要求1所述的利用GPU解码VDIF格式数据的方法，其特征在于，所述环形缓冲区被分为cqs个块，步骤S300进一步包括：
S310：从第一并发队列中读出当前帧的帧头，根据当前帧的帧头得到当前帧的帧体长度、帧秒和秒内帧序号；然后根据当前帧的帧体长度从第一并发队列中读取当前帧的帧体；
S320：判断当前帧是否为第一帧；
S330：若当前帧是第一帧，则将当前帧的帧头和帧体写入环形缓冲区的其中一块的首位，并将存储第一帧的环形缓冲区的块作为读取块，将读取块的首位存储的帧作为基准帧，然后回到步骤S310，以从第一并发队列中读取下一帧数据，并将下一帧作为当前帧；
S340：若当前帧不是第一帧，则根据当前帧的帧秒和秒内帧序号、基准帧的帧秒和秒内帧序号以及环形缓冲区的块数和块内可容纳帧数获得当前帧的相对块号和块内序号，判断相对块号是否小于0或大于cqs-1；
S350：若相对块号小于0或大于cqs-1，则回到步骤S310；否则，进一步判断相对块号是否等于cqs-1；
S360：若否，说明相对块号大于等于0且小于cqs-1，则根据相对块号和块内序号将当前帧的帧头和帧体写入环形缓冲区的对应位置，然后回到步骤S310；
若是，则根据相对块号和块内序号将当前帧的帧头和帧体写入环形缓冲区的对应位置；然后通过扫描读取块获得非法帧，并将读取块中的所有帧的帧体写入第二并发队列，在写入时，若遇到非法帧，则用非法帧的上一个合法帧的帧体替代非法帧的帧体写入第二并发队列中；写入完成后，将读取块的下一个块作为新的读取块，然后回到步骤S310。
3.根据权利要求2所述的利用GPU解码VDIF格式数据的方法，其特征在于，根据当前帧的帧秒和秒内帧序号、基准帧的帧秒和秒内帧序号以及环形缓冲区的块数和块内可容纳帧数获得当前帧的相对块号和块内序号，具体包括：
根据当前帧的帧秒和秒内帧序号、基准帧的帧秒和秒内帧序号以及环形缓冲区的块数和块内可容纳帧数获得当前帧与基准帧的帧差，将所述帧差除以块内可容纳帧数，并向下取整，得到当前帧的相对块号，将所述帧差对块内可容纳帧数取余，得到当前帧的块内序号。
4.根据权利要求2所述的利用GPU解码VDIF格式数据的方法，其特征在于，步骤S350还包括：在遇到非法帧时，将非法帧的帧秒和秒内帧序号写入丢帧文件中。
5.根据权利要求2所述的利用GPU解码VDIF格式数据的方法，其特征在于，cqs为3。
6.根据权利要求1所述的利用GPU解码VDIF格式数据的方法，其特征在于，步骤S400进一步包括：
S410：创建输入流、解码流、输出流、输入事件和解码事件；
S420：从第二并发队列中读取第一预设长度的数据到CPU输入缓冲区；
S430：CPU异步调用输入流，以通过输入流从CPU输入缓冲区中拷贝第二预设长度的数据至GPU输入缓冲区，拷贝完成后由输入流触发一次输入事件，其中第一预设长度为第二预设长度的整数倍；
S440：CPU异步调用解码流等待输入事件，在输入事件触发后通过解码流执行解码核函数操作，对GPU输入缓冲区中的第二预设长度的数据进行解码，并将解码后的数据存储在GPU输出缓冲区中，解码完成后解码流触发一次解码事件；
S450：CPU异步调用输出流等待解码事件，在解码事件触发后将解码后的数据从GPU输出缓冲区拷贝到CPU输出缓冲区；
S460：重新执行步骤S440-步骤S450，直至CPU输入缓冲区中的数据全部被解码，然后回到步骤S420。
7.根据权利要求6所述的利用GPU解码VDIF格式数据的方法，其特征在于，步骤S440进一步包括：
S441：将GPU输入缓冲区的数据与解码核函数的线程一一对应；
S442：解码核函数的各线程根据预设的量化方法对各线程对应的数据进行解码，并将解码后的数据存储至GPU输出缓冲区。
8.根据权利要求7所述的利用GPU解码VDIF格式数据的方法，其特征在于，预设的量化方法为通过预设的索引数组与VDIF格式数据进行映射或将二进制的VDIF格式数据转换为十进制。
9.根据权利要求6所述的利用GPU解码VDIF格式数据的方法，其特征在于，所述输入流、所述解码流和所述输出流设置为并行处理。
10.根据权利要求1所述的利用GPU解码VDIF格式数据的方法，其特征在于，所述第二并发队列的存储空间大于所述第一并发队列。
CN202310190655.9A 2023-03-01 2023-03-01 利用gpu解码vdif格式数据的方法 Pending CN116405155A (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN202310190655.9A CN116405155A (zh) 2023-03-01 2023-03-01 利用gpu解码vdif格式数据的方法

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN202310190655.9A CN116405155A (zh) 2023-03-01 2023-03-01 利用gpu解码vdif格式数据的方法

Publications (1)

Publication Number Publication Date
CN116405155A true CN116405155A (zh) 2023-07-07

Family

ID=87009296

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202310190655.9A Pending CN116405155A (zh) 2023-03-01 2023-03-01 利用gpu解码vdif格式数据的方法

Country Status (1)

Country Link
CN (1) CN116405155A (zh)

Cited By (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN116795756A (zh) * 2023-08-28 2023-09-22 成都量芯集成科技有限公司 一种基于串口通讯的帧数据流软件接收方法

Cited By (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN116795756A (zh) * 2023-08-28 2023-09-22 成都量芯集成科技有限公司 一种基于串口通讯的帧数据流软件接收方法
CN116795756B (zh) * 2023-08-28 2023-10-27 成都量芯集成科技有限公司 一种基于串口通讯的帧数据流软件接收方法

Similar Documents

Publication Publication Date Title
CN112084136B (zh) 队列缓存管理方法、系统、存储介质、计算机设备及应用
US6725225B1 (en) Data management apparatus and method for efficiently generating a blocked transposed file and converting that file using a stored compression method
KR20210086420A (ko) 신경망 데이터 처리 장치, 방법 및 전자 장비
CN116405155A (zh) 利用gpu解码vdif格式数据的方法
WO1994017470A1 (en) System for dynamically allocating memory registers for forming pseudo queues
CN105573711B (zh) 一种数据缓存方法及装置
US20080155353A1 (en) Data processing apparatus and method for reducing trace bandwidth
CN113015003B (zh) 视频帧缓存方法和设备
US20070283131A1 (en) Processing of high priority data elements in systems comprising a host processor and a co-processor
CN115438114B (zh) 存储格式转换方法、系统、装置、电子设备及存储介质
US7266650B2 (en) Method, apparatus, and computer program product for implementing enhanced circular queue using loop counts
WO2018022303A1 (en) Capturing commands in a multi-engine graphics processing unit
CN111343404B (zh) 成像数据处理方法及装置
Bisson et al. A cuda implementation of the pagerank pipeline benchmark
JPH0234038A (ja) データ圧縮装置
CN112256632B (zh) 一种可重构处理器中的指令分发方法及系统
CN111370070B (zh) 一种针对大数据基因测序文件的压缩处理方法
US6816923B1 (en) Arbitrating and servicing polychronous data requests in direct memory access
US20060072840A1 (en) Conversion device for performing a raster scan conversion between a JPEG decoder and an image memory
CN115328923B (zh) 时序生理数据的存储结构、查询方法、存储介质及系统
JPH02173846A (ja) 循環バツフア制御装置及び制御方法
CN117666963A (zh) 一种cpu云计算平台的数据io加速方法
CN118170702B (zh) Dma控制器以及用于广播的数据搬运方法
US6795879B2 (en) Apparatus and method for wait state analysis in a digital signal processing system
US20040111247A1 (en) Optimization of timing models using bus compression

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