Method for decoding VDIF format data using GPU

Abstract

The present invention relates to a method for decoding VDIF format data using a GPU, which includes: initializing a first concurrent queue, a second concurrent queue, a circular buffer, a CPU input buffer, a GPU input buffer, a CPU output buffer, and a GPU output buffer; reading VDIF format data and writing it into the first concurrent queue; sorting the VDIF format data in the first concurrent queue in chronological order and writing the sorted VDIF format data into the second concurrent queue; sequentially writing the sorted data into the CPU input buffer and then copying it into the GPU input buffer, where the GPU decodes the data and stores the decoded data in the GPU output buffer, and then copies it to the CPU output buffer. In the method for decoding VDIF format data using a GPU according to the present invention, frame reading, frame rearrangement, and frame decoding can be performed simultaneously, greatly reducing the processing time and improving the decoding speed.

Description

Technical Field

[0001] The present invention relates to the field of baseband signal processing in radio astronomy, and more particularly to a method for decoding VDIF format data using a GPU. Background Art

[0002] In the field of baseband signal processing in radio astronomy, common baseband data formats include MarkIV, Mark5B, and VDIF (VLBI Data Interface Format). Among them, VDIF is the most commonly used format in the international VLBI field currently. The baseband data in VDIF format follows certain rules. Each data frame contains a frame header and a frame body. The frame header is usually 32 bytes fixed, and the length of the frame body is defined in the frame header.

[0003] When observing pulsars, there are mainly two observation modes: Searching and Timing. In the search mode, usually, the dispersion (DM) effect of the pulsar signal is first eliminated, then Fourier transform (FFT) is performed to obtain the period, then folding is carried out according to the period, and finally, the candidate is verified. If the verification is successful, a new pulsar is discovered.

[0004] During the search process, a large amount of baseband data in VDIF format is generated. In order to better search for pulsars, it is necessary to perform real-time decoding on the generated baseband data in VDIF format.

[0005] In the prior art, the CPU is used to decode the baseband data in VDIF format, and its speed is slow, unable to meet the requirements of real-time decoding. Summary of the Invention

[0006] The purpose of the present invention is to provide a method for decoding VDIF format data using a GPU. Since the computing core of the GPU is much larger than that of the CPU, the VDIF format data can be decoded quickly to meet the requirements of real-time decoding.

[0007] Based on the above purpose, the present invention aims to provide a method for decoding VDIF format data using a GPU, including the steps of:

[0008] S100: Initialize the first concurrent queue, the second concurrent queue, the circular buffer, the CPU input buffer, the GPU input buffer, the CPU output buffer, and the GPU output buffer;

[0009] S200: Read the VDIF format data from the output port of the acquisition device and write the VDIF format data into the first concurrent queue;

[0010] S300: Sort the VDIF - formatted data in the first concurrent queue in chronological order, and load the sorted VDIF - formatted data into the second concurrent queue;

[0011] S400: Write the data in the second concurrent queue into the CPU input buffer in sequence, then copy the data in the CPU input buffer to the GPU input buffer. The GPU decodes the data in the GPU input buffer and stores the decoded data in the GPU output buffer, and then copies the decoded data in the GPU output buffer to the CPU output buffer.

[0012] Further, the circular buffer is divided into cqs blocks, and step S300 further includes:

[0013] S310: Read the frame header of the current frame from the first concurrent queue, obtain the frame body length, frame second, and frame sequence number within the second of the current frame according to the frame header of the current frame; then read the frame body of the current frame from the first concurrent queue according to the frame body length of the current frame;

[0014] S320: Determine whether the current frame is the first frame;

[0015] S330: If the current frame is the first frame, write the frame header and frame body of the current frame to the beginning of one of the blocks in the circular buffer, and use the block in the circular buffer storing the first frame as the read block, and use the frame stored at the beginning of the read block as the reference frame, then return to step S310 to read the next - frame data from the first concurrent queue and use the next frame as the current frame;

[0016] S340: If the current frame is not the first frame, obtain the relative block number and the sequence number within the block of the current frame according to the frame second and the frame sequence number within the second of the current frame, the frame second and the frame sequence number within the second of the reference frame, as well as the number of blocks in the circular buffer and the number of frames that can be accommodated within a block, and determine whether the relative block number is less than 0 or greater than cqs - 1;

[0017] S350: If the relative block number is less than 0 or greater than cqs - 1, return to step S310; otherwise, further determine whether the relative block number is equal to cqs - 1;

[0018] S360: If not, it means that the relative block number is greater than or equal to 0 and less than cqs - 1, then write the frame header and frame body of the current frame to the corresponding position in the circular buffer according to the relative block number and the sequence number within the block, and then return to step S310;

[0019] If so, write the frame header and frame body of the current frame to the corresponding positions in the circular buffer according to the relative block number and the in-block sequence number; then obtain illegal frames by scanning and reading blocks, and write the frame bodies of all frames in the read blocks to the second concurrent queue. When writing, if an illegal frame is encountered, replace the frame body of the illegal frame with the frame body of the previous legal frame and write it to the second concurrent queue; after writing is completed, use the next block of the read block as the new read block, and then return to step S310.

[0020] Further, obtain the relative block number and the in-block sequence number of the current frame according to the frame second and the in-second frame sequence number of the current frame, the frame second and the in-second frame sequence number of the reference frame, as well as the number of blocks in the circular buffer and the number of frames that can be accommodated in each block. Specifically, it includes:

[0021] Obtain the frame difference between the current frame and the reference frame according to the frame second and the in-second frame sequence number of the current frame, the frame second and the in-second frame sequence number of the reference frame, as well as the number of blocks in the circular buffer and the number of frames that can be accommodated in each block. Divide the frame difference by the number of frames that can be accommodated in each block and round down to obtain the relative block number of the current frame. Take the remainder of the frame difference divided by the number of frames that can be accommodated in each block to obtain the in-block sequence number of the current frame.

[0022] Further, step S350 further includes: when an illegal frame is encountered, write the frame second and the in-second frame sequence number of the illegal frame to the dropped frame file.

[0023] Further, cqs is 3.

