High-speed parallel Viterbi decoding method based on CUDA (Compute Unified Device Architecture)

Abstract

The invention discloses a high-speed parallel Viterbi decoder based on a CUDA (Compute Unified Device Architecture) and a decoding method. The decoder comprises an input cache module, a parallel data de-puncturing module, a parallel measurement calculation module, a backtracking path processing module and an output control module, and multi-level parallel processing is realized by utilizing a CUDA (Compute Unified Device Architecture) of a GPU (Graphics Processing Unit). The method comprises the following steps: segmenting a receiving sequence and then distributing the segmented receiving sequence to GPU thread blocks; rapidly reading a state transition matrix and an output matrix by adopting a pre-calculation table look-up strategy; adopting a parallel prefix scanning algorithm to accelerate state measurement calculation; and reducing the global memory access delay by using the shared memory. According to the invention, the performance bottleneck of the traditional serial Viterbi algorithm is broken through, on the premise of keeping the decoding precision, the throughput rate is obviously improved, the real-time processing of multi-code-rate and multi-code-stream convolutional codes is supported, and a remarkable acceleration effect can be obtained compared with CPU (Central Processing Unit) implementation. The method is especially suitable for high-throughput ACM scenes, such as 5G communication and satellite navigation, requiring real-time processing of multiple code streams and long constraint length convolutional codes.

Description

Technical Field

[0001] This invention belongs to the field of communication technology, and specifically discloses a high-speed parallel Viterbi decoder based on CUDA (Compute Unified Device Architecture), which is particularly suitable for adaptive coding and modulation (ACM) application scenarios such as 5G communication systems and satellite communication systems that have high requirements for decoding speed and throughput. Background Technology

[0002] The Viterbi algorithm, proposed by Andrew Viterbi in 1967, is a convolutional code decoding algorithm based on the maximum likelihood criterion. This algorithm represents the state transition process using a trellis diagram and employs dynamic programming to find the optimal path, exhibiting excellent error correction performance. Viterbi decoders are widely used in wireless communication, satellite communication, and deep space communication. The implementation of traditional Viterbi decoders faces two major challenges: first, the algorithm complexity increases exponentially with the constraint length; second, the inherent serial nature of the algorithm makes parallelization difficult. Although some research has attempted to accelerate Viterbi decoding using FPGAs or ASICs, these solutions lack flexibility and have long development cycles. On the other hand, modern GPUs possess thousands of parallel computing cores, providing an ideal hardware platform for the parallelization of Viterbi decoding.

[0003] In 5G and future communication systems, higher demands will be placed on the throughput and latency of channel encoding and decoding. Therefore, developing a high-efficiency, multi-stream parallel Viterbi decoder based on GPUs has significant practical application value. Summary of the Invention

[0004] This invention proposes a high-speed parallel Viterbi decoding method and apparatus for convolutional codes based on the CUDA architecture. Employing a two-level parallel architecture and multi-level storage optimization, it significantly improves the real-time performance and throughput of multi-rate convolutional code decoding. This method allows for flexible configuration of thread blocks and memory allocation for data streams with different code rates and segment lengths, fully utilizing the multi-core computing power of GPUs to achieve efficient decoding on a large-scale parallel hardware platform.

[0005] In the first-level parallel architecture, each data block is mapped to an independent thread block, enabling completely independent parallel decoding of different data blocks. In the second-level parallel architecture, each thread block is further allocated a corresponding thread for all states of the convolutional code, and parallel path metric updates and backtracking processes are performed for each surviving path, thereby achieving parallel computation at both the data block level and the state level, accelerating the decoding process.

[0006] Furthermore, the multi-level memory optimization strategy is as follows: frequently accessed but read-only structures such as state transition matrices, output matrices, and punch matrices are uniformly stored in GPU constant memory to fully utilize their low latency; data that needs to continuously participate in computation, such as mesh graphs and path metrics, are stored in shared memory to achieve high-speed access within thread blocks; and large data such as input sequences to be decoded and output sequences to be decoded are stored in GPU global memory to ensure overall data loading capacity and throughput efficiency.

[0007] The method first performs pre-calculation and initialization of the convolutional coding parameters, the state transition matrix `source_state`, and the output matrix `outputs` on the CPU host side. For (n, k, K) type convolutional codes, the encoder contains v = K−1 shift registers, with a total of S = 2 v There are several states. The state transition table and output table are constructed using the following lookup rules: When the current state s∈[0,S−1] and the input bit x∈{0,1}, construct the state transition matrix lookup index idx=s or s+S, obtain the predecessor state s_source from source_state[idx], and construct the output matrix lookup index idx=(s>>(K-2)<<(K−1))|s 0,1 The corresponding output codeword is obtained from outputs[idx], where s 0,1 Let s be the predecessor state, s be the current state, and S be the total number of states. This pre-computed lookup table strategy avoids the GPU kernel repeatedly performing polynomial convolution operations at runtime, significantly reducing computational load.

[0008] The host then copies the pre-calculated state transition table, output table, and index information from FrameParams to the GPU's global memory, while storing the punch matrix in constant memory. This completes the calculation and initialization operations.

