A method for implementing an SUVC enhancement layer coding based on GPU

Abstract

The application discloses a kind of SUVC enhancement layer coding methods based on GPU implementation, belong to coding field, including steps: input uncoded coefficient;Execute strip entropy calculation;Calculate the percentage of each strip entropy and the pre-allocation code rate of each strip;Calculate the coding size of all quantization coefficients of each sub-block;Multiple strips are divided into a group, according to the coding size calculation result of all sub-blocks, and the pre-allocation code rate of strip calculated, the QP decision of all strips in strip group is completed by thread bundle, the coding size of block group is calculated, and the QP decision that satisfies code control condition is found;According to the coding size of each block group calculated, the starting position of each block group in coding byte stream is calculated;According to starting position, accurate block group position is determined to carry out stream writing, and block group coding is executed;The output of coding data.The application reasonably and efficiently utilizes hardware resources, and higher production efficiency is achieved.

Description

Technical Field

[0001] This invention relates to the field of encoding, and more specifically, to a GPU-based SUVC enhancement layer encoding method. Background Technology

[0002] We propose the SUVC enhancement layer coding technique, which was designed from the outset to reduce and optimize the logical complexity of the algorithm, making its implementation on GPUs possible.

[0003] In recent years, the rapid development of GPUs has enabled them to far surpass CPUs in computing power. General-purpose computing frameworks, such as NVIDIA's CUDA, not only have the advantage of being easy to learn but also have a good ecosystem, making them a suitable parallel computing framework for encoding and decoding. In CUDA, the following are generally accepted rules: threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y, blockDim.x, and blockDim.y are all built-in constants in the CUDA programming framework, only related to the GPU code construction process; _syncthreads() is a synchronization method on a thread bundle; __syncwarp() is a synchronization method within every 32 threads; _shfl_up_sync() and __shfl_down_sync() are methods for implementing thread communication within every 32 threads in a thread bundle; atomicAdd() is an atomic addition method in the CUDA programming framework; atomicOr() is an atomic OR method in the CUDA programming framework; shared memory has a lifecycle consistent with the thread bundle and can only be accessed by all threads within the same thread bundle, requiring the __shared__ keyword; global memory has a lifecycle from the call to allocate memory to its active release, and can be accessed by all threads on the GPU.

[0004] Therefore, based on SUVC enhancement layer coding and CUDA technology, we propose a novel GPU-based enhancement layer coding method using discrete wavelet transform. Summary of the Invention

[0005] The purpose of this invention is to overcome the shortcomings of the prior art and provide a GPU-based SUVC enhancement layer encoding method that makes reasonable and efficient use of hardware resources to achieve higher production efficiency.

[0006] The objective of this invention is achieved through the following solution:

[0007] A GPU-based SUVC enhancement layer encoding method includes the following steps:

[0008] Step 1: Input the uncoded coefficients of all stripes on global memory;

[0009] Step 2: Calculate the strip entropy and sum the strip entropies;

[0010] Step 3: Based on the entropy of each stripe and the sum of the entropies of the stripes calculated in Step 2, calculate the percentage of the entropy of each stripe. The percentage data of the entropy of each stripe is used to calculate the pre-allocated code rate of each stripe.

[0011] Step 4: Calculate the encoding size of all quantization coefficients for each sub-block;

[0012] Step 5: Divide multiple stripes into a group. Based on the coding size calculation results of all sub-blocks in Step 4 and the pre-allocated code rate of the stripe calculated in Step 3, the thread bundle completes the QP decision for all stripes in the stripe group, calculates the coding size of the block group, and finds the QP decision that satisfies the code control conditions.

[0013] Step 6: Based on the encoded size of each block group calculated in Step 5, calculate the starting position of each block group in the encoded byte stream;

[0014] Step 7: Determine the accurate block group position based on the starting position and write the bitstream, then perform block group encoding;

[0015] Step 8: Output the encoded data obtained in step 7.

[0016] Further, in step 2, the strip entropy calculation includes the following sub-steps:

[0017] Step 2.1, the entropy calculation of each sub-block in the stripe, specifically includes: each thread reads the four coding coefficients of the sub-block in sequence, then performs Hadamard transform, then quantizes the result of the Hadamard transform, and sums the absolute values ​​of multiple coefficients of the quantization result to obtain the entropy of the current four coefficients;

[0018] Then, every 32 threads out of the 64 threads use __shfl_up_sync to obtain the entropy sum of the first set number of threads; __shfl_up_sync is the method for implementing thread communication within every 32 threads in a CUDA thread bundle;

[0019] Then, the last thread in every 32 threads out of the 64 threads adds the final entropy sum to the shared storage corresponding to the final entropy sum of the current sub-block using atomicAdd;

[0020] Then, use __syncthreads() to synchronize thread bundles; __syncthreads() is a synchronization method on CUDA thread bundles.

[0021] Then, the thread with threadIdx.x equal to 0 will transfer the sum of the entropy of the corresponding sub-block in the shared memory to the global memory corresponding to the sub-block index; threadIdx.x is a CUDA built-in constant;

[0022] Step 2.2, calculate the strip entropy and sum the entropies of all strips. The specific process includes:

[0023] Step 2.2.1: Use index = threadIdx.x as the initial index value for each thread, set the loop span to 128, perform one loop, traverse the entropy of all sub-blocks of the current strip, and temporarily store the sum of sub-block entropy obtained by each thread in a local variable;

[0024] Step 2.2.2: Transfer the sub-block entropy stored in the local variable to the shared memory aBlockEntropy[threadIdx.x], and call __syncthreads() to ensure the access security of the shared memory;

[0025] Step 2.2.3: Each thread loops with stride = 64 as the initial condition. Each time, stride is divided by 2 until stride is 0. Each thread whose threadIdx.x is less than stride executes aBlockEntropy[threadIdx.x] += aBlockEntropy[threadIdx.x + stride]. At the end of each loop, __syncthreads() is called to ensure the safety of access to shared memory.

[0026] Step 2.2.4: When the loop ends, the entropy of the current stripe is stored in aBlockEntropy[0]. The thread with threadIdx.x equal to 0 transfers the entropy of the stripe to the global storage corresponding to the stripe index, and uses atomicAdd to accumulate the entropy of the current stripe to the global storage corresponding to the entropy of all stripes. atomicAdd is the atomic addition method in CUDA.

[0027] Furthermore, in step 3, the pre-allocated bitrate calculation formula for each stripe is as follows:

[0028] (Expected coding size per frame) / (Number of stripes) * 0.33 + (Expected coding size per frame) * (Entropy of each stripe) / (Entropy of all stripes) * 0.67.

[0029] Further, in step 4, calculating the encoding size of all quantization coefficients for each sub-block includes the following sub-steps:

[0030] Step 4.1: Each thread sequentially accesses four consecutive coefficients in the 16x16 sub-block and performs a Hadamard transform on the four coefficients.

[0031] Step 4.2: Each thread starts the loop with qp_index=0 as the starting index and sets the loop span to 1 until qp_index reaches the maximum number of QPs. The thread then calculates the sub-block encoding and the encoding size of the four encoding modes for all QPs.

[0032] Step 4.3: Each thread performs quantization calculations on the four coefficients of the Hadamard transform according to the quantization coefficients corresponding to the current qp_index;

[0033] Step 4.4: Analyze whether the quantization result has four coefficients that are not all zero, and whether it is a 0001 coefficient group. At the same time, reset the shared storage corresponding to the non-zero flags of the 16 coefficients group and the non-zero flags of the 64 coefficients group, and use atomicOr to assign and modify these two data. atomicOr is a method in CUDA.

[0034] Step 4.5: Each thread calculates the third and fourth encoding modes respectively. The number of encoded bits consists of two parts: the number of P-encoded bits and the number of S-encoded bits, denoted as PBudget3 and SBudget3, PBudget4 and SBudget4. The first and second encoding modes consist of three parts: the number of Z-encoded bits, the number of P-encoded bits, and the number of S-encoded bits, denoted as ZBudget1, PBudget1 and SBudget1; ZBudget2, PBudget2 and SBudget2.

[0035] Step 4.6: Each thread starts the loop with stride equal to 16 as the starting index. In each loop, stride is divided by 2 until stride equals 1. Then, __shfl_down_sync is applied to the 10 variables in step 4.5 to obtain the sum of the same variable in the 32 threads. __shfl_down_sync is the method for implementing thread communication within every 32 threads in a CUDA thread bundle.

[0036] Step 4.7: After step 4.6, the 10 variables in the first thread of every 32 threads are the sum of all the variables in the 32 threads. This thread then uses the atomicAdd method on these 10 variables to finally obtain the four encoding modes of the current qp_index for the current sub-block, the encoding size of each part, and saves this part of the data in the current shared memory.