[0024] Step S400 further includes:

[0025] S410: Create an input stream, a decoding stream, an output stream, an input event, and a decoding event;

[0026] S420: Read data of a first preset length from the second concurrent queue into the CPU input buffer;

[0027] S430: The CPU asynchronously calls the input stream to copy data of a second preset length from the CPU input buffer to the GPU input buffer through the input stream. After the copying is completed, an input event is triggered by the input stream, where the first preset length is an integer multiple of the second preset length;

[0028] S440: The CPU asynchronously calls the decoding stream to wait for the input event. After the input event is triggered, perform a decoding kernel function operation through the decoding stream to decode the data of the second preset length in the GPU input buffer, and store the decoded data in the GPU output buffer. After the decoding is completed, the decoding stream triggers a decoding event;

[0029] S450: The CPU asynchronously calls the output stream to wait for the decoding event. After the decoding event is triggered, copy the decoded data from the GPU output buffer to the CPU output buffer;

[0030] S460: Re - execute steps S440 - S450 until all the data in the CPU input buffer is decoded, and then return to step S420.

[0031] Step S440 further includes:

[0032] S441: Make the data in the GPU input buffer correspond one - to - one with the threads of the decoding kernel function;

[0033] S442: Each thread of the decoding kernel function decodes the corresponding data according to a preset quantization method and stores the decoded data in the GPU output buffer.

[0034] Furthermore, the preset quantization method is to map through a preset index array with the VDIF - formatted data or convert the binary VDIF - formatted data to decimal.

[0035] Furthermore, the input stream, the decoding stream, and the output stream are set for parallel processing.

[0036] Furthermore, the storage space of the second concurrent queue is larger than that of the first concurrent queue.

[0037] The method of using GPU to decode VDIF - formatted data in the present invention has the following technical effects:

[0038] First, fully consider the time continuity of the VDIF - formatted data of pulsars, including: 1) Ensure time continuity by replacing illegal frames with the previous legal frame; 2) Record the positions of illegal frames through a dropped - frame file; The traditional frame - filling method is to simulate by calculating values such as the average and standard deviation of historical data. In comparison, the method of the present invention is faster; Since frame dropping is a local situation and the pulsar Searching mode does not require a great deal of demand for fine structures, as long as time continuity is maximally ensured, the loss of a small part of the middle data will not affect the search; In addition, when the pulsar parameters are unknown, simulating frame dropping does not achieve a good effect, and because statistical values of adjacent data are required, it will take a lot of time, resulting in an increase in frame dropping. The method of the present invention can avoid these problems.

[0039] Second, the sizes of each data buffer are fixed, which is easy to implement the method, that is, the procurement of equipment; Moreover, it avoids problems caused by an unbounded buffer, such as the system crashing due to excessive dynamically expanded space of the buffer.

[0040] Third, the input and output buffers of the CPU and GPU are page - locked memory, with fast and stable read - write speeds, avoiding the problems of unified memory and preventing security problems caused by memory expansion.

[0041] Finally, the input stream, decoding stream, and output stream of the GPU are processed in parallel, and the frame reading, frame rearrangement, and frame decoding of the CPU are also processed in parallel, greatly reducing the processing time and improving the decoding speed. Description of the Drawings

[0042] Figure 1 It is a flowchart of a method for decoding VDIF format data using a GPU according to the present invention. Detailed Embodiments

[0043] The following will give and describe in detail the preferred embodiments of the present invention in conjunction with the accompanying drawings.

[0044] As Figure 1 shown, the present invention provides a method for decoding VDIF format data using a GPU, which cooperates with a CPU (Central Processing Unit) and a GPU (Graphics Processing Unit) to quickly decode VDIF format data. It mainly includes three parallel stages (frame reading, frame rearrangement, and frame decoding) to improve the processing speed. First, in the frame reading stage, VDIF format data sent by the acquisition device is received. In the frame rearrangement stage, the VDIF format data is arranged in chronological order. In the frame decoding stage, the VDIF format data is decoded through a GPU kernel function to obtain a three-dimensional array format (number of FFTs - channels - FFT frequency points) suitable for FFT batch processing. The method includes the following steps:

[0045] S100: Initialize the first concurrent queue, the second concurrent queue, the circular buffer, the CPU input buffer, the GPU input buffer, the CPU output buffer, and the GPU output buffer.

[0046] The first concurrent queue, the second concurrent queue, the circular buffer, the CPU input buffer, the GPU input buffer, the CPU output buffer, and the GPU output buffer are 7 fixed storage spaces on the computing device and are used as data buffers for the method of the present invention.

[0047] The first concurrent queue serves as the output space for the frame reading stage and is used to store the received VDIF - formatted data. The unit of the VDIF - formatted data is a frame, and each frame of VDIF - formatted data includes a frame header and a frame body. The second concurrent queue serves as the output space for the frame rearrangement stage and is used to store the frame body data of each ordered frame of VDIF - formatted data. The data unit sizes of both the first concurrent queue and the second concurrent queue can be set to 1 byte (char type). Since the frame reading speed is set to be faster than the frame rearrangement speed, generally, the maximum available space of the second concurrent queue needs to be larger than that of the first concurrent queue. The maximum available space of the first concurrent queue needs to be at least larger than the byte number of two frames, which can reduce the probability of concurrent resource contention in frame reading and writing and frame rearrangement reading, thereby improving the processing speed. The maximum available space of the second concurrent queue can be at least equal to twice the size of the CPU input buffer, which can reduce the probability of concurrent resource contention in frame rearrangement writing and frame decoding reading, thereby improving the processing speed. The concurrent queue supports data order and concurrency, and its implementation method is well - known in the art and will not be elaborated here.