[0009] After completing pre-calculation and initialization, the CPU host parses the received bitstream information, recording the starting position and bitrate parameters of each data block in the overall data stream and writing them into the CodeRate, DataBlocks, and DataBlocksLen fields of the FrameParams structure. Then, the parallel de-puncturing GPU kernel function is started, performing parallel de-puncturing operations on each data block on the GPU device side.

[0010] During parallel de-puncturing, the receiver performs hard-information de-puncturing on each received data unit according to the normalized puncturing matrix corresponding to the current coding rate, in order to recover the complete bit sequence structure of the original master code at rate 1 / n. Specifically, for bit positions deleted during the puncturing process, the missing bits are reinserted into the specified positions in the bit stream using an index mapping method, and filled with preset hard padding values ​​to ensure that the data frame length remains consistent with the master code structure. Each data block is assigned a separate thread for independent de-puncturing processing.

[0011] After the unpunching operation is completed on the GPU device, the starting position and length of the sub-data blocks of the data block are written into the SubDataBlocks and SubDataBlocksLen fields.

[0012] In Viterbi decoding, the optimal decoding path is selected by calculating the Hamming distance between the input bits and the current state code. Specifically, at each time step, the decoder compares the input bits with the current state code and calculates the Hamming distance between them. This distance measures the degree of mismatch between the encoded output and the input bits. Through this calculation, the merits of each path at the current time step can be determined; the smaller the Hamming distance, the higher the matching degree of the path, and the more likely it is to be the correct path.

[0013] Based on this, the Viterbi decoder performs path selection and updates according to the Hamming distance calculated at each time step. The path selection at each time step is based on the matching degree between the current input bits and the encoded output of each possible state. The decoder selects the best-matching path as the surviving path at each time step and updates its path metric. The number of surviving paths is usually equal to the total number of encoder states, ensuring that at each time step, multiple paths corresponding to the current state are processed in parallel.

[0014] A key feature of this invention is that the path computation and selection processes in Viterbi decoding are completely independent. At each moment, the computation of each path can be performed independently, so path selection and updates do not depend on the computation results of other paths. This independence allows the computation of each surviving path to be parallelized, making it particularly suitable for acceleration on hardware platforms such as GPUs.

[0015] Specifically, the decoder performs parallel computation on all possible paths at each time step and selects the optimal path based on the Hamming distance of each path. Path selection and updates are independent at each time step, allowing them to be processed simultaneously on multiple processing units, significantly improving computational efficiency. This parallel computation mechanism is not only suitable for path selection but can also be effectively applied to path metric updates and subsequent backtracking operations.

[0016] One type of high-speed parallel Viterbi decoder based on the CUDA architecture is used for Viterbi decoding of (n, k, K) type convolutional codes, and includes the following steps:

[0017] Step 1: Host initialization; Allocate memory space for received signals, decoding results, state transition matrix, output matrix, puncture matrix, trellis graph, etc. Initialize the state transition matrix, output matrix, polynomial matrix information structure, and puncture matrix; Calculate parallelism; Determine thread block size, number of thread blocks, code rate information of each data block to be decoded, parallel position of the signal data stream, and shared memory space size; The parallel position of the signal data stream is the initial decoding position of each data block in the received data stream; The code rate of each data block is the puncture matrix flag bit corresponding to the code rate of the decoding block.

[0018] Step 2: GPU memory allocation and data transfer; Allocate space in the GPU global memory to store data such as state transition matrix, output matrix, punching results, and mesh graph. The host will transfer initialization data from host memory to GPU global memory. Calculate the shared memory space to store path metrics, temporary bit data, mesh graph columns, and intermediate calculation results.

[0019] The trellis graph storage array stores the corresponding trellis graph of the convolutional code. Each trellis graph column stores the state of each surviving path at a given moment during forward decoding. The trellis graph storage array is a combination of trellis graph columns at each moment. For a convolutional code of (n, k, K), based on the number of surviving paths and the total number of states S=2 according to the Viterbi algorithm. K-1 .

[0020] Step 3: Start the GPU parallel de-puncturing kernel function to perform parallel de-puncturing processing on each data block. Each thread block in the kernel function determines the position of the data in the GPU's contiguous memory and the corresponding puncturing matrix based on the location of the data block and the bitrate; it then performs the de-puncturing operation according to the puncturing matrix corresponding to the data bitrate; after de-puncturing and restoring the complete data, the data block is re-divided and the initial position of the data in memory space is recorded.

[0021] Step 4: Start the GPU parallel Viterbi decoding kernel function. First, initialize the data block path metric in the shared memory of each thread block. Then, in each thread block, calculate the branch metric for each state of the corresponding data block in parallel according to the Viterbi algorithm, and update the path metric accordingly.

[0022] Step 4.1: Data Preparation and Loading. Parallel Viterbi translation functions are started on the CPU side. Each thread is responsible for processing one sub-data block, and each thread processes one state from S (the total number of states). Based on the indices recorded in SubDataBlocks and SubDataBlocksLen, the threads read the input bitstream of the sub-data block from global memory, load it into shared memory, and map outputs and source_state to shared memory and constant memory to reduce access latency.