[0037] After repeating steps 4.8 and 4.2 to 4.7, we will obtain the four encoding modes of all QPs in the current sub-block, as well as the number of encoded bits for each part. This result is stored in the shared storage and then transferred to the global storage of each QP corresponding to the current sub-block.

[0038] Further, in step 5, the process of using thread bundles to complete QP decisions for all stripes within a stripe group until a QP decision satisfying the code control conditions is found specifically includes the following sub-steps:

[0039] Step 5.1: Use bgIdx = threadIdx.y as the initial condition for the loop, with a loop span of 32 in each iteration. The loop terminates when bgIdx is greater than or equal to the number of stripe block groups. Iterate through each block group within the stripe. threadIdx.y is a built-in constant in CUDA.

[0040] Step 5.2: Use blockIdx = threadIdx.x as the initial condition for the loop, with each loop spanning 16. The loop terminates when blockIdx is greater than or equal to the number of sub-blocks in the block group, and iterates through each word block in the block group; threadIdx.y is a built-in constant in CUDA.

[0041] Step 5.3: Each thread accesses the global storage of the four encodings of all QPs obtained in Step 4, finds the encoding size of the current QP of the stripe for the four encoding modes of the current sub-block, compares the encoding mode corresponding to the minimum number of encoding bits with the total number of encoding bits of the sub-block, and records it in the global storage as the code control result of the current sub-block; and records the Z encoding bit count, P encoding bit count, S encoding bit count and the number of empty sub-blocks of the current thread respectively.

[0042] Step 5.4: Following steps 5.2 and 5.3, a group of 16 threads in the horizontal direction completes the encoding calculation for a block group. The statistical results are then processed using `__shfl_up_sync` to obtain the final Z-coded bit count, P-coded bit count, S-coded bit count, and number of empty sub-blocks for the current block group. The final total statistical results for the block group are stored in the thread with `threaIdx.x` equal to 15. `__shfl_up_sync` is a method in CUDA for implementing thread communication within every 32 threads in a thread bundle.

[0043] Step 5.5: If the final number of empty sub-blocks is equal to the total number of sub-blocks in the current block group, it means that the current block group is full of empty sub-blocks. Only the block group header encoding size of the current block group needs to be recorded. Otherwise, the number of Z-encoded bits, P-encoded bits, and S-encoded bits are aligned by bytes, added together, and the block group header encoding size is added and saved.

[0044] Step 5.6: After steps 5.1 to 5.5, the obtained block group size is saved to the shared memory. After using __syncthreads to ensure thread bundle synchronization, all the results in the shared memory are summed up and the stripe header encoding size is added to obtain the final stripe encoding size. _syncthreads is the synchronization method on CUDA thread bundles.

[0045] Step 5.7: Repeat the above steps multiple times until a suitable QP is found that meets the code control conditions.

[0046] Further, in step 6, calculating the starting position of each block group in the encoded byte stream includes the following sub-steps:

[0047] Step 6.1: Each thread in the thread bundle uses bgIdx = threadIdx.x as the loop index, with each loop spanning 512. The loop termination condition for each thread is that bgIdx is greater than or equal to the total size of the block group; threadIdx.x is a built-in constant in CUDA.

[0048] Step 6.2: When bgIdx is greater than 0 and less than the total size of the block group, access the size of the block group at the previous index position bgIdx-1 and use it as the byte offset ByteOffset of the current block group; and when the current block group is the first block group of the stripe, add the stripe header encoding size to ByteOffset.

[0049] Step 6.3: Group the threads into 32 groups. Within each group of 32 threads, each thread starts with a stride of 1. In each loop, the stride is multiplied by 2. The loop ends when the stride equals 32. The `__shfl_up_sync()` function is used multiple times to obtain the ByteOffset values ​​of the other threads within the 32 threads. At the end of the loop, each thread's ByteOffset is the sum of the ByteOffset values ​​of all threads preceding it within the 32 threads. `__shfl_up_sync()` is the method for implementing thread communication within each group of 32 threads.

[0050] Step 6.4: The last thread in every 32 threads obtains the ByteOffset and saves it to the shared memory aBGOffset[threadIdx.x>>5]. Then, it calls __syncthreads() to ensure safe access to the shared memory. __syncthreads() is a synchronization method on the thread bundle.

[0051] Step 6.5: Obtain the byte offset preceding each index position in aBGOffset and accumulate it to its own position;

[0052] Step 6.6: If threadIdx.x divided by 32 is greater than 1, then add aBGOffset[(threadIdx.x / 32) - 1] to the index offset of the previous block group corresponding to the current bgIdx, and add the final statistical result SliceByteOffset of the size of the sub-blocks that have been traversed in the current stripe.

[0053] Step 6.7: Call __syncthreads() to synchronize and add aBGOffset[7] to SliceByteOffset to complete the calculation of the starting position of the block group encoded bitstream.

[0054] Further, in step 7, determining the accurate block group position based on the starting position for bitstream writing and performing block group encoding includes the following sub-steps:

[0055] Step 7.1: Determine if the current block group is an empty block group. If it is an empty block group, the thread with threadIdx.x equal to 0 will directly write two bytes, 0 and 2, at the beginning of the block group. If it is not empty, proceed to step 7.2.

[0056] Step 7.2: Traverse the sub-blocks in the block group and encode each sub-block in order; the sub-block encoding is divided into three parts, namely Z encoding, P encoding and S encoding.

[0057] Furthermore, Z encoding includes sub-block encoding mode encoding. The sub-block encoding mode encoding process is as follows: the sub-block encoding mode is always completed by the thread with threadIdx.x equal to 0, and the encoded data is written according to the sub-block encoding size and sub-block encoding mode determined by the code control module; if the sub-block encoding size is 0, then one bit 0 is directly written to the high bit of the currently written byte stream; if the sub-block encoding size is not 0, then the encoding write [bit value, number of bits] is determined according to the sub-block encoding mode, and the corresponding relationship is as follows: encoding mode 1 corresponds to [0b10, 2], encoding mode 2 corresponds to [0b110, 3], encoding mode 3 corresponds to [0b1110, 4], and encoding mode 4 corresponds to [0b1111, 4]; after the sub-block encoding mode is processed, the number of bits of the block header encoding is accumulated into the ZBitOffset parameter stored in the shared memory.

[0058] Furthermore, Z-coding includes encoding of the non-zero flag bit of the sub-block coding coefficients; the specific process for encoding the non-zero flag bit of the sub-block coding coefficients is as follows:

[0059] When the sub-block encoding mode is 3 or 4, all coefficients of the current sub-block need to be encoded one by one. The non-zero flags of the four coefficients of the current sub-block are all initialized to 1 to facilitate the unified processing of subsequent P encoding and S encoding.

[0060] When the sub-block encoding mode is 1, the sub-block encoding coefficients are grouped into sets of 4 coefficients; and the process is performed three times according to the index coef4Idx=threadIdx.x with a span of 32.

[0061] The first traversal writes the non-zero flag of the coefficient group, and the bit value is written at the offset position of coef4Idx + ZBitOffset; after the first traversal, ZBitOffset is incremented by 64; ZbitOffset represents the offset position of the Z-encoded bit, and "+" indicates addition;

[0062] The second traversal write is the 0001 mode flag bit. The offset position of the written bit value is calculated using __shfl_up_sync. During each loop, a thread with a threshold of 31 needs to update ZBitOffset. __shfl_up_sync is a method for implementing thread communication within every 32 threads in a thread bundle.

[0063] The third iteration writes the 3 bit values ​​corresponding to the 0001 pattern. The offset position of the written bit value is also calculated using __shfl_up_sync. In each loop, a thread with threadIdx.x equal to 31 is also needed to update ZBitOffset; threadIdx.x is a built-in constant in CUDA.

[0064] When the sub-block encoding mode is 2, the specific steps include the following six steps:

[0065] The first step is to group the sub-block coding coefficients into groups of four, and perform a traversal once according to the index coef4Idx=threadIdx.x with a span of 32; then count the non-zero flags for groups of four coefficients, groups of sixteen coefficients, and groups of sixty-four coefficients respectively. The latter two are counted using atomicOr.

[0066] The second step is to encode the non-zero flag bits, which are grouped into 64 coefficients. Here, only the first 4 of the 32 threads need to participate in writing the bit values ​​into the code stream.

[0067] The third step is to encode the non-zero flag bit with 16 coefficients as a group. Here, only the first 16 threads need to participate in writing the bit value into the code stream, and the bit value offset position needs to be statistically calculated using __shfl_up_sync.

[0068] The fourth step is to encode the non-zero flag bits, which are grouped into four coefficients. Here, the 32 threads need to perform the encoding in two loops, and the bit value offset is calculated using __shfl_up_sync.