[0048] The circular buffer is used to store the rearranged VDIF frames during the frame rearrangement stage. The circular buffer can be composed of cqs blocks. In some embodiments, cqs can be 3. To batch - process the data within 1 second by dividing it into N blocks (N is a positive integer), the number of data in each block can be set to nframes (nframes is also a positive integer, representing the number of frames in 1 / N seconds), and the data unit size is VHS + VBS (char array type), where VHS is the byte size of the frame header and VBS is the byte size of the frame body. When the observation mode is fixed, the bandwidth of the receiver is fixed. Therefore, the frame size and the number of frames per second sent by the acquisition device are generally fixed. The value of nframes depends on the bandwidth used in the pulsar observation mode, the number of VDIF bits, the byte size of the VDIF frame body, and the number of divisions N. If the processing speed is to be improved, by adjusting the value of N, the total size of the circular buffer is made not to exceed the sum of the CPU's level - 1, level - 2, and level - 3 caches. N satisfies relation (1), and nframes satisfies relation (2). In formula 1 represents rounding down a, that is, the largest integer less than a.

[0049]

[0050]

[0051] For example, in an exemplary embodiment, the bandwidth used in the pulsar observation mode is 8 GHz, the VDIF data bit nibit is 2, VHS is 32 bytes, VBS is 8000 bytes, and the total CPU cache is 64 M bytes. Then, according to Equation (1), N ≤ 62500 can be obtained, and N can take 62500. According to Equation (2), nframes = 256. nframes must be an integer. If it is not an integer, the value of N needs to be adjusted to make it an integer.

[0052] The CPU input buffer is used to store the data to be decoded (the sorted data output from the second concurrent queue) on the CPU during the frame decoding stage. The number of data is nbytes, and the data unit size is 1 byte (char type).

[0053] The GPU input buffer is used to store the data to be decoded (copied from the CPU input buffer) on the GPU during the frame decoding stage. The number of data is nbytes, and the data unit size is 1 byte (char type).

[0054] The GPU output buffer is used to store the decoded data on the GPU during the frame decoding stage (the data obtained after decoding the data in the GPU input buffer using the GPU). The number of data is 8 * nbytes / nibit, where nibit is the number of bits of the frame body data of the VDIF frame, and the data unit size is nobit bits (usually 32-bit floating-point type or 64-bit double-precision floating-point type).

[0055] The CPU output buffer is used to store the decoded data on the CPU during the frame decoding stage (copied from the GPU output buffer). The number of data is 8 * nbytes / nibit, and the data unit size is nobit bits (usually 32-bit floating-point type or 64-bit double-precision floating-point type).

[0056] S200: Frame reading stage: Read the VDIF format data from the output port of the acquisition device and write the VDIF format data into the first concurrent queue.

[0057] In some embodiments, the VDIF format data can be read from the output port of the acquisition device through the User Datagram Protocol (UDP). Specifically, the output port can be bound, and then UDP data packets (i.e., VDIF format data) can be listened for. If no data arrives, continue to wait. If data arrives, call the interface of the first concurrent queue (thread safety needs to be ensured) to write the data into the first concurrent queue.

[0058] S300: Frame Rearrangement Stage: Sort the VDIF format data in the first concurrent queue in chronological order, and load the sorted VDIF format data into the second concurrent queue.

[0059] The VDIF format data is generated in chronological order. However, when writing the read VDIF format data into the first concurrent queue, it may not be arranged in chronological order. For example, within 1 second, many frames of VDIF format data may be generated. The nth frame data within the 1st second is generated before the (n + 1)th frame data within the 1st second. But when writing into the first concurrent queue, the (n + 1)th frame data within the 1st second may be written into the first concurrent queue before the nth frame data within the 1st second. It cannot be directly used for FFT calculation. Therefore, it is necessary to re - sort it to make it arranged in the chronological order of VDFI data generation.

[0060] Step S300 further includes:

[0061] S310: Read the frame header of the current frame from the first concurrent queue, obtain the frame body length, frame second, and frame sequence number within the second of the current frame according to the frame header of the current frame; then read the frame body of the current frame from the first concurrent queue according to the frame body length of the current frame.

[0062] When reading VDIF data from the first concurrent queue, it is read in units of frames. Since the byte size of each frame body is fixed at 32 bytes, 32 - byte data can be read from the first concurrent queue first, which is the frame header of the current frame. The frame header records the frame body length, frame second, and frame sequence number within the second of the current frame. Among them, the frame second is the absolute time, indicating which second the current frame is generated in, and the frame sequence number within the second represents the frame sequence number of the current frame within that second, that is, which frame the current frame belongs to within that second. For example, if the frame second is 10000, it means the current frame is generated at the 10000th second from the reference time. If the frame sequence number within the second is 20, it means the current frame is the 20th frame within that second. According to the frame second, the frame sequence number within the second, and the number of frames generated within 1 second, the specific generation time of the current frame can be obtained. For example, if the number of frames generated within 1 second is 1000, then the generation time of the current frame with a frame second of 10000 and a frame sequence number within the second of 20 is (10000 - 1)+20 / 1000 seconds. Since the frame body length of the current frame has been obtained, read the data of the frame body length of the current frame from the first concurrent queue, which is the frame body of the current frame.

[0063] S320: Determine whether the current frame is the first frame.

[0064] The first data read from the first concurrent queue is the first frame. Generally, it is considered that the first frame data read after initializing the first concurrent queue, the second concurrent queue, the circular buffer, the CPU input buffer, the GPU input buffer, the CPU output buffer, and the GPU output buffer is the first frame, and the subsequent data does not belong to the first frame. In some embodiments, an initial value firstVDIF can be set. After initialization, firstVDIF is assigned the value 1. After reading the first frame data from the first concurrent queue, since firstVDIF is 1, it indicates that this data is the first frame data, and then firstVDIF is assigned the value 0. When reading data from the first concurrent queue subsequently, the value of firstVDIF is determined. If it is 0, it indicates that this data is not the first frame.

[0065] S330: If the current frame is the first frame, write the frame header and frame body of the current frame to the beginning of one of the blocks in the circular buffer, and use the block in the circular buffer that stores the first frame as the read block. Use the frame stored at the beginning of the read block as the reference frame, and then return to step S310 to read the next frame of data from the first concurrent queue and use the next frame as the current frame.