[0023] Step 4.2: Path Metric Calculation. Each thread performs path metric calculation in shared memory according to the Add-Compare-Select (ACS) step of the Viterbi algorithm. For the current time t, construct the state transition matrix lookup table index idx = sor s + S, obtain the predecessor state s_source from source_state[idx], construct the output matrix lookup table index idx = (s >> (K-2) << (K−1)) | s0,1, obtain the corresponding output codeword from outputs[idx], calculate the branch metric, and update the path metric table. Within the thread block, synchronization instructions ensure that all state updates are completed before proceeding to the next bit calculation, thereby maintaining the correctness of the TrellisColumn.

[0024] Step 5: Compare the cumulative metrics of the paths corresponding to each current state, select the optimal path for each state, and update the corresponding grid graph column. During this process, synchronization operations are performed on the threads within the thread block to ensure the correctness and consistency of the grid graph column updates. Specifically, when processing each bit, the surviving path information (TracebackData, TD) is synchronously updated within the thread block, including the surviving predecessor state index and its corresponding branch bit information for each state. All these intermediate results are cached in shared memory to reduce latency.

[0025] Step 6: Update the grid plot at the current moment to the grid plot, then advance to the next moment, continue to calculate and update the corresponding grid plot, until all moments (data) have been processed.

[0026] Step 7: Backtracking and Sub-block Decoding. After the last data in a sub-data block has been processed, the backtracking thread within the thread block starts, backtracking backward from the state corresponding to the minimum path metric, sequentially restoring the optimal bit sequence and generating the sub-data block decoding result. The decoding result of the sub-data block is written back to the GPU global memory output buffer according to the position information recorded in SubDataBlocks.

[0027] Step 8: Assembly and Backhaul. After the GPU completes the decoding of all sub-data blocks, it assembles the decoding results of each sub-data block according to the starting position recorded in DataBlocks to form a complete data block decoding stream, and transmits the assembled decoding output stream to the CPU host through the CUDA memory copy interface; the CPU performs the final reassembly post-processing operation according to the index relationship in FrameParams.

[0028] To improve overall performance, this invention introduces multi-level storage optimization in the GPU parallel decoding stage. This includes: storing read-only data structures such as state transition tables and output tables in constant memory; storing frequently computed data such as path metric tables, temporary bit buffers, and grid graphs in shared memory; and storing input and output bit streams in global memory. This achieves hierarchical management of data within the GPU memory system, effectively reducing memory access latency and improving bandwidth utilization. Through the architecture and steps described in this application, each thread block can independently complete the path metric, backtracking, and output concatenation operations for a sub-data block. Leveraging the advantages of the CUDA multi-core architecture, large-scale parallel computing is achieved, significantly improving the throughput and real-time performance of convolutional code decoding. Crucially, this solution is applicable to parallel decoding with different code rates and codeword lengths. This high flexibility allows it to accurately adapt to the dynamic adjustment requirements of Adaptive Coding and Modulation (ACM), making it particularly suitable for ACM application scenarios. It also possesses good scalability, ease of implementation, and hardware adaptability, providing reliable support for dynamic performance optimization of communication systems. Attached Figure Description

[0029] Figure 1 This is a schematic diagram of the overall system architecture;

[0030] Figure 2 This is a schematic diagram of the data block structure;

[0031] Figure 3 This is a diagram illustrating the four-state transitions in Viterbi.

[0032] Figure 4 This is a schematic diagram of the overall algorithm flow;

[0033] Figure 5 This is a schematic diagram of the overall decoding structure. Detailed Implementation

[0034] This invention presents a high-speed parallel Viterbi decoder and decoding method based on the CUDA architecture. Through a two-level parallel architecture and multi-level storage optimization, it achieves efficient parallel decoding of multi-rate convolutional codes. The specific implementation process is described in detail below with reference to the accompanying drawings.

[0035] This invention proposes a high-speed parallel decoding method and architecture for convolutional codes based on the CUDA framework. It is applicable to convolutional code decoding scenarios using continuous coding schemes and is compatible with the convolutional code decoding requirements of various mainstream wireless communication standards. This decoding method improves parallel efficiency by segmenting the received code stream and employs de-puncturing operations at the corresponding code rate for each data block. It can efficiently support data stream decoding processing in Adaptive Coding Modulation (ACM) and has adaptability for large-scale parallel hardware implementation. The designed decoder device is built on the CUDA architecture, effectively reducing decoding latency while maximizing the utilization of GPU resources, significantly improving decoding throughput.

[0036] Step 1: The CPU host initializes the host system threads and allocates memory space. Convolutional coding parameters are initialized.

[0037] On the CPU host side, a storage space for the data to be decoded is set up to store the received bitstream data, and a storage space for the decoded data is set up to store the decoded data results. A parameter array is set up to store the parameters of each data segment, a convolutional coding parameter structure is set up and initialized, and storage spaces for the state transition matrix and output matrix are set up and initialized.

[0038] Step 2: GPU device initialization. The CPU host sends a control signal to activate the GPU device, allocating global and shared memory on the GPU device. A puncturing matrix corresponding to the specified bitrate is defined.

[0039] On the GPU device side, a storage space for the data to be decoded is set up to store the received bitstream data, and a storage space for the decoded data is set up to store the data results after decoding. Storage spaces are set up for the state transition matrix, output matrix, de-punching results, and mesh graph array data. The host transfers the initialization data, pre-computed state transition matrix, and pre-computed output matrix from host memory to GPU global memory, starts the GPU's parallel decoding thread, allocates shared memory space to store path metrics, temporary bit data, mesh graph arrays, and intermediate calculation results, and initializes the path metrics in the shared memory.