[0069] The fifth step is to encode the 0001 pattern flag bit with four coefficients as a group. Here, the 32 threads need to perform the encoding in two loops, and the bit value offset position needs to be statistically calculated using __shfl_up_sync.

[0070] The sixth step is to encode the 3 bit values ​​corresponding to the 0001 pattern. Here, the 32 threads need to perform the encoding in two loops, and the bit value offset position needs to be statistically calculated using __shfl_up_sync.

[0071] Furthermore, P-encoding and S-encoding are completed by encoding coefficients through a single loop of encoding sub-blocks. The specific process includes: using coefIdx=threadIdx.x as the initial index, performing a loop traversal with a span of 32, accessing one coefficient each time, and each thread calculating the P-encoded bit value and number of bits, S-encoded bit value and number of bits of the current coefficient. If the non-zero flag bit of the four coefficients corresponding to the current coefficient is 1, the current calculation result is retained; otherwise, the current calculation result is set to zero.

[0072] Then, __shfl_up_sync is used to count the bit positions of P and S respectively, and atomicOr is used to write the calculation results into the bitstream. At the same time, PBitOffset and SbitOffset are updated after each loop. PBitOffset and SbitOffset represent the offset positions of P and S encoded bits respectively.

[0073] The beneficial effects of this invention include:

[0074] This invention can fully utilize the heterogeneous parallel computing of CPU+GPU to schedule different computing needs, assigning highly parallel and computationally intensive tasks to GPUs for execution, while assigning tasks with low parallelism and high logical complexity to CPUs for execution, thus making reasonable and efficient use of hardware resources and achieving higher production efficiency. Attached Figure Description

[0075] To more clearly illustrate the technical solutions in the embodiments of the present invention or the prior art, the drawings used in the description of the embodiments or the prior art will be briefly introduced below. Obviously, the drawings described below are only some embodiments of the present invention. For those skilled in the art, other drawings can be obtained based on these drawings without creative effort.

[0076] Figure 1 This is a flowchart of a method according to an embodiment of the present invention. Detailed Implementation

[0077] All features disclosed in all embodiments of this specification, or steps in all methods or processes implied in the disclosure, may be combined and / or extended or replaced in any way, except for mutually exclusive features and / or steps.

[0078] The specific implementation steps of this invention are as follows:

[0079] The input consists of the uncoded coefficients of all stripes in global video memory.

[0080] The entropy calculation of each sub-block in the stripe is taken as an example of a 16x16 sub-block. The entropy calculation of each sub-block can be completed by 64 threads, and each thread can calculate the entropy of four coefficients.

[0081] A thread bundle consisting of 512 threads is used as a concurrent scheduling module on the GPU. One thread bundle can perform entropy calculation for 8 sub-blocks. The thread distribution in the thread bundle is two-dimensional, so the value range of threadIdx.x is [0, 63] and the value range of threadIdx.y is [0, 7].

[0082] The process of 64 threads calculating a sub-block is as follows: each thread reads the four coding coefficients of the sub-block in sequence, then performs a Hadamard transform, then quantizes the result of the Hadamard transform, and sums the absolute values ​​of the four coefficients of the quantization result to obtain the entropy of the current four coefficients.

[0083] The entropy of each 32 threads out of the 64 threads is obtained by using __shfl_up_sync.

[0084] The last thread in every 32 threads out of the 64 threads adds the final entropy sum to the shared storage corresponding to the final entropy sum of the current sub-block using atomicAdd.

[0085] Use __syncthreads() to synchronize thread bundles and ensure safe access to shared memory.

[0086] Threads with threadIdx.x equal to 0 will transfer the sum of the entropy of the corresponding sub-block in the shared memory to the global memory corresponding to the sub-block index.

[0087] 2.2 Calculate the strip entropy and sum the entropies of all stripes.

[0088] On the GPU, 128 threads are grouped into a thread bundle to complete the entropy statistics of one strip; the thread distribution of the thread bundle is divided in one dimension, that is, the value range of threadIdx.x is [0, 127].

[0089] 2.2.1 Using `index = threadIdx.x` as the initial index for each thread, and a loop span of 128, perform one iteration to traverse the entropy of all sub-blocks in the current stripe. The sum of the sub-block entropy obtained by each thread is temporarily stored in a local variable.

[0090] 2.2.2 Transfer the sub-block entropy stored in the local variable to the shared memory aBlockEntropy[threadIdx.x], and call __syncthreads() to ensure the access safety of the shared memory.

[0091] 2.2.3 Each thread uses a loop index stride = 64 as the initial condition. Each time stride is divided by 2 until stride is 0. Each thread whose threadIdx.x is less than stride executes aBlockEntropy[threadIdx.x] += aBlockEntropy[threadIdx.x + stride]. At the end of each loop, __syncthreads() is called to ensure the safety of access to shared memory.

[0092] 2.2.4 When the loop ends, the entropy of the current stripe is stored in aBlockEntropy[0]. The thread with threadIdx.x equal to 0 transfers the entropy of the stripe to the global storage corresponding to the stripe index, and uses atomicAdd to accumulate the entropy of the current stripe to the global storage corresponding to the entropy of all stripes.

[0093] The stripe rate pre-allocation is based on the entropy of each stripe calculated in the previous step, and the sum of the entropies of all stripes. The percentage of entropy for each stripe is used to calculate the pre-allocated code rate for each stripe.

[0094] The formula for calculating the pre-allocated bit rate for each stripe is as follows:

[0095] (Expected coding size per frame) / (Number of stripes) * 0.33f + (Expected coding size per frame) * (Entropy of each stripe) / (Entropy of all stripes) * 0.67f;

[0096] The encoding size of all quantization coefficients within each sub-block is calculated by enumerating the encoding size of all quantization coefficients in the sub-block. Each sub-block is completed by 64 threads, and each thread bundle consists of 128 threads. The thread distribution within each thread bundle is two-dimensional, so the value range of threadIdx.x is [0, 63], and the value range of threadIdx.y is [0, 1]. The following process is illustrated using the enumeration calculation of one sub-block.

[0097] 4.1 Each thread accesses four consecutive coefficients in a 16x16 sub-block in sequence and performs a Hadamard transformation on the four coefficients.

[0098] 4.2 Each thread starts the loop with qp_index=0 as the starting index, with a loop span of 1, until qp_index reaches the maximum number of QPs, and then calculates the sub-block encoding corresponding to all QPs. The encoding size for the four encoding modes is also considered.

[0099] 4.3 Each thread performs quantization calculations on the four coefficients of the Hadamard transform according to the quantization coefficients corresponding to the current qp_index.

[0100] 4.4 Analyze whether the quantization result has four coefficients that are not all zero, and whether it is a 0001 coefficient group. At the same time, reset the shared storage corresponding to the non-zero flags of the 16 coefficients group and the non-zero flags of the 64 coefficients group, and use atomicOr to assign and modify these two data.

[0101] 4.5 Each thread calculates the third and fourth encoding modes separately. Since all coefficients need to be encoded one by one, the number of encoded bits only includes two parts: the number of P-encoded bits and the number of S-encoded bits, which can be denoted as PBudget3 and SBudget3, PBudget4 and SBudget4. The first and second encoding modes, due to the presence of all-zero coefficient groups that are not encoded, contain three parts: the number of Z-encoded bits, the number of P-encoded bits, and the number of S-encoded bits, which can be denoted as ZBudget1, PBudget1 and SBudget1; ZBudget2, PBudget2 and SBudget2, respectively.

[0102] 4.6 Each thread starts the loop with stride equal to 16 as the starting index. In each loop, stride is divided by 2 until stride equals 1. Then, __shfl_down_sync is applied to the 10 variables in step 4.5 to obtain the sum of the same variable in the 32 threads.

[0103] 4.7 After step 3.6, the 10 variables in the first thread of every 32 threads are the sum of all the variables in the 32 threads. This thread then uses atomicAdd on these 10 variables to finally obtain the four encoding modes of the current qp_index for the current sub-block and the encoding size of each part. This data is stored in the current shared memory.

[0104] After repeating steps 4.84.2 to 4.7, you will obtain the four encoding modes of all QPs in the current sub-block and the number of encoded bits for each part. This result is stored in shared storage and needs to be transferred to the global storage corresponding to each QP in the current sub-block.

[0105] In stripe code control, multiple stripes are grouped together. Based on the coding size calculation results of all sub-blocks in step 4 and the stripe pre-allocated code rate calculated in step 2, a single thread bundle completes the QP decision for all stripes within a stripe group. The QP decision method is to find the first QP such that the stripe coding size is less than the stripe pre-allocated code rate, or the maximum QP value is reached.