[0066] Since the current frame is the first frame, its frame header and frame body can be directly written into the circular buffer. The absolute time of the first frame is the benchmark for sorting and also the starting point of the data. Since the circular buffer is divided into cqs blocks, each block has its own number, such as 0, 1..cqs - 1 respectively. Each block can store nframes frame data, and the storage position of each frame data also has its own in-block sequence number, such as 1, 2...nframes respectively, indicating the storage positions of the 1st, 2nd...nframes in the block. When reading or writing data to the circular buffer, it will start from the position with in-block sequence number 1 first, and then read sequentially. The address where the data is read or written for the first time is the beginning of the block. For example, when the read pointer points to block 0 of the buffer, it will start reading data from the beginning of block 0 (i.e., number 1), and then read sequentially to nframes to complete the reading of the data in this block.

[0067] When the current frame is the first frame, the block that stores the first frame is the read block. When data needs to be retrieved from the circular buffer, start reading from the read block. The frame stored at the beginning of the read block can be used as the reference frame. The reference frame is used to determine whether the data in the circular buffer needs to be written to the second concurrent queue and to locate the address (i.e., position) in the circular buffer where the subsequent frames to be read are stored, so as to ensure that the VDIF format data written to the second concurrent queue is arranged in chronological order.

[0068] S340: If the current frame is not the first frame, obtain the relative block number and the in-block sequence number of the current frame based on the frame second and the in-second frame sequence number of the current frame, the frame second and the in-second frame sequence number of the reference frame, as well as the number of blocks in the circular buffer and the number of frames that can be accommodated in each block, and determine whether the relative block number is less than 0 or greater than cqs - 1.

[0069] S350: If the relative block number is less than 0 or greater than cqs - 1, go back to step S310; otherwise, further determine whether the relative block number is equal to cqs - 1.

[0070] Since the storage space of the circular buffer is limited, if the frame difference between the current frame and the reference frame exceeds its storage space, then the current frame cannot be stored in the circular buffer, indicating data abnormality. The current frame cannot be sorted and needs to be discarded, and then go back to step S310 to read the next frame from the first concurrent queue. Specifically, the frame difference between the current frame and the reference frame can be calculated through the frame second and the in-second frame sequence number of the current frame and the frame second and the in-second frame sequence number of the reference frame, and then divide the frame difference by the number of frames nframes that can be accommodated in each block. The remainder is used as the in-block sequence number, and the result is rounded down to obtain the relative block number of the current frame. If the relative block number is less than 0, it means that the current frame was generated before the reference frame. If the relative block number is greater than cqs - 1, it means that the current frame was generated after the reference frame and the number of frames in between exceeds the storage space of the circular buffer and cannot be stored in the circular buffer.

[0071] S360: If not, it means that the relative block number is greater than or equal to 0 and less than cqs - 1, then write the frame header and frame body of the current frame to the corresponding position in the circular buffer according to the relative block number and the in-block sequence number, and then go back to step S310; if so (i.e., the relative block number is equal to cqs - 1), then write the frame header and frame body of the current frame to the corresponding position in the circular buffer according to the relative block number and the in-block sequence number; then obtain the illegal frames by scanning and reading the block, and write the frame bodies of all the frames in the read block to the second concurrent queue. When writing, if an illegal frame is encountered, replace the frame body of the illegal frame with the frame body of the previous legal frame and write it to the second concurrent queue; after writing is completed, use the next block of the read block as the new read block, and then go back to step S310.

[0072] If the relative block number is greater than or equal to 0 and less than cqs - 1, it indicates that the frame difference between the current frame and the reference frame does not exceed its storage space and can be stored in the circular buffer. At this time, the relative block number represents the number of blocks between the block storing the current frame and the block storing the reference frame. The corresponding position of the current frame stored in the circular buffer can be obtained based on the relative block number and the in - block sequence number. For example, if the reference frame is stored in block 0 and the relative block number is 1, then there is a difference of 1 block between the block storing the current frame and the block storing the reference frame. Therefore, the current frame needs to be stored in block 1. If the reference frame is stored in block cqs - 1 and the relative block number is 2, then there is a difference of 2 blocks between the block storing the current frame and the block storing the reference frame. Since the blocks form a circle, the 2 blocks after block cqs - 1 are block 0 and block 1 in sequence. Therefore, the current frame needs to be stored in block 1. And so on. The in - block sequence number is the sequence number of the current frame stored within the block. If the in - block sequence number is 2, then the current frame needs to be stored at the position of the second frame within the block.

[0073] If the relative block number is equal to cqs - 1, it indicates that there are already cqs - 1 blocks storing VDIF data. The current frame can be stored in the circular buffer. However, at the same time, the data in the read block needs to be written into the second concurrent queue to release the read block, facilitating subsequent data storage. Since when reading data from the first concurrent queue, some frames that do not meet the requirements may be skipped (see step 330), the data in the read block may not be complete, that is, there is no data written in some positions. For example, assume N = 2 and 30 frames of data are generated within 1 second, then nframes = 15, and each block can store 15 frames of data. The reference frame is stored at the beginning of the read block. Usually, it is expected to store the next frame of the reference frame (the next bit of the frame sequence number within the second) at the position of the 2nd frame in the read block to arrange them in chronological order. However, the next frame of the reference frame may appear before the reference frame and has thus been discarded. Then, there will be no data stored at the position of the 2nd frame in the read block. At this time, the next frame of the reference frame is a missing frame or an illegal frame, and the other frames stored in the read block are legal frames. By scanning the read block, all illegal frames can be obtained, and then the information of the illegal frames (including the frame second and the frame sequence number within the second) is recorded in the dropped frame file. Since the data in the second concurrent queue is used for subsequent decoding, it is necessary to ensure that its data is arranged in chronological order and is continuous without missing. When writing the data in the read block into the second concurrent queue, since there is no data at the storage position corresponding to the illegal frame, if directly written according to the data in the read block, it will cause data missing in the second concurrent queue, resulting in discontinuous time. Therefore, when writing the data in the read block into the second concurrent queue, if an illegal frame is encountered, the previous legal frame of the illegal frame can be used to replace the illegal frame and written into the second concurrent queue. For example, if the positions of the 2nd frame and the 10th frame in the read block are empty and the rest of the positions store data, then when writing the data in the read block into the second concurrent queue, start from the reference frame (i.e., the frame stored at the beginning of the read block), write the frame body of the reference frame into the second concurrent queue, then use the frame body of the reference frame to replace the frame body of the illegal frame corresponding to the position of the 2nd frame in the read block and write it into the second concurrent queue. Then, write the frame bodies of the frames stored at the positions of the 3rd frame, the 4th frame... the 9th frame into the second concurrent queue in turn. Then, use the frame body of the frame stored at the position of the 9th frame to replace the frame body of the illegal frame corresponding to the position of the 10th frame and write it into the second concurrent queue, and so on. Finally, the arrangement order in the second concurrent queue is the frame body of the reference frame, the frame body of the reference frame, the frame body of the 3rd frame... the frame body of the 9th frame, the frame body of the 9th frame, the frame body of the 11th frame... the frame body of the Nth frame. In this way, the data in the second concurrent queue is strictly arranged in chronological order and there is no missing data.