[0040] Step 3: Based on the core metadata pre-stored on the host side—including the precise offset (position information) of each data block in the overall data stream and the corresponding data transmission rate parameters—the dedicated de-puncturing kernel function on the GPU side is initiated. Leveraging the advantages of the GPU's multi-core parallel computing architecture, synchronous parallel de-puncturing operations are initiated on all data blocks, accelerating data integrity recovery through hardware computing power. After de-puncturing is completed, the data block segmentation process is immediately executed: it is recommended to use a segment length of 512-1024 symbols (this range can be dynamically adjusted according to the cache capacity and computational efficiency requirements of subsequent data processing modules) to evenly split the recovered complete data blocks. Simultaneously, the initial starting address of each segmented data in the physical memory space is recorded in real time through the memory management interface, providing precise memory location data for subsequent data scheduling, transmission, and computation operations.

[0041] Step 4: Pass the de-punctured and segmented data into the Viterbi decoding kernel function. For each segmented data block with a length of 512-1024 symbols, calculate the path metric and branch metric forward according to the Viterbi algorithm.

[0042] Step 5: Compare the cumulative metrics of the paths corresponding to each current state, select the optimal path for each state, update the corresponding grid graph, and accumulate the metrics of each surviving path; perform synchronization operations on the threads within the thread block to ensure the correctness and consistency of the grid graph updates.

[0043] Step 6: Update the grid plot at the current moment to the grid plot, then proceed to the next moment, continue to calculate and update the corresponding grid plot, until all moments (data) have been processed.

[0044] Step 7: Based on the final path accumulation metric of each surviving path, backtrack the optimal path in the mesh graph to generate decoding results and merge them to generate the final decoded output stream.

[0045] Step 8: The host transfers the merged decoded result from GPU memory back to host memory.

[0046] Below is a comparison Figure 1 as well as Figure 5 Describe the overall decoding process structure, such as Figure 1 as well as Figure 5 As shown, initialization is performed on the CPU host side, including setting the storage space for each variable, pre-calculating the state transition matrix and output matrix; simultaneously, initialization is performed on the GPU device side, including setting the GPU global memory and shared memory, defining the corresponding bitrate puncturing matrix, and then enabling the GPU parallel de-puncturing kernel function to perform de-puncturing operations on the data blocks, restoring the original encoded data, and then further segmenting the de-punctured data to improve parallelism (e.g., ...). Figure 2(As shown), finally, the parallel decoding kernel function is activated to perform Viterbi decoding, merge, output, and backhaul on the re-segmented data blocks. For this segmented parallel decoding process, steps 4, 5, 6, and 7 are performed.

[0047] In step 4, a corresponding de-puncturing operation is performed on each convolutionally encoded data block in the received data stream. By reversing the puncturing rules during encoding, the encoded data block is restored to its original encoded data length L. Then, the complete data block after de-puncturing is re-divided into blocks, and the length of the blocks satisfies L = L0 + L1 + ... + L i (where L) i The recommended value range is 1024. <L i <2048, adapting to parallel computing efficiency requirements). In steps 5 and 6, the block length is L. i The forward path metric is calculated on a data block-by-data basis. Relying on the "add-compare-select" (ACS) operation logic of the Viterbi algorithm, the cumulative path metric (PM) and surviving path metric (TD) corresponding to each surviving state at the current stage are updated in real time. In step 7, after the path metric calculation of all symbols in the data block is completed, the state corresponding to the minimum surviving path metric is selected as the backtracking starting point. The optimal decoding path is obtained by backtracking from this state along the grid diagram. Finally, the backtracking results of each block are combined and spliced ​​to form a complete decoding data stream and stored in the designated space.

[0048] Below is a comparison Figure 4 The entire parallel decoding method is described in detail.

[0049] like Figure 4 As shown, the state transition matrix and output matrix are pre-calculated during the initialization phase on the CPU host side. Since the polynomial and constraint length of the convolutional encoding are determined, the matrix sequence of the state transition matrix and output matrix of the convolutional encoding is uniquely determined. Therefore, this sequence is calculated in advance during the initialization phase, and a table lookup is performed directly during the decoding process, improving decoding efficiency.

[0050] The specific rules for constructing the state matrix and output matrix are described below.

[0051] Define the symbol for the state transition matrix. For an (n, k, K) type convolutional code, the encoder contains v = K-1 shift registers, for a total of S = 2... v There are several states. The state transition process is as follows: In the current state S... d When the input bit x enters the encoder, the original v bits in the shift register shift one bit to the right, the rightmost bit is shifted out, and the newly input bit x becomes the leftmost bit of the shift register, thus forming a new state S. d' This transfer process can be represented as:

[0052]

[0053] Where d and d' represent the numbers of the current state and the next state, respectively.

[0054] The output matrix (outputs) is constructed based on the encoder's output generation characteristics. In state S... d When the input bit x is given, the encoder outputs an n-bit codeword c(S). d Output each bit c of the codeword. (r) It is generated by the following formula:

[0055]

[0056] Where b i G represents the K-bit sequence involved in encoding (including the current input x and the shift register state). i (r) represents the i-th coefficient of the r-th generator polynomial, and ⊕ represents modulo-2 addition.