[0106] On the GPU, each thread bundle completes the QP decision for one stripe. Each thread bundle consists of 512 threads, which can be distributed in two dimensions. The value of threadIdx.x ranges from [0, 15], and the value of threadIdx.y ranges from [0, 31]. Sixteen threads in the horizontal direction each complete the optimal encoding mode decision for one sub-block and calculate the encoding size of the block group. In the vertical direction, the stripe block group can be traversed according to threadIdx.y. The steps for calculating the stripe encoding size are as follows:

[0107] 5.1 The loop starts with bgIdx = threadIdx.y as the initial condition, with a loop span of 32 in each iteration. The loop terminates when bgIdx is greater than or equal to the number of stripe block groups, and each block group within the stripe is traversed.

[0108] 5.2 The loop starts with blockIdx = threadIdx.x as the initial condition, with each loop spanning 16. The loop terminates when blockIdx is greater than or equal to the number of sub-blocks in the block group, and each word block in the block group is traversed.

[0109] 5.3 For each thread, the global storage of all four encodings of the QPs in the sub-block is accessed. The encoding size of the current QP for the current sub-block in each of the four encoding modes is found. The encoding mode corresponding to the minimum number of encoding bits is compared with the total number of encoding bits in the sub-block, and recorded in the global storage as the code control result of the current sub-block. The Z-encoding bit count, P-encoding bit count, S-encoding bit count, and number of empty sub-blocks for the current thread are also recorded.

[0110] 5.4 Following steps 5.2 and 5.3, a group of 16 threads in the horizontal direction completed the encoding calculation for a block group. The statistical results were processed using __shfl_up_sync to obtain the final Z-coded bit count, P-coded bit count, S-coded bit count, and number of empty sub-blocks for the current block group. The final total statistical results for the block group were then stored in the thread with `threaIdx.x` equal to 15.

[0111] 5.5 If the final number of empty sub-blocks is equal to the total number of sub-blocks in the current block group, it means that the current block group is full of empty sub-blocks. In this case, it is only necessary to record the block group header encoding size of the current block group. Otherwise, the number of Z-encoded bits, P-encoded bits, and S-encoded bits are aligned by bytes, added together, and the block group header encoding size is added and saved.

[0112] 5.6 After steps 4.1 to 4.5, the obtained block group size is saved to the shared memory. After using __syncthreads to ensure thread bundle synchronization, all the results in the shared memory are summed up, and the stripe header encoding size is added to obtain the final stripe encoding size.

[0113] 5.7 Repeat the above steps multiple times until a suitable QP is found that meets the code control conditions.

[0114] The starting position calculation of the block group encoding bitstream is based on the encoding size of each block group obtained by the code control module. It calculates the starting position of each block group in the encoded byte stream to ensure that the subsequent encoding module can write the bitstream from the accurate block group position. 512 threads on the GPU form a thread bundle to complete the calculation of the byte starting position of all sub-blocks. The thread division in the thread bundle is one-dimensional, that is, the value range of threadIdx.x is [0, 511].

[0115] 6.1 Each thread in the thread bundle uses bgIdx = threadIdx.x as the loop index, and each loop spans 512. The loop termination condition for each thread is that bgIdx is greater than or equal to the total size of the block group.

[0116] 6.2 When bgIdx is greater than 0 and less than the total block size, access the block size at the previous index position bgIdx-1 and use it as the byte offset ByteOffset for the current block. Furthermore, if the current block is the first block of a stripe, add the stripe header encoding size to ByteOffset.

[0117] 6.332 threads are grouped together. Within each group of 32 threads, each thread starts with a stride of 1. In each loop, the stride is multiplied by 2. The loop ends when the stride equals 32. After multiple loops, __shfl_up_sync() is used to obtain the ByteOffset values ​​of other threads in the group of 32 threads. At the end of the loop, the ByteOffset of each thread is the sum of the ByteOffsets of all threads preceding the current thread in the group of 32 threads.

[0118] 6.4 The last thread in every 32 threads receives the ByteOffset and saves it to the shared memory aBGOffset[threadIdx.x>>5]. __syncthreads() is called to ensure safe access to the shared memory.

[0119] 6.5 Following the same method as step 5.3, aBGOffset obtains the byte offset before each index position in aBGOffset and accumulates it into its own position, that is, aBGOffset[n] += aBGOffset[n-1]+ aBGOffset[n-2]+...+aBGOffset[0], where the value of n is in the range of [1, 7].

[0120] 6.6 If threadIdx.x divided by 32 is greater than 1, then the index offset of the preceding block group corresponding to the current bgIdx needs to be added to aBGOffset[(threadIdx.x / 32) - 1], and the final statistical result of the size of the sub-blocks that have been traversed by the current stripe, SliceByteOffset, also needs to be added.

[0121] 6.7 Call __syncthreads() to synchronize and add aBGOffset to SliceByteOffset[7].

[0122] Block group encoding: A block group is the smallest encoding unit that can be scheduled in parallel in SUVC. The GPU implementation can encode one block group using 32 threads. Multiple block groups can be processed simultaneously within a thread bundle. Given the limited shared memory on the GPU, processing two block groups simultaneously is more appropriate, meaning a thread bundle consists of 64 threads, with 32 threads encoding one block group. Here, we assume all threads in the thread bundle are distributed in two dimensions, so the value range of `threadIdx.x` is [0, 31], and the value range of `threadIdx.y` is [0, 1]. The encoding process for each block group is as follows:

[0123] 7.1 First, determine whether the current block group is an empty block group. If it is an empty block group, the thread with threadIdx.x equal to 0 will directly write two bytes 0 and 2 at the beginning of the block group. If it is not empty, continue to execute the following steps.

[0124] 7.2 Traverse the sub-blocks in the block group and encode each sub-block in order. The sub-block encoding is mainly divided into three parts: Z encoding, P encoding, and S encoding.

[0125] 7.2.1Z encoding consists of two parts: sub-block encoding mode encoding and sub-block encoding coefficient non-zero flag bit encoding.

[0126] 7.2.1.1 Block Header Encoding

[0127] Sub-block encoding is always performed by the thread with threadIdx.x equal to 0, and the encoded data is written based on the sub-block encoding size and sub-block encoding mode determined by the code control module. If the sub-block encoding size is 0, one bit 0 is directly written to the high bit of the currently written byte stream; if the sub-block encoding size is not 0, the encoding write [bit value, number of bits] is determined according to the sub-block encoding mode, with the following correspondences: encoding mode 1 corresponds to [0b10, 2], encoding mode 2 corresponds to [0b110, 3], encoding mode 3 corresponds to [0b1110, 4], and encoding mode 4 corresponds to [0b1111, 4]. After the sub-block encoding mode is processed, the number of bits in the block header encoding is accumulated and added to the ZBitOffset parameter stored in shared memory.

[0128] 7.2.1.2 Encoding of Non-zero Flag Bits for Sub-block Coding Coefficients

[0129] 7.2.1.2.1 The sub-block encoding mode is 3 or 4. All coefficients of the current sub-block need to be encoded one by one. The non-zero flag bits of the four coefficients of the current sub-block can be initialized to 1 to facilitate the unified processing of subsequent P encoding and S encoding.

[0130] 7.2.1.2.2 When the sub-block coding mode is 1, the sub-block coding coefficients are grouped into sets of four coefficients.

[0131] The algorithm is traversed three times based on the index coef4Idx=threadIdx.x and the span is 32.

[0132] The first iteration writes the non-zero flag of the coefficient group, with the bit value offset being coef4Idx + ZBitOffset. After the first iteration, ZBitOffset is incremented by 64.

[0133] The second traversal write is the 0001 mode flag bit. The offset position of the written bit value needs to be calculated using __shfl_up_sync. During each loop, a thread with a threshold of 31 needs to update ZBitOffset.

[0134] The third iteration writes the 3 bit values ​​corresponding to the 0001 pattern. The offset position of the written bit value also needs to be calculated using __shfl_up_sync. In each loop, the thread with thresholdx.x equal to 31 also needs to update ZBitOffset.

[0135] 7.2.1.2.3 When the sub-block encoding mode is 2, it is completed in 6 steps.

[0136] The first step is to group the sub-block coding coefficients into groups of four and iterate through them once using the index coef4Idx=threadIdx.x with a span of 32. Then, count the non-zero flags for groups of four, sixteen, and sixty-four coefficients respectively. The latter two require the use of atomicOr to perform the counting.

[0137] The second step is to encode the non-zero flag bits, which are grouped into 64 coefficients. Here, only the first 4 of the 32 threads need to participate in writing the bit values ​​into the code stream.

[0138] The third step is to encode the non-zero flag bits, which are grouped into 16 coefficients. Here, only the first 16 threads need to participate in writing the bit values ​​into the code stream, and the bit value offset position needs to be statistically calculated using __shfl_up_sync.