[0074] The recording format of the dropped frame file can be set as follows: one dropped frame record per line, with the frame seconds and frame number in the line separated by commas. The writing method can be achieved by using the aio method of the operating system, which is a well-known method and will not be elaborated here. Through the dropped frame file, the time information of the dropped frames can be determined, facilitating the identification of which frames at which time points are dropped when processing the decoded data (such as FFT).

[0075] After the data in the read block is completely read, take the next block of the read block as the new read block, and then return to step S310 until the observation ends.

[0076] In some embodiments, if the block numbers of the circular buffer are sequentially 0, 1..cqs - 1, steps S310 - S350 can be implemented by the following method (the "=" sign in the method represents assignment):

[0077] S310a: Read 32 bytes of data from the first concurrent queue, which is the frame header of the current frame. Obtain the frame body length, frame seconds fsec, and frame sequence number within the second fnum of the current frame based on the frame header, and then read data with the size of the frame body length of the current frame from the first concurrent queue, which is the frame body of the first concurrent queue.

[0078] S320a: If the current frame is the first frame of this observation, then set fsec0 = fsec, fnum0 = fnum, blk = 0, and write the frame header and frame body of the current frame to the beginning of the 0th block of the circular buffer. blk0 is used to mark the read block of the circular buffer in the current state and count the number of blocks rearranged in this observation.

[0079] S330a: If the current frame is not the first frame of this observation, then calculate the relative block number iblk and the sequence number within the block iframe of the current frame through the following formulas (3) and (4), where MOD represents the modulo operator;

[0080]

[0081] iframe = frame difference MOD number of frames that can be accommodated within the block = [(fsec - fsec0) × nframes × N + fnum - fnum0] MOD [nframes] (4)

[0082] where fsec is the frame seconds of the current frame, fsec0 is the frame seconds of the reference frame, fnum is the frame sequence number within the second of the current frame, and fnum0 is the frame sequence number within the second of the reference frame

[0083] If the relative block number iblk is less than 0 or greater than cqs - 1, it indicates that the data is abnormal. Skip this frame and return to step S310a again.

[0084] S340a: If the relative block number iblk is greater than or equal to 0 and less than cqs - 1, calculate the index position blkn of the block storing the current frame as blkn = [blk0 + iblk] MOD cqs. Then write the frame header and frame body of the current frame to the position of the iframes frame in the blkn-th block of the circular buffer, and then return to step S310a;

[0085] S350a: If the relative block number iblk is equal to cqs - 1, calculate the index position blkn of the block storing the current frame as blkn = [blk0 + iblk] MOD cqs. Then write the frame header and frame body of the current frame to the position of the iframes frame in the blkn-th block of the circular buffer. Then calculate the actual index number of the read block as (blk0 + iblk) MOD cqs, scan and read the block according to the actual index number to obtain an illegal frame, and record the illegal frame information in the dropped frame file. Then write the frame bodies of all frames in the read block to the second concurrent queue. When writing, if an illegal frame is encountered, replace the frame body of the illegal frame with the frame body of the previous legal frame of the illegal frame and write it to the second concurrent queue; after writing, set blk0 = blk0 + 1, fsec0 = fsec0 + (fnum0 + nframes) / (nframes * N), fnum0 = (fnum0 + nframes) mod (nframes * N), and then return to step S310a. Repeat this process until the end of this observation.

[0086] In an exemplary embodiment, the total number of frames in 1 second is 30, N = 2, cqs = 3, the number of frames that each block of the circular buffer can accommodate is nframes = 30 / N = 15, the block numbers (i.e., actual index numbers) of the circular buffer are 0, 1, and 2 respectively, the frame sequence numbers within each second are 0 = 29, the frame second fsec1 of the first frame (reference frame) is 0, the frame sequence number fnum1 is 0, then fsec0 = fsec1 = 0, fnum0 = fnum1 = 0, blk0 = 0. Five frames arrive in sequence, which are frame second 0 and frame sequence number within the second 9, frame second 1 and frame sequence number within the second 25, frame second 1 and frame sequence number within the second 11, frame second 1 and frame sequence number within the second 21, frame second 1 and frame sequence number within the second 22. The processing conditions are as follows:

[0087] 1) The frame second fsec of the current frame is 0, the frame sequence number within the second fnum is 9, the frame difference between the current frame and the reference frame is [(0 - 0) × 15 × 2 + 9 - 0] = 9, so iframe = 9 mod 15 = 9, Since 0 is not equal to cqs - 1, it is not released. The actual index block number is (0 + 0) mod 3 = 0, and it is directly stored in the 0-th block, the position of the 9-th frame; then continue to read the next frame: frame second 1 and frame sequence number within the second 25.

[0088] 2) If the frame second fsec of the current frame is 1 and the frame sequence number fnum within the second is 25, then the frame difference from the reference frame is [(1 - 0)×15×2 + 25 - 0] = 55. So, iframe = 55 mod 10 = 5. Since 3 is greater than cqs - 1, exceeding the current storage range, it is abnormal data. Skip this frame and continue to read the next frame: frame second 1 and frame sequence number 11 within the second.