[0057] To facilitate fast lookup by the decoder, the aforementioned state transition and output relationships are pre-calculated and stored in a matrix. The state transition matrix `source_state` has a dimension of 2S×1 and uses a specific storage structure: the first S elements store the transition relationships when x=0, and the last S elements store the transition relationships when x=1. The output matrix `outputs` has a dimension of 2S×1 and uses an index mapping method for storage: the index value `idx` is derived from the current state `s` and the previous state `s`. 0,1 The decision is made jointly, i.e., idx = (s >> (K-2) << (K-1)) | s 0,1 .

[0058] The predecessor state is obtained by looking up the state transition matrix, and the index is determined by the current state, i.e., idx = s or s + S.

[0059] During the GPU initialization phase, the puncturing matrix is ​​declared using the `__constant__` keyword and stored in constant memory. Simultaneously, the following settings are made in the GPU's global memory: using the CUDA memory copy function (cudaMemcpy), the convolutional coding parameter structure defined on the CPU host (including constraint length, number of polynomials, and specific polynomials), the position information of the data block in the data stream, the bitrate information of each data point, and the pre-calculated state transition matrix `source_state` and output matrix `outputs` are copied uniformly to the GPU device.

[0060] During the de-punching stage of the decoded data, the host first copies the data to be decoded to the device's global memory, and then starts the de-punching kernel function. This kernel function restores the original encoded length of the data block in parallel according to the starting position of each data block in the data stream, and adopts a single-threaded processing mode for a single data block.

[0061] The specific implementation logic is as follows: The length of the punctured sequence for decoding is DataBlocksLen. Based on the puncturing period P and the number of bits deleted in each period D, the length of the original sequence recovered after de-puncturing can be calculated. During the restoration process, each position of the original sequence is iterated over in a loop. ,pass Determine the index of the current position in the punching cycle: When the value of the punching matrix Punc[i%P] is 1, read the corresponding bit from the punching sequence according to the location index (the location index is the corresponding position of the data bit in the data block); when the value of Punc[i%P] is 0 (that is, the bit at this position has been deleted during the punching process), fill it with a preset placeholder (such as INT_MAX).

[0062] Because the input data and the de-punched data differ in the number of bits, position mapping is required during parallel de-punching. Within a thread block, the thread flag is first mapped to data position i, and then the exact corresponding positions of the original data and the de-punched data are calculated. Establishing this mapping relationship relies on the total number of effective bits in the punch matrix (Punc_count) and the cumulative position value of the effective bits (puncU), thereby deriving the position mapping rules between the input data and the de-punched data.

[0063] The relationship between the total number of significant digits in the punch matrix, the cumulative sum of the positions of the significant digits, and the punch period and the punch matrix is ​​as follows:

[0064] Punc[]={1,1,0,1,1,0};P=10;puncU=6;Punc_count[]={0,1,0,2,3,0};The punch matrix and matrix parameters are stored in constant memory in advance.

[0065] The specific implementation of parallelism is as follows:

[0066] int32_ttmp1=L / blockDim.x, tmp2 = L % blockDim.x;

[0067] int32_t b_index, bits_locate = DataBlocks[blockIdx.x] , depuncbits_locate = 65536 * blockIdx.x;

[0068] for (i = 0; i < tmp1; i++){

[0069] b_index = Punc_count[code_rate][threadIdx.x % P [code_rate]] +

[0070] (threadIdx.x + i * blockDim.x) / P[code_rate] * puncU[code_rate];

[0071] depuncbits[depuncbits_locate + threadIdx.x + i *blockDim.x]=Punc[code_rate][threadIdx.x % P [code_rate]] == 0 ? INT_MAX :bits[bits_locate + b_index];

[0072] }

[0073] if (threadIdx.x < tmp2){

[0074] b_index = Punc_count[code_rate][threadIdx.x % P [code_rate]] +

[0075] (threadIdx.x + i * blockDim.x) / P [code_rate] *puncU[code_rate];

[0076] depuncbits[depuncbits_locate + threadIdx.x + i *blockDim.x]=Punc[code_rate][threadIdx.x % P [code_rate]] == 0 ? INT_MAX :bits[bits_locate + b_index];

[0077] }

[0078] Where L = DataBlocksLen / CodeRate (CodeRate is the coding rate), is used to characterize the target length of the punched data, code_rate is the data bit rate flag, blockDim.x, blockIdx.x and threadIdx.x are the number of threads set in the thread block, the current thread block ID and the thread ID in the current thread block, respectively, and tmp1 and tmp2 are obtained by dividing L by the thread block dimension blockDim.x and taking the remainder, respectively, as the basis for calculating the number of loops. `bits_locate` is the starting position of the input data frame corresponding to the current thread block (determined by `DataBlocks [blockIdx.x]`), and `depuncbits_locate` is the starting position of the de-punching data frame corresponding to the current thread block (calculated with an offset of 65536 bits, i.e., 65536 * blockIdx.x). `b_index` is used to store the position mapping relationship between the input bit data and the de-punching data. Its value is calculated by combining the total number of effective bits in the punch matrix, the cumulative position value of the effective bits, the thread index `threadIdx.x`, and the loop variable `i`.