[0139] The fourth step is to encode the non-zero flag bits, which are grouped into four coefficients. Here, the 32 threads need to perform the encoding in two loops, and the bit offset position needs to be statistically calculated using __shfl_up_sync.

[0140] The fifth step is to encode the 0001 pattern flag bit, which consists of four coefficients. Here, the 32 threads need to perform the encoding in two loops, and the bit offset position needs to be calculated using __shfl_up_sync.

[0141] The sixth step is to encode the 3 bit values ​​corresponding to the 0001 pattern. Here, the 32 threads need to perform the encoding in two loops, and the bit value offset position needs to be statistically calculated using __shfl_up_sync.

[0142] 7.2.2 P and S encoding

[0143] P-encoding and S-encoding can be completed by encoding coefficients in a single loop. Using `coefIdx=threadIdx.x` as the initial index, a loop with a span of 32 is performed, iterating through one coefficient at a time. Each thread calculates the P-encoded bit value and number of bits, and the S-encoded bit value and number of bits for the current coefficient. If the non-zero flag corresponding to a group of four coefficients is 1, the current calculation result is retained; otherwise, the result is set to zero. Then, `__shfl_up_sync` is used to count the P and S encoded bit positions, and `atomicOr` is used to write the calculation results into the bitstream. Simultaneously, `PBitOffset` and `SBitOffset` are updated after each loop.

[0144] The output will be the encoded data obtained in step 7.

[0145] It should be noted that, within the scope of protection defined in the claims of this invention, the following embodiments can be combined and / or extended or replaced in any logical manner from the above specific embodiments, such as the disclosed technical principles, disclosed technical features or implicitly disclosed technical features.

[0146] Example 1

[0147] A GPU-based SUVC enhancement layer encoding method includes the following steps:

[0148] Step 1: Input the uncoded coefficients of all stripes on global memory;

[0149] Step 2: Calculate the strip entropy and sum the strip entropies;

[0150] Step 3: Based on the entropy of each stripe and the sum of the entropies of the stripes calculated in Step 2, calculate the percentage of the entropy of each stripe. The percentage data of the entropy of each stripe is used to calculate the pre-allocated code rate of each stripe.

[0151] Step 4: Calculate the encoding size of all quantization coefficients for each sub-block;

[0152] Step 5: Divide multiple stripes into a group. Based on the coding size calculation results of all sub-blocks in Step 4 and the pre-allocated code rate of the stripe calculated in Step 3, the thread bundle completes the QP decision for all stripes in the stripe group, calculates the coding size of the block group, and finds the QP decision that satisfies the code control conditions.

[0153] Step 6: Based on the encoded size of each block group calculated in Step 5, calculate the starting position of each block group in the encoded byte stream;

[0154] Step 7: Determine the accurate block group position based on the starting position and write the bitstream, then perform block group encoding;

[0155] Step 8: Output the encoded data obtained in step 7.

[0156] Example 2

[0157] Based on Example 1, step 2, the strip entropy calculation includes the following sub-steps:

[0158] Step 2.1, the entropy calculation of each sub-block in the stripe, specifically includes: each thread reads the four coding coefficients of the sub-block in sequence, then performs Hadamard transform, then quantizes the result of the Hadamard transform, and sums the absolute values ​​of multiple coefficients of the quantization result to obtain the entropy of the current four coefficients;

[0159] Then, every 32 threads out of the 64 threads use __shfl_up_sync to obtain the entropy sum of the first set number of threads; __shfl_up_sync is the method for implementing thread communication within every 32 threads in a CUDA thread bundle;

[0160] Then, the last thread in every 32 threads out of the 64 threads adds the final entropy sum to the shared storage corresponding to the final entropy sum of the current sub-block using atomicAdd;

[0161] Then, use __syncthreads() to synchronize thread bundles; __syncthreads() is a synchronization method on CUDA thread bundles.

[0162] Then, the thread with threadIdx.x equal to 0 will transfer the sum of the entropy of the corresponding sub-block in the shared memory to the global memory corresponding to the sub-block index; threadIdx.x is a CUDA built-in constant;

[0163] Step 2.2, calculate the strip entropy and sum the entropies of all strips. The specific process includes:

[0164] Step 2.2.1: Use index = threadIdx.x as the initial index value for each thread, set the loop span to 128, perform one loop, traverse the entropy of all sub-blocks of the current strip, and temporarily store the sum of sub-block entropy obtained by each thread in a local variable;

[0165] Step 2.2.2: Transfer the sub-block entropy stored in the local variable to the shared memory aBlockEntropy[threadIdx.x], and call __syncthreads() to ensure the access security of the shared memory;

[0166] Step 2.2.3: Each thread loops with stride = 64 as the initial condition. Each time, stride is divided by 2 until stride is 0. Each thread whose threadIdx.x is less than stride executes aBlockEntropy[threadIdx.x] += aBlockEntropy[threadIdx.x + stride]. At the end of each loop, __syncthreads() is called to ensure the safety of access to shared memory.

[0167] Step 2.2.4: When the loop ends, the entropy of the current stripe is stored in aBlockEntropy[0]. The thread with threadIdx.x equal to 0 transfers the entropy of the stripe to the global storage corresponding to the stripe index, and uses atomicAdd to accumulate the entropy of the current stripe to the global storage corresponding to the entropy of all stripes. atomicAdd is the atomic addition method in CUDA.

[0168] Example 3

[0169] Based on Example 1, in step 3, the pre-allocated code rate calculation formula for each stripe is as follows:

[0170] (Expected coding size per frame) / (Number of stripes) * 0.33 + (Expected coding size per frame) * (Entropy of each stripe) / (Entropy of all stripes) * 0.67.

[0171] Example 4

[0172] Based on Example 1, step 4, which involves calculating the encoding size of all quantization coefficients for each sub-block, includes the following sub-steps:

[0173] Step 4.1: Each thread sequentially accesses four consecutive coefficients in the 16x16 sub-block and performs a Hadamard transform on the four coefficients.

[0174] Step 4.2: Each thread starts the loop with qp_index=0 as the starting index and sets the loop span to 1 until qp_index reaches the maximum number of QPs. The thread then calculates the sub-block encoding and the encoding size of the four encoding modes for all QPs.

[0175] Step 4.3: Each thread performs quantization calculations on the four coefficients of the Hadamard transform according to the quantization coefficients corresponding to the current qp_index;

[0176] Step 4.4: Analyze whether the quantization result has four coefficients that are not all zero, and whether it is a 0001 coefficient group. At the same time, reset the shared storage corresponding to the non-zero flags of the 16 coefficients group and the non-zero flags of the 64 coefficients group, and use atomicOr to assign and modify these two data. atomicOr is a method in CUDA.

[0177] Step 4.5: Each thread calculates the third and fourth encoding modes respectively. The number of encoded bits consists of two parts: the number of P-encoded bits and the number of S-encoded bits, denoted as PBudget3 and SBudget3, PBudget4 and SBudget4. The first and second encoding modes consist of three parts: the number of Z-encoded bits, the number of P-encoded bits, and the number of S-encoded bits, denoted as ZBudget1, PBudget1 and SBudget1; ZBudget2, PBudget2 and SBudget2.

[0178] Step 4.6: Each thread starts the loop with stride equal to 16 as the starting index. In each loop, stride is divided by 2 until stride equals 1. Then, __shfl_down_sync is applied to the 10 variables in step 4.5 to obtain the sum of the same variable in the 32 threads. __shfl_down_sync is the method for implementing thread communication within every 32 threads in a CUDA thread bundle.

[0179] Step 4.7: After step 4.6, the 10 variables in the first thread of every 32 threads are the sum of all the variables in the 32 threads. This thread then uses the atomicAdd method on these 10 variables to finally obtain the four encoding modes of the current qp_index for the current sub-block, the encoding size of each part, and saves this part of the data in the current shared memory.

[0180] After repeating steps 4.8 and 4.2 to 4.7, we will obtain the four encoding modes of all QPs in the current sub-block, as well as the number of encoded bits for each part. This result is stored in the shared storage and then transferred to the global storage of each QP corresponding to the current sub-block.

[0181] Example 5

[0182] Based on Example 1, in step 5, the process of using a thread bundle to complete the QP decision for all stripes within the stripe group until a QP decision that satisfies the code control conditions specifically includes the following sub-steps:

[0183] Step 5.1: Use bgIdx = threadIdx.y as the initial condition for the loop, with a loop span of 32 in each iteration. The loop terminates when bgIdx is greater than or equal to the number of stripe block groups. Iterate through each block group within the stripe. threadIdx.y is a built-in constant in CUDA.