[0089] 3) If the frame second fsec of the current frame is 1 and the frame sequence number fnum within the second is 11, then the frame difference from the reference frame is [(1 - 0)×15×2 + 11 - 0] = 41. So, iframe = 41 mod 15 = 11. Since 2 is equal to cqs - 1, indicating that two blocks are full, it is necessary to release the read block. The actual block number of the read block is blk0 mod cqs = 0 mod 3 = 0. The data in the read block is written into the second concurrent queue in chronological order. Then, clear block 0, and make blk0 = 0 + 1 = 1, fesc0 = 0 + (0 + 15) / (15*2) = 0, fnum0 = 0 + 15 = 15. Then continue to read the next frame: frame second 1 and frame sequence number 21 within the second.

[0090] 4) If the frame second fsec of the current frame is 1 and the frame sequence number fnum within the second is 21, the frame difference between this frame and the reference frame is [(1 - 0)×15×2 + 21 - 15] = 36. So, iframe = 36 mod 15 = 6. Since 2 is equal to cqs - 1, indicating that two blocks are full, it is necessary to release the read block. The actual block number of the read block is blk0 mod cqs = 1 mod 3 = 1. The data in the read block is written into the second concurrent queue in chronological order. Then, clear block 0, and make blk0 = 1 + 1 = 2, fesc0 = 0 + (15 + 15) / (15*2) = 1, fnum0 = (15 + 15) mod (15*2) = 0. Then continue to read the next frame: frame second 1 and frame sequence number 22 within the second.

[0091] 5) If the frame second fsec of the current frame is 1 and the frame sequence number fnum is 22, then the frame difference between this frame and the first frame is [(1 - 1)×15×2 + 22 - 0] = 22. So, iframe = 22 mod 15 = 7. Since 1 is not equal to cqs - 1, there is no need to release. The actual index block number is (2 + 1) mod 3 = 0. Directly store the current frame in block 0, at the position of the 7th frame. Then continue to read the next frame.

[0092] Step S400: Frame decoding stage: The data in the second concurrent queue is sequentially written into the CPU input buffer, and then the data in the CPU input buffer is copied to the GPU input buffer. The GPU decodes the data in the GPU input buffer and stores the decoded data in the GPU output buffer, and then the decoded data in the GPU output buffer is copied to the CPU output buffer.

[0093] The frame decoding stage mainly decodes the continuous VDIF format data in the second concurrent queue. Since the GPU cannot directly read data from the second concurrent queue, the data in the second concurrent queue is first sequentially written into the CPU input buffer and then copied to the GPU input buffer to utilize the GPU to decode the data in the GPU input buffer. The decoded data is stored in the GPU output buffer, which can be copied to the CPU output buffer for subsequent use.

[0094] Step S400 further includes:

[0095] S410: Create an input stream, a decoding stream, an output stream, an input event, and a decoding event.

[0096] A stream, also known as a CUDA (Compute Unified Device Architecture) stream, is a general-purpose parallel computing architecture introduced by NVIDIA. This architecture enables the GPU to solve complex computing problems. A stream is a work queue on the GPU, and the work in the queue will be executed in a specific order. This work can include: calls to kernel functions, cudaMemcpy series transfers, and operations on CUDA events. The order in which each stream is added to the queue will determine their execution order. Each CUDA stream can be regarded as a single task, so multiple streams can be started to execute multiple tasks in parallel. In many cases, the work in multiple streams may be executed simultaneously.

[0097] In some embodiments, the input stream is used to execute a kernel function for memory copy operations from the CPU to the GPU to copy the data in the CPU input buffer to the GPU input buffer; the decoding stream is used to execute kernel functions related to decoding to decode the data in the GPU input buffer; the output stream is used to execute a kernel function for memory copy operations from the GPU to the CPU to copy the decoded data in the GPU output buffer to the CPU output buffer; the input event is an event triggered by the input stream after the input is completed and waited for by the decoding stream; the decoding event is an event triggered by the decoding stream after decoding is completed and waited for by the output stream.

[0098] S420: Read data of a first preset length from the second concurrent queue into the CPU input buffer.

[0099] The byte size of the data that can be stored in the CPU input buffer is nbytes, so the first preset length is nbytes bytes.

[0100] S430: The CPU asynchronously calls the input stream to copy data of the second preset length from the CPU input buffer to the GPU input buffer through the input stream. After the copying is completed, an input event is triggered by the input stream, where the first preset length is an integer multiple of the second preset length.

[0101] To improve the decoding speed, the data in the CPU input buffer can be divided into m parts to achieve parallel processing, that is, the second preset length can be nbytes / m bytes. In some embodiments, nbytes and m need to meet the following conditions: 1) nbytes can be divided evenly by m; 2) nbytes / m can be divided evenly by nbin, where nbin is the number of FFT points, that is, the number of data points used when performing FFT on the decoded data later. Usually, nbin = 65536 (or other powers of 2), and nfft = nbytes / m / nbin represents the number of times the data of the current row in parallel processing can be used to execute FFT.

[0102] S440: The CPU asynchronously calls the decoding stream to wait for the input event. After the input event is triggered, the decoding kernel function operation is executed through the decoding stream to decode the data of the second preset length in the GPU input buffer, and the decoded data is stored in the GPU output buffer. After the decoding is completed, the decoding stream triggers a decoding event.

[0103] Asynchronous call means that after the CPU calls the input stream, instead of waiting for the input stream to complete the copying, it directly calls the decoding stream. Therefore, during the process of the input stream copying data, the decoding stream has been started and is waiting for the input event. After the input event is triggered, the decoding stream will immediately start decoding.

[0104] The decoding stream decodes through a preset decoding kernel function. In some embodiments, the thread block of the decoding kernel function can be the three-dimensional thread block and three-dimensional grid block structure provided by CUDA. The thread block format is (blocksize.x, blocksize.y, blocksize.z), and the specific values are determined by the specific graphics card model used (ensuring not exceeding the maximum supported amount). Usually, the values can be (1024, 32, 32). The grid block format of the GPU is (gridsize.x, gridsize.y, gridsize.z), and its x can be set to (rounded up), y can be z can be