[0079] The specific processing logic is as follows: First, the de-punching mapping of the batch data is completed through a loop (a total of tmp1 rounds). If the corresponding position of the punch matrix, Punc [code_rate][threadIdx.x % P [code_rate]], is 0, then the de-punching output position is assigned the value INT_MAX; otherwise, the value is taken from the input data bits [bits_locate + b_index] and filled into the corresponding de-punching output position. After the loop ends, if the current thread index threadIdx.x is less than the remainder tmp2, the remaining scattered data needs to be processed additionally, and the mapping assignment logic consistent with that in the loop is executed.

[0080] It should be noted that the number of bits in a single frame of punch-down data is limited to 65536. If the actual requirement exceeds this limit, the unit memory length can be adjusted to the maximum frame length during the memory allocation stage and when calculating the starting position of the punch-down data frame (depuncbits_locate).

[0081] After the punching operation is completed, the data block is further divided into smaller blocks, such as... Figure 2As shown, the data stream obtained from the punching process is first divided into blocks according to the data block dimension. Then, each data block is further subdivided, with the length of each subdivided block ranging from 1024 to 2048. The length of each subdivided data block and its specific location within the total data memory space are recorded in detail in a dedicated data parameter structure, providing accurate indexing and positioning for subsequent GPU-based parallel decoding and overall data concatenation operations.

[0082] The data parameter structure is stored as follows:

[0083] typedef struct{

[0084] int* CodeRate;

[0085] int* DataBlocks;

[0086] int* DataBlocksLen

[0087] int* SubDataBlocks;

[0088] int* SubDataBlocksLen;

[0089] }FrameParams

[0090] The FrameParams structure needs to be declared separately on both the CPU and GPU sides. CodeRate represents the bitrate of the data block, DataBlocks records the starting position of each data block in the overall data stream, and DataBlocksLen records the length of each data block—the memory size of these three corresponds to the size of the data blocks in the data stream.

[0091] In addition, SubDataBlocks records the starting position of each sub-block in the overall data stream after the data block is further subdivided, and SubDataBlocksLen is the length of these sub-blocks. Specifically, CodeRate, DataBlocks, and DataBlocksLen are updated after the CPU host receives the data stream; while SubDataBlocks and SubDataBlocksLen are updated after the GPU device completes the unpunching operation.

[0092] During the parallel decoding phase, the initialization of thread blocks and the number of threads needs to be completed, and a shared memory space needs to be defined before starting the parallel Viterbi decoding kernel function. Each thread block is responsible for processing one data segment, and S (total number of states) threads are allocated to each thread block.

[0093] After the parallel decoding kernel function starts, each thread on the GPU device reads the starting position and length information of the corresponding sub-data block in SubDataBlocks and SubDataBlocksLen according to its index in the thread block, and loads the corresponding input bit sequence into shared memory; at the same time, it maps the pre-calculated state transition matrix source_state and output matrix outputs from global memory to constant memory to reduce global memory access latency. For the state transition matrix index x∈{0,1}:

[0094]

[0095] Where s0 and s1 are the two predecessor states of the current state, s is the current state, and S is the total number of states.

[0096] For x∈{0,1}, output the matrix index:

[0097]

[0098] Where x is the input bit, s 0,1 s represents the predecessor state, and s represents the current state.

[0099] Each thread executes the Viterbi algorithm's Add-Compare-Select (ACS) path metric calculation process in shared memory, performing branch metrics and path metric updates on all states of the current bit, and writing the calculation results to the path metric table in shared memory. Within each thread block, synchronization instructions ensure that all threads complete the state update of the current symbol before proceeding to the next bit's calculation, guaranteeing the consistency and correctness of the Trellis column. Path metric updates follow:

[0100]

[0101] Where BM(s) = Hamming distance(r) t c out PM t (s) represents the path metric for the current state at time s, PM t-1 (s) represents the path metric of the previous state s. After the thread completes the path metric calculation for the current symbol, it writes the path metric of the optimal surviving path and its corresponding state back to shared memory.

[0102] path_metrics[threadIdx.x] = PM t (s)

[0103] rellis_column[threadIdx.x] = source_state

[0104] During parallel processing, s = threadIdx.x (where threadIdx.x is the current thread ID in the thread block), and to reduce memory access latency, the mesh graph and survival path are declared and stored using shared memory.

[0105] __ shared__ int rellis_column[]

[0106] __ shared__ int path_metrics[]

[0107] Shared memory offers architectural advantages such as low latency and high bandwidth, enabling access to path metrics and state indexes within a thread block to be completed within one clock cycle, thereby effectively avoiding the access overhead and performance bottlenecks caused by frequent access to global memory.

[0108] As all symbols gradually complete path metric calculations, the GPU retains surviving path information (mesh graph, TracebackData, TD) within each thread block, including the surviving path predecessor state index and its corresponding branch bit information for each state.

[0109] After the last symbol of the sub-data block is processed, the backtracking thread within the thread block starts, backtracking backward from the state corresponding to the minimum path metric, sequentially restoring the optimal bit sequence and generating the decoding result of the sub-data block in shared memory.

[0110] The decoding results of the sub-data blocks are written back to the output buffer of the GPU's global memory via thread blocks and stored according to the position information recorded in SubDataBlocks. After all thread blocks have completed decoding, the GPU summarizes the output results of each sub-data block and splices them together according to the starting position recorded in DataBlocks to form a complete data block decoding result.