[0184] Step 5.2: Use blockIdx = threadIdx.x as the initial condition for the loop, with each loop spanning 16. The loop terminates when blockIdx is greater than or equal to the number of sub-blocks in the block group, and iterates through each word block in the block group; threadIdx.y is a built-in constant in CUDA.

[0185] Step 5.3: Each thread accesses the global storage of the four encodings of all QPs obtained in Step 4, finds the encoding size of the current QP of the stripe for the four encoding modes of the current sub-block, compares the encoding mode corresponding to the minimum number of encoding bits with the total number of encoding bits of the sub-block, and records it in the global storage as the code control result of the current sub-block; and records the Z encoding bit count, P encoding bit count, S encoding bit count and the number of empty sub-blocks of the current thread respectively.

[0186] Step 5.4: Following steps 5.2 and 5.3, a group of 16 threads in the horizontal direction completes the encoding calculation for a block group. The statistical results are then processed using `__shfl_up_sync` to obtain the final Z-coded bit count, P-coded bit count, S-coded bit count, and number of empty sub-blocks for the current block group. The final total statistical results for the block group are stored in the thread with `threaIdx.x` equal to 15. `__shfl_up_sync` is a method in CUDA for implementing thread communication within every 32 threads in a thread bundle.

[0187] Step 5.5: If the final number of empty sub-blocks is equal to the total number of sub-blocks in the current block group, it means that the current block group is full of empty sub-blocks. Only the block group header encoding size of the current block group needs to be recorded. Otherwise, the number of Z-encoded bits, P-encoded bits, and S-encoded bits are aligned by bytes, added together, and the block group header encoding size is added and saved.

[0188] Step 5.6: After steps 5.1 to 5.5, the obtained block group size is saved to the shared memory. After using __syncthreads to ensure thread bundle synchronization, all the results in the shared memory are summed up and the stripe header encoding size is added to obtain the final stripe encoding size. _syncthreads is the synchronization method on CUDA thread bundles.

[0189] Step 5.7: Repeat the above steps multiple times until a suitable QP is found that meets the code control conditions.

[0190] Example 6

[0191] Based on Example 1, step 6, which involves calculating the starting position of each block group in the encoded byte stream, includes the following sub-steps:

[0192] Step 6.1: Each thread in the thread bundle uses bgIdx = threadIdx.x as the loop index, with each loop spanning 512. The loop termination condition for each thread is that bgIdx is greater than or equal to the total size of the block group; threadIdx.x is a built-in constant in CUDA.

[0193] Step 6.2: When bgIdx is greater than 0 and less than the total size of the block group, access the size of the block group at the previous index position bgIdx-1 and use it as the byte offset ByteOffset of the current block group; and when the current block group is the first block group of the stripe, add the stripe header encoding size to ByteOffset.

[0194] Step 6.3: Group the threads into 32 groups. Within each group of 32 threads, each thread starts with a stride of 1. In each loop, the stride is multiplied by 2. The loop ends when the stride equals 32. The `__shfl_up_sync()` function is used multiple times to obtain the ByteOffset values ​​of the other threads within the 32 threads. At the end of the loop, each thread's ByteOffset is the sum of the ByteOffset values ​​of all threads preceding it within the 32 threads. `__shfl_up_sync()` is the method for implementing thread communication within each group of 32 threads.

[0195] Step 6.4: The last thread in every 32 threads obtains the ByteOffset and saves it to the shared memory aBGOffset[threadIdx.x>>5]. Then, it calls __syncthreads() to ensure safe access to the shared memory. __syncthreads() is a synchronization method on the thread bundle.

[0196] Step 6.5: Obtain the byte offset preceding each index position in aBGOffset and accumulate it to its own position;

[0197] Step 6.6: If threadIdx.x divided by 32 is greater than 1, then add aBGOffset[(threadIdx.x / 32) - 1] to the index offset of the previous block group corresponding to the current bgIdx, and add the final statistical result SliceByteOffset of the size of the sub-blocks that have been traversed in the current stripe.

[0198] Step 6.7: Call __syncthreads() to synchronize and add aBGOffset[7] to SliceByteOffset to complete the calculation of the starting position of the block group encoded bitstream.

[0199] Example 7

[0200] Based on Example 1, step 7, which involves determining the accurate block group position based on the starting position for bitstream writing and performing block group encoding, includes the following sub-steps:

[0201] Step 7.1: Determine if the current block group is an empty block group. If it is an empty block group, the thread with threadIdx.x equal to 0 will directly write two bytes, 0 and 2, at the beginning of the block group. If it is not empty, proceed to step 7.2.

[0202] Step 7.2: Traverse the sub-blocks in the block group and encode each sub-block in order; the sub-block encoding is divided into three parts, namely Z encoding, P encoding and S encoding.

[0203] Example 8

[0204] Based on Example 7, Z encoding includes sub-block encoding mode encoding. The sub-block encoding mode encoding process includes: the sub-block encoding mode is always completed by the thread with threadIdx.x equal to 0, and the encoded data is written by the sub-block encoding size and sub-block encoding mode determined by the code control module; if the sub-block encoding size is 0, then one bit 0 is directly written to the high bit of the currently written byte stream; if the sub-block encoding size is not 0, then the encoding write [bit value, number of bits] is determined according to the sub-block encoding mode, and the corresponding relationship is as follows: encoding mode 1 corresponds to [0b10, 2], encoding mode 2 corresponds to [0b110, 3], encoding mode 3 corresponds to [0b1110, 4], and encoding mode 4 corresponds to [0b1111, 4]; after the sub-block encoding mode is processed, the number of bits of the block header encoding is accumulated to the ZBitOffset parameter stored in the shared storage.

[0205] Example 9

[0206] Based on Example 7, Z-coding includes encoding of the non-zero flag bit of the sub-block coding coefficients; the encoding of the non-zero flag bit of the sub-block coding coefficients specifically includes the following process:

[0207] When the sub-block encoding mode is 3 or 4, all coefficients of the current sub-block need to be encoded one by one. The non-zero flags of the four coefficients of the current sub-block are all initialized to 1 to facilitate the unified processing of subsequent P encoding and S encoding.

[0208] When the sub-block encoding mode is 1, the sub-block encoding coefficients are grouped into sets of 4 coefficients; and the process is performed three times according to the index coef4Idx=threadIdx.x with a span of 32.

[0209] The first traversal writes the non-zero flag of the coefficient group, and the bit value is written at the offset position of coef4Idx + ZBitOffset; after the first traversal, ZBitOffset is incremented by 64; ZbitOffset represents the offset position of the Z-encoded bit, and "+" indicates addition;

[0210] The second traversal write is the 0001 mode flag bit. The offset position of the written bit value is calculated using __shfl_up_sync. During each loop, a thread with a threshold of 31 needs to update ZBitOffset. __shfl_up_sync is a method for implementing thread communication within every 32 threads in a thread bundle.

[0211] The third iteration writes the 3 bit values ​​corresponding to the 0001 pattern. The offset position of the written bit value is also calculated using __shfl_up_sync. In each loop, a thread with threadIdx.x equal to 31 is also needed to update ZBitOffset; threadIdx.x is a built-in constant in CUDA.

[0212] When the sub-block encoding mode is 2, the specific steps include the following six steps:

[0213] The first step is to group the sub-block coding coefficients into groups of four, and perform a traversal once according to the index coef4Idx=threadIdx.x with a span of 32; then count the non-zero flags for groups of four coefficients, groups of sixteen coefficients, and groups of sixty-four coefficients respectively. The latter two are counted using atomicOr.

[0214] The second step is to encode the non-zero flag bits, which are grouped into 64 coefficients. Here, only the first 4 of the 32 threads need to participate in writing the bit values ​​into the code stream.

[0215] The third step is to encode the non-zero flag bit with 16 coefficients as a group. Here, only the first 16 threads need to participate in writing the bit value into the code stream, and the bit value offset position needs to be statistically calculated using __shfl_up_sync.

[0216] The fourth step is to encode the non-zero flag bits, which are grouped into four coefficients. Here, the 32 threads need to perform the encoding in two loops, and the bit value offset is calculated using __shfl_up_sync.

[0217] The fifth step is to encode the 0001 pattern flag bit with four coefficients as a group. Here, the 32 threads need to perform the encoding in two loops, and the bit value offset position needs to be statistically calculated using __shfl_up_sync.

[0218] The sixth step is to encode the 3 bit values ​​corresponding to the 0001 pattern. Here, the 32 threads need to perform the encoding in two loops, and the bit value offset position needs to be statistically calculated using __shfl_up_sync.

[0219] Example 10