[0105] The decoding method of the decoding kernel function is as follows:

[0106] S441: Correspond one by one the data in the GPU input buffer with the thread blocks of the decoding kernel function.

[0107] Specifically, the channel number ichan to which the data calculated by the current thread belongs, the sequence number ibin within the FFT to which the data belongs, and the FFT number ifft to which the data point belongs can be calculated according to the thread block format and grid block format of the decoding kernel function. The calculation formulas are as follows:

[0108] ibin = blockIdx.x * blockDim.x + threadIdx.x (5)

[0109] ifft = blockIdx.y * blockDim.y + threadIdx.y (6)

[0110] ichan = blockIdx.z * blockDim.z + threadIdx.z (7)

[0111] S442: Each thread of the decoding kernel function decodes the data corresponding to each thread according to a preset quantization method, and stores the decoded data in the GPU output buffer.

[0112] Specifically, the index position idx1 of the decoded data and the index position idx2 where the data to be decoded is stored can be obtained according to ichan, ibin, and ifft, and then the intermediate value ivalue is calculated through idx1 and idx2. The calculation formulas for idx1, idx2, and ivalue are as follows:

[0113] idx1 = ibin + nbin * ichan + nchan * nbin * ifft (8)

[0114] idx2 = ichan + nchan * ibin + nchan * nbin * ifft (9)

[0115] ibyte = idx2 >> nibit (10)

[0116] ibits = nibit * (idx2 & ((1 << nibit) - 1))) (11)

[0117] ivalue = (nbytesBuf[ibyte] >> ibits) & ((1 << nibit) - 1)) (12) ]>

[0118] Among them, >>, <<, and & are all bit operators. >> is right shift, << is left shift, and & is bitwise AND.

[0119] The ivlaue can be decoded by a preset quantization method, and then the decoded data is mapped to the GPU output buffer through idx1. The quantization method refers to mapping the established index array with the VDIF format data to decode the VDIF format data. For example, the VDIF format data is usually 2 bits, which is binary. The quantized values are 0, 1, 2, or 3, and its index array mask can be [-3.3359, -1, 1, 3.3359]. Then, during decoding, the 0 in the VDIF format data can be decoded as -3.3359, 1 as -1, 2 as 1, and 3 as 3.3359 to obtain the decoded data. In some embodiments, the quantization method can also directly quantize the binary VDIF format data into a decimal value as the decoded data. For example, if the VDIF format data is 01, the decoded data is 1, and if it is 11, the decoded data is 3.

[0120] When decoding using the index array, the decoded data can be obtained by the following formula and mapped to the GPU output buffer:

[0121] Buffer go [idx1] = mask[ivalue] (13)

[0122] Buffer go [idx1] refers to the index position of idx1 that maps the value on the right side of the formula to the GPU output buffer.

[0123] When directly quantizing the binary VDIF format data into decimal for decoding, the decoded data can be obtained by the following formula and mapped to the GPU output buffer:

[0124] Buffer go [idx1] = ivalue - (1 << (nibit - 1)) (14)

[0125] S450: The CPU asynchronously calls the output stream to wait for the decoding event, and after the decoding event is triggered, copies the decoded data from the GPU output buffer to the CPU output buffer.

[0126] After the CPU calls the decoding stream, it immediately calls the output stream, and the work of the output stream can be carried out simultaneously with the decoding stream and the input stream.

[0127] S460: Re-execute steps S440 - S450 until all the data in the CPU input buffer is decoded, and then return to step S420.

[0128] During the frame decoding stage, the data in the CPU input buffer is copied in integer parts to the GPU input buffer. Since the input stream, decoding stream, and output stream can be processed in parallel, input, decoding, and output can be performed simultaneously, reducing the computation time and improving the computation speed.

[0129] In some embodiments, the frame reading stage can be implemented by a frame reading module, the frame rearrangement stage can be implemented by a frame rearrangement module, and the frame decoding stage can be implemented by a frame decoding module. In some embodiments, when decoding VDIF format data using the method of the present invention, the frame decoding module, the frame rearrangement module, and the frame reading module can be started sequentially. When the frame decoding module cannot obtain the data in the CPU input buffer, it will wait and send a signal to the frame reading module, causing the frame rearrangement module to store the arranged data in the CPU input buffer. When the frame rearrangement module cannot obtain data from the first concurrent queue, it will wait and send a signal to the frame reading module, causing the frame reading module to read data from the acquisition device.

[0130] Frame reading, frame rearrangement, and frame decoding can also be implemented in parallel. That is, during frame decoding, the frame reading module is simultaneously reading new data, and the frame rearrangement module is simultaneously arranging the data, thus greatly overlapping the reading, rearrangement, and decoding times and improving the decoding efficiency.

[0131] The method of decoding VDIF format data using a GPU according to the present invention has the following technical effects:

[0132] First, the time continuity of the VDIF format data of pulsars is fully considered, including: 1) ensuring time continuity by replacing illegal frames with the previous legal frame; 2) recording the positions of illegal frames through dropped frame files; the traditional frame filling method is to simulate by calculating values such as the average and standard deviation of historical data. In comparison, the method of the present invention is faster; since frame dropping is a local situation and the pulsar Searching mode does not require a high demand for fine structure, as long as time continuity is ensured to the greatest extent, the loss of a small part of data in the middle will not affect the search; in addition, when the pulsar parameters are unknown, simulating frame dropping does not achieve a good effect, and because statistical values of adjacent data are required, it will take a lot of time and lead to an increase in frame dropping. The method of the present invention can avoid these problems.

[0133] Second, the sizes of the respective data buffers are fixed, which is easy to implement the method, that is, the procurement of equipment; moreover, problems caused by buffer sizes without upper limits are avoided, such as system crashes due to excessive dynamically expanded space of the buffer.

[0134] Third, the input and output buffers of the CPU and GPU are page-locked memory, with fast and stable read and write speeds, avoiding the problems of unified memory and preventing security problems caused by memory expansion.