[0111] After the GPU completes the decoding and splicing of all sub-data blocks, it transmits the index recorded in FrameParams and the spliced ​​complete decoded output stream to the CPU host's memory space through the CUDA memory copy interface. The CPU then performs the final decoding data reassembly based on the mapping relationship between DataBlocks and SubDataBlocks in FrameParams, and outputs a complete continuous decoded data stream.

[0112] To improve overall performance, this invention introduces multi-level storage optimization in the GPU parallel decoding stage: read-only data structures such as state transition matrix and output matrix are stored in constant memory; path metric table, temporary bit cache, and grid graph are stored in shared memory; and input bit sequence and output bit sequence are stored in global memory, realizing hierarchical management of data in the GPU memory system, thereby effectively reducing memory access latency and improving access bandwidth utilization.

[0113] Meanwhile, by introducing a segmented parallel decoding mechanism, parallel processing of different data blocks is supported, enabling synchronous and efficient decoding of data under multiple code rates and channels on the GPU. This is suitable for scenarios with high real-time requirements, such as adaptive coding and modulation (ACM), high-speed mobile communication, and long codeword low-latency processing. Through thread-block-level independent computation and a multi-core collaborative backtracking strategy, the integrity and continuity of the decoding results are maintained even under large-scale parallel design, ultimately outputting the correct surviving path sequence. The method in this application follows standard convolutional code decoding rules and possesses good algorithmic versatility, ease of implementation, and hardware portability.

Claims

1. A high-speed parallel Viterbi decoding method based on CUDA architecture for Viterbi decoding of (n, k, K) type convolutional codes, characterized in that, Includes the following steps: Step 1: Host initialization; Allocate memory space for received signals, decoding results, state transition matrix, output matrix, puncturing matrix, trellis graph, etc., initialize the state transition matrix, output matrix, polynomial matrix information structure and puncturing matrix, calculate parallelism, determine thread block size, number of thread blocks, code rate information of each data block to be decoded, parallel position of signal data stream and size of shared memory space; Step 2: GPU memory allocation and data transfer; Space is allocated in the GPU global memory, and the host transfers initialization data from the host memory to the GPU global memory to calculate the shared memory space. Step 3: Start the GPU parallel unpunch kernel function to perform parallel unpunch processing on each data block; Each thread block in the kernel function determines the location of the data in the GPU's contiguous memory and the corresponding puncturing matrix based on the location and bitrate of the data block, and performs de-puncturing operation according to the puncturing matrix corresponding to the data bitrate. After the punching process is completed and the complete data is recovered, the data block is re-divided and the data is recorded in the memory space at its initial position. Step 4: Start the GPU parallel Viterbi decoding kernel function; initialize the data block path metric in the shared memory of each thread block, calculate the branch metric of each state of the corresponding data block in parallel according to the Viterbi algorithm in each thread block, and update the path metric accordingly; Step 5: Compare the cumulative metrics of the paths corresponding to each current state, select the optimal path for each state, and update the corresponding grid graph. Perform synchronization operations on the threads within the thread block; Step 6: Update the grid graph of the current moment to the grid graph, advance to the next moment, continue to calculate and update the corresponding grid graph, until the processing of all moments is completed; Step 7: Based on the cumulative metric of the final path, backtrack the optimal path recorded in the mesh graph, generate the decoding results, and merge them into the final decoding output sequence; The host transfers the decoded output from the GPU memory back to the host memory.

2. The high-speed parallel Viterbi decoder based on CUDA architecture according to claim 1, characterized in that: In step 1, the state transition matrix and output matrix are pre-calculated and initialized based on the convolutional coding polynomial, the number of output symbols, the number of input symbols, and the constraint length. When calculating the state transition and the encoded output, fast output is achieved by looking up a table. The state transition matrix and output matrix of a (n, k, K) type convolutional code are constructed as follows: For the output matrix: Given the current state s (0 ≤ s < S, where S is the total number of states) and the input bit u (0 or 1), the output sequence c is uniquely determined by the generating polynomial; define the index , represent the current state s as a (K - 1)-bit binary number, where K is the constraint length, and concatenate it with the input bit u to form a K-bit sequence; for the j-th generating polynomial G j , , the output bit c j is calculated by taking the bit-weighted sum of this polynomial and modulo 2: , and all c j are combined in order to form the output sequence (where g j,k , c j are stored in binary form); For the state transition matrix: given the current state s (K-1 bits) and input bits u (0 or 1), the next state s' is determined by the following relation: The only certainty is that this relation is equivalent in binary terms to the binary representation of the current state s. Then, the input bit u is used as the least significant bit and concatenated to obtain the binary representation of the next state s'. Where, source_state[i] represents "the target state is..." Input bits are The source state s corresponding to the input u=0 is the source state when u=0. When the input u=1, the source state The source_state array is stored in the order of [all source states corresponding to u=0, all source states corresponding to u=1]. Each target state corresponds to two source states corresponding to inputs 0 and 1 respectively.