[0220] Based on Example 7, P-encoding and S-encoding are completed by encoding coefficients through a single loop of encoding sub-blocks. The specific process includes: using coefIdx=threadIdx.x as the initial index, performing a loop traversal with a loop span of 32, accessing one coefficient each time, and each thread calculating the P-encoding bit value and number of bits, S-encoding bit value and number of bits of the current coefficient. If the non-zero flag bit of the four coefficients corresponding to the current coefficient is 1, the current calculation result is retained; otherwise, the current calculation result is set to zero.

[0221] Then, __shfl_up_sync is used to count the bit positions of P and S respectively, and atomicOr is used to write the calculation results into the bitstream. At the same time, PBitOffset and SbitOffset are updated after each loop. PBitOffset and SbitOffset represent the offset positions of P and S encoded bits respectively.

[0222] The units described in the embodiments of the present invention can be implemented in software or hardware, and the described units can also be located in a processor. The names of these units do not necessarily limit the specific unit itself.

[0223] According to one aspect of the present invention, a computer program product or computer program is provided, the computer program product or computer program including computer instructions stored in a computer-readable storage medium. A processor of a computer device reads the computer instructions from the computer-readable storage medium, and executes the computer instructions, causing the computer device to perform the methods provided in the various optional implementations described above.

[0224] In another aspect, embodiments of the present invention also provide a computer-readable medium, which may be included in the electronic device described in the above embodiments; or it may exist independently and not assembled into the electronic device. The computer-readable medium carries one or more programs, which, when executed by the electronic device, cause the electronic device to perform the methods described in the above embodiments.

Claims

1. A GPU-based SUVC enhancement layer encoding method, characterized in that, Includes the following steps: Step 1: Input the uncoded coefficients of all stripes on global memory; Step 2: Calculate the strip entropy and sum the strip entropies; Step 3: Based on the entropy of each stripe and the sum of the entropies of the stripes calculated in Step 2, calculate the percentage of the entropy of each stripe. The percentage data of the entropy of each stripe is used to calculate the pre-allocated code rate of each stripe. Step 4: Calculate the encoding size of all quantization coefficients for each sub-block; Step 5: Divide multiple stripes into a group. Based on the coding size calculation results of all sub-blocks in Step 4 and the pre-allocated code rate of the stripe calculated in Step 3, the thread bundle completes the QP decision for all stripes in the stripe group, calculates the coding size of the block group, and finds the QP decision that satisfies the code control conditions. Step 6: Based on the encoded size of each block group calculated in Step 5, calculate the starting position of each block group in the encoded byte stream; Step 7: Determine the accurate block group position based on the starting position and write the bitstream, then perform block group encoding; Step 8: Output the encoded data obtained in step 7.

2. The GPU-based SUVC enhancement layer encoding method according to claim 1, characterized in that, In step 2, the strip entropy calculation includes the following sub-steps: Step 2.1, the entropy calculation of each sub-block in the stripe, specifically includes: each thread reads the four coding coefficients of the sub-block in sequence, then performs Hadamard transform, then quantizes the result of the Hadamard transform, and sums the absolute values ​​of multiple coefficients of the quantization result to obtain the entropy of the current four coefficients; Then, every 32 threads out of the 64 threads use __shfl_up_sync to obtain the entropy sum of the first set number of threads; __shfl_up_sync is the method for implementing thread communication within every 32 threads in a CUDA thread bundle; Then, the last thread in every 32 threads out of the 64 threads adds the final entropy sum to the shared storage corresponding to the final entropy sum of the current sub-block using atomicAdd; Then, use __syncthreads() to synchronize thread bundles; __syncthreads() is a synchronization method on CUDA thread bundles. Then, the thread with threadIdx.x equal to 0 will transfer the sum of the entropy of the corresponding sub-block in the shared memory to the global memory corresponding to the sub-block index; threadIdx.x is a CUDA built-in constant; Step 2.2, calculate the strip entropy and sum the entropies of all strips. The specific process includes: Step 2.2.1: Use index = threadIdx.x as the initial index value for each thread, set the loop span to 128, perform one loop, traverse the entropy of all sub-blocks of the current strip, and temporarily store the sum of sub-block entropy obtained by each thread in a local variable; Step 2.2.2: Transfer the sub-block entropy stored in the local variable to the shared memory aBlockEntropy[threadIdx.x], and call __syncthreads() to ensure the access security of the shared memory; Step 2.2.3: Each thread loops with stride = 64 as the initial condition. Each time, stride is divided by 2 until stride is 0. Each thread whose threadIdx.x is less than stride executes aBlockEntropy[threadIdx.x] += aBlockEntropy[threadIdx.x + stride]. At the end of each loop, __syncthreads() is called to ensure the safety of access to shared memory. Step 2.2.4: When the loop ends, the entropy of the current stripe is stored in aBlockEntropy[0]. The thread with threadIdx.x equal to 0 transfers the entropy of the stripe to the global storage corresponding to the stripe index, and uses atomicAdd to accumulate the entropy of the current stripe to the global storage corresponding to the entropy of all stripes. atomicAdd is the atomic addition method in CUDA.

3. The GPU-based SUVC enhancement layer encoding method according to claim 1, characterized in that, In step 3, the pre-allocated bitrate calculation formula for each stripe is as follows: (Expected coding size per frame) / (Number of stripes) * 0.33 + (Expected coding size per frame) * (Entropy of each stripe) / (Entropy of all stripes) * 0.

67.

4. The GPU-based SUVC enhancement layer encoding method according to claim 1, characterized in that, In step 4, calculating the encoding size of all quantization coefficients for each sub-block includes the following sub-steps: Step 4.1: Each thread sequentially accesses four consecutive coefficients in the 16x16 sub-block and performs a Hadamard transform on the four coefficients. Step 4.2: Each thread starts the loop with qp_index=0 as the starting index and sets the loop span to 1 until qp_index reaches the maximum number of QPs. The thread then calculates the sub-block encoding and the encoding size of the four encoding modes for all QPs. Step 4.3: Each thread performs quantization calculations on the four coefficients of the Hadamard transform according to the quantization coefficients corresponding to the current qp_index; Step 4.4: Analyze whether the quantization result has four coefficients that are not all zero, and whether it is a 0001 coefficient group. At the same time, reset the shared storage corresponding to the non-zero flags of the 16 coefficients group and the non-zero flags of the 64 coefficients group, and use atomicOr to assign and modify these two data. atomicOr is a method in CUDA. Step 4.5: Each thread calculates the third and fourth encoding modes respectively. The number of encoded bits consists of two parts: the number of P-encoded bits and the number of S-encoded bits, denoted as PBudget3 and SBudget3, PBudget4 and SBudget4. The first and second encoding modes consist of three parts: the number of Z-encoded bits, the number of P-encoded bits, and the number of S-encoded bits, denoted as ZBudget1, PBudget1 and SBudget1; ZBudget2, PBudget2 and SBudget2. Step 4.6: Each thread starts the loop with stride equal to 16 as the starting index. In each loop, stride is divided by 2 until stride equals 1. Then, __shfl_down_sync is applied to the 10 variables in step 4.5 to obtain the sum of the same variable in the 32 threads. __shfl_down_sync is the method for implementing thread communication within every 32 threads in a CUDA thread bundle. Step 4.7: After step 4.6, the 10 variables in the first thread of every 32 threads are the sum of all the variables in the 32 threads. This thread then uses the atomicAdd method on these 10 variables to finally obtain the four encoding modes of the current qp_index for the current sub-block, the encoding size of each part, and saves this part of the data in the current shared memory. After repeating steps 4.8 and 4.2 to 4.7, we will obtain the four encoding modes of all QPs in the current sub-block, as well as the number of encoded bits for each part. This result is stored in the shared storage and then transferred to the global storage of each QP corresponding to the current sub-block.