[0135] Finally, the input stream, decoding stream, and output stream of the GPU are processed in parallel, and the frame reading, frame rearrangement, and frame decoding of the CPU are also processed in parallel, greatly reducing the processing time and improving the decoding speed.

[0136] The above are only the preferred embodiments of the present invention and are not intended to limit the scope of the present invention. Various changes can be made to the above embodiments of the present invention. That is, all simple, equivalent changes and modifications made according to the claims and the content of the specification of the present invention application fall within the scope of the claims of the present invention patent. The content not described in detail in the present invention is all conventional technical content.

Claims

1. A method for decoding VDIF format data using a GPU, characterized in that, Including the steps: S100: Initialize the first concurrent queue, the second concurrent queue, the circular buffer, the CPU input buffer, the GPU input buffer, the CPU output buffer, and the GPU output buffer; S200: Read the VDIF format data from the output port of the acquisition device and write the VDIF format data into the first concurrent queue; S300: Sort the VDIF format data in the first concurrent queue in chronological order and load the sorted VDIF format data into the second concurrent queue; S400: Write the data in the second concurrent queue into the CPU input buffer in sequence, then copy the data in the CPU input buffer to the GPU input buffer. The GPU decodes the data in the GPU input buffer and stores the decoded data in the GPU output buffer, and then copies the decoded data in the GPU output buffer to the CPU output buffer; The circular buffer is divided into cqs blocks. Step S300 further includes: S310: Read the frame header of the current frame from the first concurrent queue, obtain the frame body length, frame second, and frame sequence number within the second of the current frame according to the frame header of the current frame; then read the frame body of the current frame from the first concurrent queue according to the frame body length of the current frame; S320: Determine whether the current frame is the first frame; S330: If the current frame is the first frame, write the frame header and frame body of the current frame to the beginning of one of the blocks in the circular buffer, and use the block in the circular buffer storing the first frame as the read block. Use the frame stored at the beginning of the read block as the reference frame, then return to step S310 to read the next frame of data from the first concurrent queue and use the next frame as the current frame; S340: If the current frame is not the first frame, obtain the relative block number and the sequence number within the block of the current frame according to the frame second and the frame sequence number within the second of the current frame, the frame second and the frame sequence number within the second of the reference frame, as well as the number of blocks in the circular buffer and the number of frames that can be accommodated within the block, and determine whether the relative block number is less than 0 or greater than cqs - 1; S350: If the relative block number is less than 0 or greater than cqs - 1, return to step S310; otherwise, further determine whether the relative block number is equal to cqs - 1; S360: If not, it means that the relative block number is greater than or equal to 0 and less than cqs - 1. Write the frame header and frame body of the current frame to the corresponding position in the circular buffer according to the relative block number and the sequence number within the block, and then return to step S310; If so, write the frame header and frame body of the current frame to the corresponding position in the circular buffer according to the relative block number and the sequence number within the block; then obtain the illegal frames by scanning the read block and write the frame bodies of all the frames in the read block into the second concurrent queue. When writing, if an illegal frame is encountered, replace the frame body of the illegal frame with the frame body of the previous legal frame and write it into the second concurrent queue; after writing is completed, use the next block of the read block as the new read block, and then return to step S310.

2. The method for decoding VDIF format data using a GPU according to claim 1, characterized in that, Obtain the relative block number and in-block sequence number of the current frame based on the frame second and in-second frame sequence number of the current frame, the frame second and in-second frame sequence number of the reference frame, as well as the number of blocks in the circular buffer and the number of frames that can be accommodated in each block, specifically including: Obtain the frame difference between the current frame and the reference frame based on the frame second and in-second frame sequence number of the current frame, the frame second and in-second frame sequence number of the reference frame, as well as the number of blocks in the circular buffer and the number of frames that can be accommodated in each block. Divide the frame difference by the number of frames that can be accommodated in each block and round down to obtain the relative block number of the current frame. Take the remainder of the frame difference with respect to the number of frames that can be accommodated in each block to obtain the in-block sequence number of the current frame.

3. The method for decoding VDIF format data using a GPU according to claim 1, characterized in that, Step S350 further includes: when an illegal frame is encountered, write the frame second and in-second frame sequence number of the illegal frame into the dropped frame file.

4. The method for decoding VDIF format data using a GPU according to claim 1, wherein cqs is 3.

5. The method for decoding VDIF format data using a GPU according to claim 1, characterized in that Step S400 further includes: S410: Create an input stream, a decoding stream, an output stream, an input event, and a decoding event; S420: Read data of a first preset length from the second concurrent queue into the CPU input buffer; S430: The CPU asynchronously calls the input stream to copy data of a second preset length from the CPU input buffer to the GPU input buffer through the input stream. After the copy is completed, the input stream triggers an input event, where the first preset length is an integer multiple of the second preset length; S440: The CPU asynchronously calls the decoding stream to wait for the input event. After the input event is triggered, perform decoding kernel function operations through the decoding stream to decode the data of the second preset length in the GPU input buffer and store the decoded data in the GPU output buffer. After the decoding is completed, the decoding stream triggers a decoding event; S450: The CPU asynchronously calls the output stream to wait for the decoding event. After the decoding event is triggered, copy the decoded data from the GPU output buffer to the CPU output buffer; S460: Re-execute steps S440 - S450 until all the data in the CPU input buffer is decoded, and then return to step S420.

6. The method for decoding VDIF format data using a GPU according to claim 5, characterized in that, Step S440 further includes: S441: Make the data in the GPU input buffer correspond one-to-one with the threads of the decoding kernel function; S442: Each thread of the decoding kernel function decodes the data corresponding to each thread according to a preset quantization method and stores the decoded data in the GPU output buffer.

7. The method for decoding VDIF format data using a GPU according to claim 6, characterized in that, The preset quantization method is to map through a preset index array with the VDIF format data or convert the binary VDIF format data to decimal.

8. The method for decoding VDIF format data using a GPU according to claim 5, characterized in that, The input stream, the decoding stream, and the output stream are set for parallel processing.

9. The method for decoding VDIF format data using a GPU according to claim 1, characterized in that, The storage space of the second concurrent queue is larger than that of the first concurrent queue.

Images