3. A high-speed parallel Viterbi decoder based on CUDA architecture according to claim 1, characterized in that: In step 2, the mesh graph storage memory is located in the GPU's global memory, and its storage structure is as follows: ； In this three-dimensional structure, the first dimension is the segmented data block (num = number of segmented data blocks), the second dimension is the grid pattern when each data block is computed forward (L = length of each segmented data block), and the third dimension is the specific survival state selected during forward computation. In the forward path metric calculation, after the forward calculation kernel cycle of a certain thread block is completed and the mesh graph is updated, the mesh graph for that step is merged into the mesh graph storage array: TD[num][L i ][s 0-S ] = grid array [0-S] (where L i (Indicates the current grid plot column index); The mesh graph column storage area and path metric storage area are located in shared memory, with a lifecycle based on the forward computation kernel and a scope based on thread blocks. For each thread block of the forward computation kernel, the storage structure of the path metric storage area is as follows: 。 4. A high-speed parallel Viterbi decoder based on CUDA architecture according to claim 1, characterized in that: In step 3, the corresponding de-puncturing operation is performed on each data block in the data stream using the de-puncturing function, so that data blocks with different code rates are all restored to the original standard-length data blocks with a code rate of 1 / n, providing an input sequence that conforms to the convolutional code structure for subsequent Viterbi decoding. The specific implementation process is as follows: For a (n, k, K) type convolutional code, its original coding code rate is In actual transmission, the bit rate can be increased by periodically punching (removing some parity bits). Let the punching period be P, and the punching mode be defined by the binary array Punc[P] (1 indicates retention, 0 indicates deletion). The bit rate after punching is... (D is the number of bits deleted per cycle); During the decoding stage, the original code rate structure needs to be restored through a de-puncturing operation: Let the length of the punctured sequence to be decoded be L, then the length of the original sequence recovered after de-puncturing is... The process based on the punch-down function is as follows: iterate through each position of the recovered sequence. ,pass Determine the index of the current position within the punching cycle, when At that time, the corresponding bits are read from the input punch sequence (values ​​are taken sequentially through the locate index). When the bit at that position is deleted during punching, it is filled with a preset placeholder (such as INT_MAX).

5. A high-speed parallel Viterbi decoder based on CUDA architecture according to claim 1, characterized in that: In step 4, when calculating the forward path metric, the forward calculation kernel is started. There are 10 threads, including num thread blocks (num is the number of data blocks), each thread block contains S threads (S is the total number of encoder states, and S=2^(K-1), K is the constraint length), where every 32 lines form a thread bundle; in step 5, during path backtracking, each thread block enables one thread to perform backtracking operation, realizing parallel backtracking of segmented data of each data block.

6. The high-speed parallel Viterbi decoder according to any one of claims 1 to 5, characterized in that: The decoder includes a hierarchical memory management and data cache optimization module. This module dynamically allocates shared memory, global memory, and constant memory according to the scale of the decoding task, the number of thread blocks, and the degree of parallelism. The hierarchical memory management module automatically adjusts the memory allocation strategy based on the data access frequency, so that high-frequency access data is preferentially stored in shared memory, and reduces memory access conflicts through memory alignment and cache prefetching technology, thereby improving memory bandwidth utilization and reducing global memory access latency.

7. The high-speed parallel Viterbi decoder according to any one of claims 1 to 5, characterized in that: The decoder employs a multi-level parallel computing strategy, including a cooperative scheduling mechanism at the task level, data block level, and thread level, wherein: Task-level parallelism is used to perform decoding operations simultaneously across multiple independent data streams or frames, thereby improving overall throughput. Block-level parallelism is used to divide a bit sequence into several blocks within a single data stream and assign them to different thread blocks for independent computation. Thread-level parallelism is used to simultaneously compute multi-state branch metrics and path updates within a single data block, reducing branch divergence with the help of SIMT architecture and Warp-level optimization techniques.

8. The high-speed parallel Viterbi decoder according to any one of claims 1 to 5, characterized in that: The decoding method further includes a dynamic path filtering and merging step. This step introduces a path metric difference threshold judgment mechanism in the backtracking stage to eliminate low-confidence paths and merge similar paths. Specifically, when the cumulative metric difference between two candidate paths is less than a preset threshold, the system merges them into a single backtracking path, thereby reducing redundant calculations.

9. The high-speed parallel Viterbi decoder according to any one of claims 1 to 5, characterized in that: The decoder includes a hierarchical synchronization and control mechanism, which employs the following strategies at different parallelism levels: Within a thread block, synchronization is achieved through shared memory barriers to ensure that the path metric calculation for the same bit processing step is completed by all threads before proceeding to the next step. Between thread blocks, global data consistency is achieved through synchronization functions or atomic operations to avoid race conditions between different blocks during mesh graph updates or path merging.

10. The high-speed parallel Viterbi decoder according to any one of claims 1 to 5, characterized in that: The decoder features multi-rate convolutional code adaptation capabilities. It generates polynomials, constraint lengths, state transition matrices, and puncturing patterns through parameterized configuration, enabling rapid switching and parallel decoding of different communication standards. The adaptation module dynamically loads the corresponding puncturing matrix based on the bitrate identifier of the input data frame and automatically adjusts the thread block size and parallelism to accommodate data blocks of different lengths and bitrates. Furthermore, this adaptation capability supports simultaneous decoding of multiple bitstreams, ensuring high throughput processing and stable decoding performance in heterogeneous communication environments.

Images