5. The GPU-based SUVC enhancement layer encoding method according to claim 1, characterized in that, In step 5, the QP decision for all stripes within the stripe group is completed by the thread bundle until a QP decision that satisfies the code control conditions is found. This specifically includes the following sub-steps: Step 5.1: Use bgIdx=threadIdx.y as the initial condition for the loop, with a loop span of 32 each time. The loop terminates when bgIdx is greater than or equal to the number of stripe block groups, and iterates through each block group within the stripe. threadIdx.y is a built-in constant in CUDA; Step 5.2: Use blockIdx = threadIdx.x as the initial condition for the loop, with a loop span of 16 each time. The loop terminates when blockIdx is greater than or equal to the number of sub-blocks in the block group, and iterates through each word block of the block group. threadIdx.y is a built-in constant in CUDA; Step 5.3: Each thread accesses the global storage of the four encodings of all QPs obtained in Step 4, finds the encoding size of the current QP of the stripe for the four encoding modes of the current sub-block, compares the encoding mode corresponding to the minimum number of encoding bits with the total number of encoding bits of the sub-block, and records it in the global storage as the code control result of the current sub-block; and records the Z encoding bit count, P encoding bit count, S encoding bit count and the number of empty sub-blocks of the current thread respectively. Step 5.4: Following steps 5.2 and 5.3, a group of 16 threads in the horizontal direction completed the encoding calculation of a block group; the final Z-encoded bit count, P-encoded bit count, S-encoded bit count, and number of empty sub-blocks of the current block group were obtained by using __shfl_up_sync on the statistical results, and the final total statistical results of the block group were saved in the thread where threaIdx.x equals 15; _shfl_up_sync is a method in CUDA that enables thread communication within every 32 threads in a thread bundle. Step 5.5: If the final number of empty sub-blocks is equal to the total number of sub-blocks in the current block group, it means that the current block group is full of empty sub-blocks. Only the block group header encoding size of the current block group needs to be recorded. Otherwise, the number of Z-encoded bits, P-encoded bits, and S-encoded bits are aligned by bytes, added together, and the block group header encoding size is added and saved. Step 5.6: After steps 5.1 to 5.5, the obtained block group size is saved to the shared memory. After using __syncthreads to ensure thread bundle synchronization, all the results in the shared memory are summed up and the stripe header encoding size is added to obtain the final stripe encoding size. _syncthreads is the synchronization method on CUDA thread bundles. Step 5.7: Repeat the above steps multiple times until a suitable QP is found that meets the code control conditions.

6. The GPU-based SUVC enhancement layer encoding method according to claim 1, characterized in that, In step 6, calculating the starting position of each block group in the encoded byte stream includes the following sub-steps: Step 6.1: Each thread in the thread bundle uses bgIdx = threadIdx.x as the loop index, with a loop span of 512. The loop termination condition for each thread is that bgIdx is greater than or equal to the total size of the block group. threadIdx.x is a built-in constant in CUDA; Step 6.2: When bgIdx is greater than 0 and less than the total size of the block group, access the size of the block group at the previous index position bgIdx-1 and use it as the byte offset ByteOffset of the current block group; and when the current block group is the first block group of the stripe, add the stripe header encoding size to ByteOffset. Step 6.3: Group the threads into 32 groups. Within each group of 32 threads, each thread starts with a stride of 1. In each loop, the stride is multiplied by 2. The loop ends when the stride equals 32. The `__shfl_up_sync()` function is used multiple times to obtain the ByteOffset values ​​of the other threads within the 32 threads. At the end of the loop, each thread's ByteOffset is the sum of the ByteOffset values ​​of all threads preceding it within the 32 threads. `__shfl_up_sync()` is the method for implementing thread communication within each group of 32 threads. Step 6.4: The last thread in every 32 threads obtains the ByteOffset and saves it to the shared memory aBGOffset[threadIdx.x>>5]. Then, it calls __syncthreads() to ensure safe access to the shared memory. __syncthreads() is a synchronization method on the thread bundle. Step 6.5: Obtain the byte offset preceding each index position in aBGOffset and accumulate it to its own position; Step 6.6: If threadIdx.x divided by 32 is greater than 1, then add aBGOffset[(threadIdx.x / 32) - 1] to the index offset of the previous block group corresponding to the current bgIdx, and add the final statistical result SliceByteOffset of the size of the sub-blocks that have been traversed in the current stripe. Step 6.7: Call __syncthreads() to synchronize and add aBGOffset[7] to SliceByteOffset to complete the calculation of the starting position of the block group encoded bitstream.

7. The GPU-based SUVC enhancement layer encoding method according to claim 1, characterized in that, In step 7, determining the accurate block group position based on the starting position for bitstream writing and performing block group encoding includes the following sub-steps: Step 7.1: Determine if the current block group is an empty block group. If it is an empty block group, the thread with threadIdx.x equal to 0 will directly write two bytes, 0 and 2, at the beginning of the block group. If it is not empty, proceed to step 7.

2. Step 7.2: Traverse the sub-blocks in the block group and encode each sub-block in order; The sub-block encoding is divided into three parts: Z encoding, P encoding, and S encoding.

8. The GPU-based SUVC enhancement layer encoding method according to claim 7, characterized in that, Z-encoding includes sub-block encoding mode encoding. The sub-block encoding mode encoding process is as follows: the sub-block encoding mode is always completed by the thread with threadIdx.x equal to 0, and the encoded data is written according to the sub-block encoding size and sub-block encoding mode determined by the code control module; if the sub-block encoding size is 0, then one bit 0 is directly written to the high bit of the currently written byte stream; if the sub-block encoding size is not 0, then the encoding write [bit value, number of bits] is determined according to the sub-block encoding mode, and the corresponding relationship is as follows: encoding mode 1 corresponds to [0b10, 2], encoding mode 2 corresponds to [0b110, 3], encoding mode 3 corresponds to [0b1110, 4], and encoding mode 4 corresponds to [0b1111, 4]; after the sub-block encoding mode is processed, the number of bits of the block header encoding is accumulated into the ZBitOffset parameter stored in the shared memory.

9. The GPU-based SUVC enhancement layer encoding method according to claim 7, characterized in that, Z-coding includes encoding of the non-zero flag bit of the sub-block encoding coefficients; the specific process for encoding the non-zero flag bit of the sub-block encoding coefficients is as follows: When the sub-block encoding mode is 3 or 4, all coefficients of the current sub-block need to be encoded one by one. The non-zero flags of the four coefficients of the current sub-block are all initialized to 1 to facilitate the unified processing of subsequent P encoding and S encoding. When the sub-block encoding mode is 1, the sub-block encoding coefficients are grouped into sets of 4 coefficients; and the process is performed three times according to the index coef4Idx=threadIdx.x with a span of 32. The first traversal writes the non-zero flag of the coefficient group, and the bit value is written at the offset position of coef4Idx + ZBitOffset; after the first traversal, ZBitOffset is incremented by 64; ZbitOffset represents the offset position of the Z-encoded bit, and "+" indicates addition; The second traversal write is the 0001 mode flag bit. The offset position of the written bit value is calculated using __shfl_up_sync. During each loop, a thread with a threshold of 31 needs to update ZBitOffset. __shfl_up_sync is a method for implementing thread communication within every 32 threads in a thread bundle. The third iteration writes the 3 bit values ​​corresponding to the 0001 pattern. The offset position of the written bit value is also calculated using __shfl_up_sync. In each loop, a thread with threadIdx.x equal to 31 is also needed to update ZBitOffset; threadIdx.x is a built-in constant in CUDA. When the sub-block encoding mode is 2, the specific steps include the following six steps: The first step is to group the sub-block coding coefficients into groups of four, and perform a traversal once according to the index coef4Idx=threadIdx.x with a span of 32; then count the non-zero flags for groups of four coefficients, groups of sixteen coefficients, and groups of sixty-four coefficients respectively. The latter two are counted using atomicOr. The second step is to encode the non-zero flag bits, which are grouped into 64 coefficients. Here, only the first 4 of the 32 threads need to participate in writing the bit values ​​into the code stream. The third step is to encode the non-zero flag bit with 16 coefficients as a group. Here, only the first 16 threads need to participate in writing the bit value into the code stream, and the bit value offset position needs to be statistically calculated using __shfl_up_sync. The fourth step is to encode the non-zero flag bits, which are grouped into four coefficients. Here, the 32 threads need to perform the encoding in two loops, and the bit value offset is calculated using __shfl_up_sync. The fifth step is to encode the 0001 pattern flag bit with four coefficients as a group. Here, the 32 threads need to perform the encoding in two loops, and the bit value offset position needs to be statistically calculated using __shfl_up_sync. The sixth step is to encode the 3 bit values ​​corresponding to the 0001 pattern. Here, the 32 threads need to perform the encoding in two loops, and the bit value offset position needs to be statistically calculated using __shfl_up_sync.

10. The GPU-based SUVC enhancement layer encoding method according to claim 7, characterized in that, P-encoding and S-encoding are completed by encoding coefficients through a single loop of encoding sub-blocks. The specific process includes: using coefIdx=threadIdx.x as the initial index, performing a loop traversal with a span of 32, accessing one coefficient at a time, and each thread calculating the P-encoding bit value and number of bits, S-encoding bit value and number of bits of the current coefficient. If the non-zero flag bit of the four coefficients corresponding to the current coefficient is 1, the current calculation result is retained; otherwise, the current calculation result is set to zero. Then, __shfl_up_sync is used to count the bit positions of P and S respectively, and atomicOr is used to write the calculation results into the bitstream. At the same time, PBitOffset and SbitOffset are updated after each loop. PBitOffset and SbitOffset represent the offset positions of P and S encoded bits respectively.

Images