A general-purpose GPU-oriented JPEG decoding batch processing and double-buffer pipeline scheduling method

Abstract

The application relates to a general-purpose GPU-oriented JPEG decoding batch processing and double-buffer pipeline scheduling method, belongs to the field of heterogeneous computing and image processing, and comprises the following steps: obtaining code stream data of a to-be-decoded JPEG image; performing entropy decoding and inverse quantization on the code stream data by a CPU to obtain a plurality of DCT coefficient blocks; if the size of the to-be-decoded JPEG image is not smaller than a preset size threshold, the DCT coefficient blocks are cached in a ring buffer located in a host memory; the state of the ring buffer is continuously monitored; when a preset batch triggering condition is met and a GPU is detected to be available, the DCT coefficient blocks in the ring buffer are batch copied to GPU display memory; based on the DCT coefficient blocks in the GPU display memory, integer IDCT transformation and color space conversion calculation are performed by the GPU to obtain an original image corresponding to the to-be-decoded JPEG image, and the original image is read back to the host memory; wherein the double-buffer pipeline is adopted to perform batch copying of the DCT coefficient blocks, integer IDCT transformation, color space conversion calculation and reading back of the original image. Efficient real-time decoding is realized.

Description

Technical Field

[0001] This invention relates to the field of heterogeneous computing and image processing technology, and in particular to a batch processing and double-buffered pipeline scheduling method for JPEG (Joint Photographic Experts Group, an image format that achieves high compression ratio by discarding information that is not sensitive to the human eye) decoding for general-purpose GPUs (Graphics Processing Units). Background Technology

[0002] JPEG images are widely used in digital photography, medical imaging, security monitoring, and web page images due to their high compression ratio and good visual quality. IDCT (Inverse Discrete Cosine Transform) and color conversion are computationally intensive operations in JPEG decoding, accounting for more than 60% of the total decoding time and are key aspects for performance optimization.

[0003] Currently, mainstream high-performance JPEG decoding methods heavily rely on dedicated hardware or specific ecosystems. For example, NVIDIA's nvJPEG library utilizes its dedicated NVDEC video decoding unit and CUDA (Compute Unified Device Architecture) parallel computing architecture on its GPUs to achieve millisecond-level batch decoding on data center GPUs. However, this method is heavily tied to NVIDIA's closed-source drivers and hardware, making it impossible to run on domestic GPU platforms (such as domestically produced GPUs). Similarly, Intel's QSV (Quick Sync Video, Intel's hardware acceleration technology) and AMD VCE (Video Coding Engine, AMD's hardware acceleration technology) also rely on fixed-function hardware modules within their respective GPUs.

[0004] For general-purpose GPUs that lack dedicated encoding / decoding units (such as some domestically produced GPUs that only support the OpenCL 1.2 general-purpose computing software standard and do not yet have dedicated JPEG hardware decoding units), existing JPEG decoding methods face significant bottlenecks: (1) Excessive data transfer overhead: The traditional method of using the OpenCL interface to call the GPU hardware unit to complete JPEG decoding usually triggers a data transfer from the host to the device after decoding each minimum coding unit (MCU). This results in the PCIe (Peripheral Component Interconnect Express) bus bandwidth not being fully utilized and transmission delays, causing a performance inversion phenomenon where "hardware acceleration is actually slower". (2) Low efficiency of floating-point operation: Most open-source algorithms are implemented using floating-point IDCT algorithm, but the floating-point unit performance of domestic GPUs is limited, and the JPEG standard itself allows for fully integer implementation. Floating-point schemes waste computing power and increase power consumption.

[0005] In summary, existing technologies either rely on dedicated hardware, focus on algorithm accuracy / structural innovation, or lack in-depth optimization of the data transfer bottleneck required for JPEG decoding on general-purpose GPUs. They have not yet provided a decoding method that can achieve efficient, compatible, and stable acceleration of standard JPEG decoding on a domestic CPU+GPU platform without a dedicated JPEG codec unit. Summary of the Invention

[0006] Based on the above analysis, the embodiments of the present invention aim to provide a JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs, in order to solve the technical problems of excessive PCIe transmission overhead and low floating-point operation efficiency caused by frequent small-batch data copying on domestic CPU+GPU heterogeneous platforms that only support OpenCL general-purpose computing in the absence of dedicated JPEG encoding and decoding hardware.

[0007] The objective of this invention is mainly achieved through the following technical solutions: This invention provides a batch processing and double-buffered pipeline scheduling method for JPEG decoding on general-purpose GPUs, comprising the following steps: Step S1: Obtain the bitstream data of the JPEG image to be decoded. The CPU performs entropy decoding and dequantization on the bitstream data to obtain multiple DCT coefficient blocks. If the size of the JPEG image to be decoded is not less than a preset size threshold, the DCT coefficient blocks are cached in a circular buffer located in the host memory. Step S2: Continuously monitor the status of the circular buffer. When the preset batch triggering conditions are met and the GPU is detected to be available, copy the DCT coefficient blocks in the circular buffer to the GPU memory in batches. Step S3: Based on the DCT coefficient block in the GPU memory, the GPU performs integer IDCT transformation and color space conversion calculations to obtain the original image corresponding to the JPEG image to be decoded, and then reads the original image back to the host memory. The process employs a double-buffered pipeline for batch copying of DCT coefficient blocks, integer IDCT transformation, color space conversion calculation, and reading back of the original image.

[0008] Furthermore, a double-buffered pipeline is employed for batch copying of DCT coefficient blocks, integer IDCT transformation, color space conversion calculation, and readback of the original image, including: Create a first OpenCL command queue and a second OpenCL command queue; the first OpenCL command queue is used for batch copying from the host memory's circular buffer to the GPU's video memory, and the second OpenCL command queue is used for performing integer IDCT transformation and color space conversion calculations and reading back the original image; The first OpenCL command queue is used to perform a batch copy operation of the first batch of DCT coefficient blocks from the circular buffer in the host memory to the GPU memory; after the first batch of copying is completed, the second OpenCL command queue performs integer IDCT transformation and color space conversion calculation operations for the first batch of DCT coefficient blocks. For batch serial number , Execute the first OpenCL command queue Batch copy operation of DCT coefficient blocks from the host memory's circular buffer to the GPU's video memory; In the After the batch copy is completed, the second OpenCL command queue is used to execute the first... Integer IDCT transform and color space conversion calculation operations for batch DCT coefficient blocks; In the After the batch calculation is completed, the second OpenCL command queue is used to execute the first... The corresponding raw image data is read back from GPU memory to host memory. For the first Batch copy operation, for the first Batch calculation operations, for the first Batch readback operations are executed in parallel.

[0009] Furthermore, the preset batch triggering conditions include any one or a combination of the following: Condition 1: The number of DCT coefficient blocks cached in the circular buffer reaches or exceeds the preset number of blocks; Condition 2: The waiting time since the last copy of DCT coefficient block data from the circular buffer to the GPU has reached or exceeded the preset time.

[0010] Furthermore, detecting that the size of the JPEG image to be decoded is not less than a preset size threshold includes: The CPU is used to obtain the SOF of the JPEG image to be decoded, and the width and height of the JPEG image to be decoded are obtained based on the SOF. If the width and height of the JPEG image to be decoded are both less than the preset size threshold, then the size of the JPEG image to be decoded is determined to be less than the preset size threshold; the CPU is directly used to perform integer IDCT transformation and color space conversion calculation operations on the DCT coefficient block. Otherwise, if the size of the JPEG image to be decoded is determined to be no greater than a preset size threshold, the DCT coefficient block is cached in a circular buffer located in the host memory.

[0011] Furthermore, the detection GPU is available, including: Detect the GPU's OpenCL version to determine if it supports OpenCL 1.2 or higher; Attempt to create a GPU context; If the GPU's OpenCL version supports OpenCL 1.2 or higher, and the GPU context is successfully created, the GPU is considered usable; otherwise, the GPU is considered unusable.

[0012] Furthermore, both the first OpenCL command queue and the second OpenCL command queue are in-order command queues.

[0013] Furthermore, the circular buffer employs a lock-free queue mechanism to achieve thread safety.

[0014] Furthermore, the integer IDCT transform employs fixed-point integer arithmetic, and the DCT coefficient block obtained after dequantization on the CPU is stored in the circular buffer in 16-bit signed integer format.

[0015] Furthermore, the GPU is a general-purpose GPU and does not contain a dedicated JPEG hardware encoding / decoding unit; The CPU and GPU communicate in full-duplex mode via the PCIe bus.

[0016] Further, the bitstream data of the JPEG image to be decoded is obtained, including: The JPEG image to be decoded is read into the host memory; Using JPEG decoding library functions, the marked segments of the JPEG image to be decoded are parsed, and the scan start mark SOS is located. From the SOS marker to the end-of-image marker (EOI), continuous binary data is extracted as the bitstream data of the JPEG image to be decoded.

[0017] Compared with the prior art, the present invention can achieve at least one of the following beneficial effects: 1. Existing technologies, when accelerating decoding on general-purpose GPU platforms without dedicated JPEG hardware, often result in excessive PCIe transmission overhead due to frequent small-batch data copying, leading to a performance inversion where "hardware acceleration is actually slower." This invention constructs a circular buffer in host memory and designs a dual-trigger mechanism based on a quantity threshold (≥128 blocks) and a time threshold (≤5ms) to aggregate scattered DCT coefficient blocks into large data blocks and submit them to GPU memory in batches. This allows the PCIe bus to transmit over 256KB of data at a time, effectively distributing transmission overhead and improving bandwidth utilization. Real-world testing shows that in high-resolution image batch decoding scenarios, this invention improves decoding throughput by 3-5 times compared to traditional block-by-block transmission schemes, completely resolving the acceleration failure problem of JPEG decoding on general-purpose GPU platforms. It addresses the core pain point of "performance inversion" and significantly improves decoding throughput. 2. This invention employs a double-buffered pipeline for batch copying of DCT coefficient blocks, integer IDCT transformation, color space conversion calculations, and original image readback, achieving saturated utilization of CPU and GPU resources. Addressing the hardware limitation of domestic general-purpose GPUs not supporting out-of-order execution, this invention proposes a dual-command queue collaborative scheduling mechanism: creating two independent in-order OpenCL command queues, dedicated to data copying and computation readback respectively, and constructing a "copy-first" mechanism through event synchronization. Batch, calculation Batch, reread The design utilizes a "batch" pipeline architecture. It fully leverages the full-duplex nature of the PCIe bus, enabling uplink data transmission, GPU core computation, and downlink result readback to proceed in parallel. This improves the utilization of GPU computing units and simultaneously saturates the bidirectional bandwidth of the PCIe bus, achieving maximum utilization of both computing and communication resources. 3. This invention addresses the hardware and software compatibility challenges of domestic platforms by constructing a multi-layered runtime fallback mechanism: Before decoding begins, small images are filtered using an image size threshold (640×480 pixels). If an image is detected as too small, the system seamlessly switches to the CPU SIMD decoding path (such as libjpeg-turbo), ensuring uninterrupted decoding functionality. During the initialization phase, dual detection via OpenCL version lookup and context creation accurately identifies GPU unavailability anomalies. This design gives the solution the robustness of industrial-grade products, achieving a perfect balance between high-efficiency acceleration and reliable operation on domestic CPU+GPU heterogeneous platforms; the comprehensive robust design ensures stable availability in complex environments.

[0018] In this invention, the above-described technical solutions can be combined with each other to achieve more preferred combinations. Other features and advantages of this invention will be set forth in the following description, and some advantages may become apparent from the description or be learned by practicing the invention. The objects and other advantages of this invention can be realized and obtained from what is particularly pointed out in the description and drawings. Attached Figure Description

[0019] The accompanying drawings are for illustrative purposes only and are not intended to limit the invention. Throughout the drawings, the same reference numerals denote the same parts.

[0020] Figure 1 This is a flowchart of the JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs in an embodiment of the present invention; Figure 2 This is a flowchart of the JPEG image decoding process on a CPU+GPU heterogeneous platform in an embodiment of the present invention; Figure 3 This is a flowchart illustrating the copying, computation, and readback process using a double-buffered pipeline in an embodiment of the present invention. Figure 4 This is a schematic diagram of the physical components of a JPEG image file in an embodiment of the present invention. Detailed Implementation

[0021] Preferred embodiments of the present invention will now be described in detail with reference to the accompanying drawings, which form part of this application and are used together with the embodiments of the present invention to illustrate the principles of the present invention, but are not intended to limit the scope of the present invention.

[0022] To address the performance bottleneck caused by frequent small-batch data copying during JPEG decoding on general-purpose GPUs in existing technologies, this invention provides a batch processing and double-buffered pipeline scheduling method for JPEG decoding on general-purpose GPUs. This method does not rely on a specific IDCT algorithm implementation, focusing on task scheduling, data flow management, and heterogeneous resource coordination between CPU and GPU in the decoding process. It is suitable for scenarios where computationally intensive image processing tasks are offloaded from the CPU to general-purpose GPUs without dedicated hardware accelerators.

[0023] During JPEG encoding, image data undergoes the following process: Original image → Color space conversion / sampling → FDCT (forward DCT) → Quantization → Entropy coding (Huffman) → Compressed JPEG image file.

[0024] JPEG decoding is the reverse process of JPEG encoding: JPEG file → Entropy decoding (Huffman) → Inverse quantization → IDCT → Color space conversion → Original image.

[0025] The main process of JPEG decoding is: entropy decoding (Huffman decoding), inverse quantization, inverse discrete cosine transform (IDCT), and color space conversion. The technical solution of this invention mainly focuses on the optimization and improvement of the offloading and scheduling of existing IDCT transform and color space conversion on domestic general-purpose GPUs.

[0026] A specific embodiment of the present invention discloses a JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs, such as... Figure 1 , Figure 2 As shown, it includes the following steps: Step S1: Obtain the bitstream data of the JPEG image to be decoded. The CPU performs entropy decoding and inverse quantization on the bitstream data to obtain multiple DCT coefficient blocks. If the size of the JPEG image to be decoded is not less than a preset size threshold, the DCT coefficient blocks are cached in a circular buffer located in the host memory. Step S2: Continuously monitor the status of the circular buffer. When the preset batch triggering conditions are met and the GPU is detected to be available, copy the DCT coefficient blocks in the circular buffer to the GPU memory in batches. Step S3: Based on the DCT coefficient block in the GPU memory, the GPU performs integer IDCT transformation and color space conversion calculations to obtain the original image corresponding to the JPEG image to be decoded, and then reads the original image back to the host memory. The process employs a double-buffered pipeline for batch copying of DCT coefficient blocks, integer IDCT transformation, color space conversion calculation, and reading back of the original image.

[0027] Step S1 includes steps S11-S12.

[0028] Step S11: Obtain the bitstream data of the JPEG image to be decoded. The CPU performs entropy decoding and inverse quantization on the bitstream data to obtain multiple DCT coefficient blocks.

[0029] Obtain the bitstream data of the JPEG image to be decoded, including: The JPEG image to be decoded is read into the host memory; Using JPEG decoding library functions, the marked segments of the JPEG image to be decoded are parsed, and the scan start mark SOS is located. From the SOS marker to the end-of-image marker (EOI), continuous binary data is extracted as the bitstream data of the JPEG image to be decoded.

[0030] The central processing unit (CPU) calls the JPEG decoding library libjpeg to perform Huffman entropy decoding and inverse quantization on the streaming JPEG bitstream data, and obtains multiple discrete cosine transform (DCT) coefficient blocks with 8×8 pixels as the basic unit.

[0031] JPEG streams are 01 bitstreams.

[0032] The number of DCT coefficient blocks obtained after entropy decoding and inverse quantization of a JPEG image is calculated as follows: Formula (1) in, The number of color components. For the first The image width corresponding to each color component For the first The image height corresponding to each color component.

[0033] Example: A common 1920×1080 pixel color JPEG image, YUV, 4:2:0 sampling, DCT coefficient block calculation example, as shown in Table 1.

[0034] Table 1: Example of DCT coefficient block calculation for JPEG images

[0035] The total number of DCT coefficient blocks in this JPEG image is 32400+8100+8100=48600 8×8 pixel DCT coefficient blocks.

[0036] Step S12: If the size of the current JPEG image to be decoded is not less than (greater than or equal to) a preset size threshold, the DCT coefficient block is cached in a circular buffer located in the host memory.

[0037] At the start of the decoding task, during the initialization phase, the size attributes of the input JPEG image to be decoded and the availability of the GPU are determined. If the JPEG image size is smaller than a preset size threshold (such as 640×480 pixels) or the GPU is unavailable, the CPU is used directly for decoding instead of submitting the GPU for hardware-accelerated decoding.

[0038] The detection process includes ensuring that the size of the JPEG image to be decoded is not less than a preset size threshold, including: The CPU is used to obtain the SOF (Start of Frame) of the JPEG image to be decoded, and the width and height of the JPEG image to be decoded are obtained based on the SOF. If the width and height of the JPEG image to be decoded are both less than the preset size threshold, then the size of the JPEG image to be decoded is determined to be less than the preset size threshold; the CPU is directly used to perform integer IDCT transformation and color space conversion calculation operations on the DCT coefficient block. Otherwise, if the size of the JPEG image to be decoded is determined to be no greater than a preset size threshold, the DCT coefficient block is cached in a circular buffer located in the host memory.

[0039] For example, the CPU calls the `jpeg_read_header()` function of the JPEG decoding library libjpeg to parse the marked segments of the JPEG image to be decoded and obtain the SOF of the JPEG image to be decoded, such as... Figure 4 As shown, the width and height are obtained based on SOF.

[0040] If the size (width, height) of the JPEG image to be decoded is smaller than the preset size threshold and the GPU is unavailable, then it is not necessary to cache the DCT coefficient blocks generated by the CPU to the circular buffer, and seamlessly switch to the CPU SIMD decoding path (such as libjpeg-turbo). Otherwise, continue to try the GPU path to ensure that the decoding function is not interrupted.

[0041] For example, the preset size threshold is 640×480 pixels.

[0042] In traditional GPU acceleration solutions, the GPU is usually directly called to process each independent MCU (smallest coding unit), which will cause a fatal problem: an explosion of PCIe bus transaction count.

[0043] Since domestically produced GPUs lack dedicated JPEG image decoding hardware, this solution uses a software-based approach to overcome hardware limitations.

[0044] When a CPU processes a JPEG image file, it is a serial, block-by-block process; a GPU is a massively parallel computing process.

[0045] CPUs process data serially, continuously generating DCT coefficient blocks; GPUs perform parallel computing, requiring multiple DCT blocks to be acquired at once for efficient execution. Therefore, this invention establishes a cache between the CPU and GPU, creating a ring buffer in host memory (not the CPU's built-in cache). The DCT coefficient blocks generated by the CPU are written to this ring buffer. Circular buffer Host memory, which includes a circular buffer, is the parent set of physical memory managed and used by the CPU, i.e., the RAM of the computer host. The circular buffer is a fixed-size area specifically initialized and allocated in the host memory. It is a subset of the host memory and is used to cache DCT coefficient blocks.

[0046] While GPU acceleration is fast, it comes with fixed overhead costs.

[0047] ① Kernel startup overhead: Each time a decoding task is submitted to the GPU, the driver layer has a fixed delay, ranging from tens to hundreds of microseconds. ② PCIe transmission overhead: Even with batch processing, for extremely small images, the entire image may only require a few MCUs, and the data volume may not even fill the payload of a single PCIe transaction. ③ Context switching overhead: Calling the OpenCL API involves switching from user mode to kernel mode.

[0048] When a JPEG image is smaller than 640x480 pixels, the CPU's SIMD instruction set (such as SSE / AVX) can complete the decoding in milliseconds, while the GPU's fixed overhead may cause the total time to exceed that of the CPU. Setting a preset image size threshold ensures that only large images "worth" being processed by the GPU will take the accelerated decoding path.

[0049] The circular buffer is used to cache the DCT coefficient block obtained by the CPU after completing JPEG entropy decoding. It is created by the host in the host memory during the initialization process when JPEG decoding begins and destroyed after the JPEG decoding process is completed.

[0050] The circular buffer uses a lock-free queue mechanism to achieve thread safety.

[0051] The circular buffer is protected by a thread-safe mechanism (such as a lock-free queue mechanism) to ensure data consistency in a multi-threaded environment.

[0052] Entropy decoding is a computationally intensive task. If a mutex lock is used, when the buffer is in a boundary state (such as almost full or almost empty), lock contention will cause the CPU core to get stuck in a wait, resulting in huge context switching overhead.

[0053] Lock-free queues (based on CAS operations, i.e., Compare-and-swap atomic operations) can ensure that decoding threads and data copying threads can proceed without blocking at the most critical moments, which is crucial for maintaining the continuity of PCIe bus transmission.

[0054] For example, the size of the circular buffer is initialized to be no less than 256KB, which can buffer no less than 2048 8×8 pixel DCT coefficient blocks. The size of the circular buffer can accommodate no less than 8 copy submissions of DCT coefficient block data to the GPU.

[0055] Each 8×8 pixel DCT coefficient block, if stored as a 16-bit integer, has the following size: 8 × 8 × 2 = 128 bytes 256KB / 128 bytes = 2048 blocks.

[0056] 256KB is the size of many CPU L2 caches and is also the threshold that can trigger the MRRS (Maximum Read Request Size) limit in PCIe bus transmission.

[0057] The ring buffer is less than 256KB, which cannot fully utilize the bandwidth of the PCIe bus link; The design logic of a circular buffer to accommodate at least eight GPU-submitted DCT coefficient blocks is to provide a resilient buffer, absorbing speed fluctuations between the CPU and GPU. This means that when the GPU is temporarily unable to retrieve data due to computational demands (e.g., processing high-resolution image batches), the circular buffer can still continue to receive new DCT coefficient block data decoded by the CPU. This prevents the CPU from being blocked and waiting due to a full buffer, ensuring the smoothness of the entire decoding process.

[0058] On domestically produced platforms, hardware resources are prioritized, and software design must be meticulously planned, using circular buffers and the simplest data structures to solve complex problems.

[0059] In real-time decoding scenarios, if dynamic data structures (such as linked lists or dynamic arrays) are used, malloc / free operations may be triggered every time new DCT coefficient blocks arrive. These operations are unpredictable and may lead to: ① Memory fragmentation; ② Unpredictable delays (ranging from tens of microseconds to milliseconds); ③ In a multi-threaded environment, lock contention intensifies.

[0060] Advantages of a circular buffer: The circular buffer allocates at least 256KB of host memory during initialization, and dynamically allocates and releases memory throughout the decoding process in a circular manner. This is crucial for video decoding or real-time image processing that requires a stable frame rate.

[0061] The CPU decoding thread continuously stores DCT coefficient blocks into a circular buffer, while the monitoring thread (or submission thread) reads DCT coefficient blocks from the circular buffer. If a traditional queue is used, it typically requires a mutex lock for protection, which will lead to: ① The context switching overhead caused by lock contention; ② CPU or GPU is blocked and waiting.

[0062] Advantages of a circular buffer: A circular buffer, combined with atomic variables (such as read / write pointers), can perfectly implement a lock-free queue: Write pointer: The CPU atomically updates the circular buffer after completing the write operation. Read pointer: Data is copied to GPU memory in batches and then updated atomically; Full / empty checks: These are determined by pointer relationships and do not require lock protection.

[0063] Step S1 converts the JPEG bitstream to be decoded into DCT coefficient blocks that can be directly computed by the GPU. By judging the preset size threshold of the JPEG image, it realizes the intelligent switching between CPU local decoding and GPU accelerated decoding paths, ensuring low latency for decoding small-sized images and high efficiency for GPU acceleration of large-sized images. At the same time, by using the lock-free queue mechanism of the circular buffer, it avoids fragmentation and unpredictable latency caused by dynamic memory allocation, and provides stable and efficient data buffer support for the subsequent batch copying of DCT coefficient blocks and double-buffered pipeline scheduling. It effectively bridges the difference in capabilities between CPU serial processing and GPU parallel computing, laying the foundation for improving the overall decoding performance.

[0064] Step S2, specifically.

[0065] This step is the process of dynamically submitting the DCT coefficient block data of the circular buffer to the GPU memory in batches.

[0066] The status of the circular buffer is continuously monitored in real time. The system determines whether the amount of data or the data waiting time in the circular buffer meets preset batch triggering conditions, and also checks if the GPU is available. These two steps are logically ANDed and both are indispensable.

[0067] The preset batch triggering conditions include any one or a combination of the following: Condition 1: The number of DCT coefficient blocks cached in the circular buffer reaches or exceeds the preset number of blocks; Condition 2: The waiting time since the last copy of DCT coefficient block data from the circular buffer to the GPU has reached or exceeded the preset time.

[0068] For example, the preset number of blocks is 128, and the preset time is 5 milliseconds; in actual applications, this can be changed according to specific needs.

[0069] The detection of GPU availability includes: Detect the GPU's OpenCL version to determine if it supports OpenCL 1.2 or higher; Attempt to create a GPU context; If the GPU's OpenCL version supports OpenCL 1.2 or higher, and the GPU context is successfully created, the GPU is considered usable; otherwise, the GPU is considered unusable.

[0070] The availability of the GPU is determined by checking the OpenCL environment.

[0071] If (Condition 1 or Condition 2) and GPU is available is true, that is, the preset batch triggering condition is met and the GPU is detected to be available, then the OpenCL interface clEnqueueWriteBuffer function is used to trigger the task of batch copying the DCT coefficient block data in the circular buffer to the GPU memory.

[0072] The preset batch triggering conditions are a dual triggering logic of "quantity and time"; the quantity threshold ensures high efficiency, and the timeout threshold ensures real-time performance.

[0073] If either condition 1 or condition 2 is met, and the GPU is available, a task will be triggered to batch copy the DCT coefficient block data of the circular buffer to the GPU memory.

[0074] ①Condition 1: Threshold condition for the number of DCT coefficient blocks in the circular buffer. This refers to the number of DCT coefficient blocks currently backed up in the circular buffer, i.e., how many 8×8 pixel DCT coefficient blocks are currently cached in the circular buffer. When the number of DCT coefficient blocks in the circular buffer reaches a preset threshold for the number of DCT coefficient blocks. , ≥128 blocks, corresponding to approximately 64KB of data, is below the threshold. Each PCIe bus transfer has a large bandwidth overhead, and using the GPU for DCT coefficient block calculation has no significant effect; therefore, it is necessary to ensure the utilization rate of PCIe bus bandwidth.

[0075] 128 blocks × 128 bytes per block (16-bit DCT coefficient block) = 16384 bytes, which is 16KB; When the amount of data transmitted in a single transaction is less than 16KB, the overhead of the PCIe bus link layer and data link layer increases significantly. When the amount of data exceeds 16KB, the effective bandwidth utilization begins to approach saturation.

[0076] Only by accumulating enough DCT coefficient block data and initiating a "bulk transfer" can the expensive PCIe bus overhead be spread out by the massive amount of data, truly leveraging the advantages of DMA (Direct Memory Access).

[0077] ②Condition 2: Timeout threshold for submitting DCT coefficient blocks to GPU memory: The timer starts after the last submission of DCT coefficient blocks to GPU memory (i.e., the last execution of clEnqueueWriteBuffer). When the waiting time exceeds the preset time threshold... , Greater than or equal to 5 milliseconds, even if the number of DCT coefficient blocks does not reach It also forces the submission of DCT coefficient blocks in the circular buffer to the GPU memory to avoid excessive latency in the first frame and ensure real-time performance.

[0078] The human eye typically perceives image latency in the 100-millisecond range, but the latency of interactive systems (such as clicking to view the next image) needs to be kept even lower. If only a quantity threshold is relied upon, accumulating 128 blocks could take hundreds of milliseconds or even longer when processing small images or in scenarios with slow low-entropy decoding, resulting in a noticeable "lag" for the user.

[0079] This mechanism ensures that decoding can proceed in a timely manner even when there are insufficient DCT coefficient blocks to fill the circular buffer. Sacrificing some transmission efficiency for real-time decoding responsiveness is crucial, especially when decoding the first frame of the image.

[0080] The purpose of setting the time threshold is to avoid excessive latency in the first frame due to waiting for the number of DCT coefficient blocks in the circular buffer to reach the threshold of 128 blocks. This ensures that even if the number of DCT coefficient blocks in the circular buffer does not reach the preset number, the DCT coefficient block data can still be submitted to the GPU memory in a timely manner to start the GPU decoding of the first frame.

[0081] This mechanism ensures that the DCT coefficient block data size is large enough for each task submitted to the GPU for computation, maximizing PCIe bus transmission efficiency while also ensuring real-time performance.

[0082] clEnqueueWriteBuffer is a function interface in OpenCL used to write DCT coefficient block data from a circular buffer in host memory to a GPU device buffer.

[0083] The device buffer is a buffer located in the GPU's video memory.

[0084] When (Condition 1 or Condition 2) and GPU availability is true, data transfer on the PCIe bus is triggered. The following operations are performed: ① Locking and fetching: Fetch all currently available data from the circular buffer. A block of DCT coefficients is packaged into a GPU task batch; ② Call the API to call the OpenCL clEnqueueWriteBuffer command, which submits the data pointer and size of this batch to the first OpenCL command queue (copy queue).

[0085] ③ Status Reset: Reset the timer =0, and update the read pointer of the circular buffer, releasing the space occupied by the circular buffer of the DCT coefficient block that has been copied to the GPU cache.

[0086] Regardless of the threshold for trigger quantity and time, and provided the GPU is available, the PCIe bus transfer task for DCT coefficient block data is successfully initiated, preparing data for GPU decoding.

[0087] The availability of the GPU is checked as follows: Detection A: Call the API functions clGetPlatformIDs() and clGetDeviceInfo() in OpenCL to query the OpenCL version number of the GPU and determine whether it supports OpenCL 1.2 or higher.

[0088] The clGetPlatformIDs function is an API in OpenCL for obtaining platform information; The clGetDeviceInfo function is an API interface in OpenCL for obtaining GPU information.

[0089] The results of the examination are as follows: (1) Does it support OpenCL? (2) Does a GPU exist? (3) Does the GPU support the minimum OpenCL version 1.2? If the GPU cannot be found, it means that OpebCL is not installed or the GPU driver is abnormal. Directly revert to the CPU's SIMD decoding path (such as libjpeg-turbo).

[0090] Detection B: Context and resource creation.

[0091] The standard API function clCreateContext for OpenCL contexts attempts to initialize the OpenCL GPU context.

[0092] If the driver is not installed or the context initialization fails and the clCreateContext function returns CL_DEVICE_NOT_AVAILABLE or CL_INVALID_DEVICE, the CPU directly calls the IDCT function in the libjpeg library to avoid performance degradation caused by fixed overhead for small images.

[0093] During context creation, an attempt is made to establish a connection with the GPU and allocate a basic resource handle.

[0094] If the result is CL_DEVICE_NOT_AVAILABLE, the GPU device is currently unavailable, possibly being exclusively used by another process, or is in a hibernation / suspended state. If CL_INVALID_DEVICE is returned, the GPU is invalid.

[0095] Once the fallback condition is triggered, the system switches from the GPU decoding path to the CPU path.

[0096] When there are compatibility issues with the GPU OpenCL driver, hardware resources are occupied by other processes, or the GPU device itself is not ready, the system can fall back to the CPU decoding solution to ensure that the JPEG decoding function is always available and that the application does not crash or lag due to GPU abnormalities.

[0097] Through the above technical solution, the present invention significantly reduces the data transfer frequency between the host CPU and the device GPU, hides the transmission delay, and achieves stable JPEG decoding acceleration on a domestic platform without dedicated hardware, thus solving the core pain point of "acceleration is actually slower" in the prior art.

[0098] Step S2 enables efficient, real-time batch transfer scheduling of DCT coefficient blocks from the host memory's circular buffer to the GPU's video memory. Through a dual triggering logic of "quantity threshold ensuring bandwidth efficiency + time threshold ensuring real-time response," it ensures that the amount of data transferred in a single PCIe bus is large enough to amortize link overhead and maximize the advantages of DMA transfer, while avoiding latency and stuttering in the first frame or real-time scenarios due to insufficient data accumulation. Simultaneously, GPU availability status detection ensures the feasibility of task submission. Through operations such as batch packaging, submission to the OpenCL copy queue, and status reset, it provides a continuous and stable data source input for subsequent parallel computing by the GPU. This effectively connects the entropy decoding output on the CPU side with the parallel processing stage on the GPU side, making it a crucial intermediate step for achieving CPU-GPU collaboration and improving overall JPEG decoding throughput and real-time performance.

[0099] Step S3, specifically.

[0100] The process of copying DCT coefficient block data from the circular buffer to GPU memory in batches, and performing IDCT transformation calculations and color space conversions on the GPU, requires serial processing. This is because the OpenCL of domestic general-purpose GPUs does not support out-of-order command queues. As a result, the copying of DCT coefficient block data from the circular buffer to GPU memory, the calculation of integer IDCT transformations and color space conversions, and the reading back of the original image decoded by the GPU to host memory all need to be processed serially.

[0101] The out-of-order command queue is an out-of-order queue for execution. Domestic general-purpose GPUs require tasks submitted to the queue to be executed in a specific order, and cannot dynamically adjust the execution order based on hardware resources or output tasks that have no dependencies in parallel.

[0102] In existing technologies, a queue is created, and all commands (batch copy → computation → readback) must be executed serially in strict accordance with the order of enqueueing. The next task will never start until the previous task is completed, which will cause the PCIe bus and GPU computing units to be idle and waiting.

[0103] The GPU is a general-purpose GPU and does not contain a dedicated JPEG hardware encoding / decoding unit; The CPU and GPU communicate in full-duplex mode via the PCIe bus.

[0104] Create two in-order queues; during the initialization phase, create two independent in-order OpenCL command queues and assign them completely different responsibilities.

[0105] A double-buffered pipeline is used for batch copying of DCT coefficient blocks, integer IDCT transformation, color space conversion calculation, and readback of the original image, including: Create a first OpenCL command queue and a second OpenCL command queue; the first OpenCL command queue is used for batch copying from the host memory's circular buffer to the GPU's video memory, and the second OpenCL command queue is used for performing integer IDCT transformation and color space conversion calculations and reading back the original image; The first OpenCL command queue is used to perform a batch copy operation of the first batch of DCT coefficient blocks from the circular buffer in the host memory to the GPU memory; after the first batch of copying is completed, the second OpenCL command queue performs integer IDCT transformation and color space conversion calculation operations for the first batch of DCT coefficient blocks. For batch serial number , Execute the first OpenCL command queue Batch copy operation of DCT coefficient blocks from the host memory's circular buffer to the GPU's video memory; In the After the batch copy is completed, the second OpenCL command queue is used to execute the first... Integer IDCT transform and color space conversion calculation operations for batch DCT coefficient blocks; In the After the batch calculation is completed, the second OpenCL command queue is used to execute the first... The corresponding raw image data is read back from GPU memory to host memory. For the first Batch copy operation, for the first Batch calculation operations, for the first Batch readback operations are executed in parallel.

[0106] This invention creates two independent in-order OpenCL command queues to solve the problem of low serial efficiency caused by the inability of general-purpose GPUs to support out-of-order execution.

[0107] In-order OpenCL command queues are a type of command queue in OpenCL (Open Computing Language). Their core characteristic is that commands submitted to this queue (such as data copying, kernel execution, or reading back decoded raw image data) must be executed serially in the strict order of submission. That is, the next command can only begin execution after the previous one is completed. It is impossible to dynamically adjust the execution order based on hardware resources or to process unrelated tasks in parallel. This queue type is the opposite of out-of-order command queues.

[0108] Both the first OpenCL command queue and the second OpenCL command queue are in-order command queues.

[0109] During image decoding initialization, a first OpenCL command queue and a second OpenCL command queue are created; The first OpenCL command queue is a copy queue (copy_queue); The second OpenCL command queue is the compute queue (compute_queue).

[0110] The copy queue is used to execute the clEnqueueWriteBuffer function, which copies DCT coefficient blocks in batches from the host memory circular buffer to the GPU video memory buffer.

[0111] The clEnqueueWriteBuffer function is an OpenCL function interface used to write data from host memory to a device buffer.

[0112] The computation queue is used to execute the clEnqueueNDRangeKernel function and the clEnqueueReadBuffer.

[0113] The clEnqueueNDRangeKernel function is used to submit OpenCL kernels to the command queue for execution. Here, it mainly submits the kernels for IDCT and color conversion calculations.

[0114] The clEnqueueReadBuffer function is an OpenCL function used to read decoded raw image data from GPU memory into host memory.

[0115] By assigning commands of different natures to different queues, the limitations of existing single-queue technology are broken. The two queues are independent at the GPU driver level, proceeding in parallel without blocking each other.

[0116] like Figure 2 As shown, the copy queue and the compute queue are synchronized through OpenCL's event wait list mechanism. When the preset batch triggering conditions in step S2 are met and the GPU is available, the copy queue performs a data copy operation from the circular buffer in the host memory to the DCT coefficient block in the GPU memory. After the copy is completed, the write operation of the copy queue is completed. Once the compute queue detects that the write operation of the current copy queue has been completed, it starts the kernel compute (kernel function compute) of the compute queue.

[0117] The next batch of copy queues can be initiated in parallel during the computation of the current batch, thus forming a "copy queue". →Calculation queue → Calculate queue back "The assembly line".

[0118] copy It uses a copy queue to store the next batch The DCT coefficient block is copied from the host memory's circular buffer to the GPU's video memory; Simultaneously calculate the queue for the current batch. DCT coefficient block calculation; Use the computation queue to process the previous batch The calculation results are read back from the GPU memory to the host memory.

[0119] The GPU is working at full capacity to calculate the first... At the same time as the approval, the first The batch-decoded raw image is read back from GPU memory to host memory; the PCIe bus is simultaneously processing the first... Batch DCT coefficient block data is copied from the host memory's circular buffer to the GPU's video memory.

[0120] By using pipeline overlap, the waiting time of traditional "copy → compute → readback" serial execution is avoided, making full use of the parallelism of the GPU and keeping the GPU and PCIe bus busy at all times, significantly improving the JPEG decoding throughput on general-purpose GPU platforms.

[0121] For example: when the first During GPU computation, the first batch The first batch of data has been copied to the GPU memory, while the second batch... The decoded raw image results are read back to the host memory. The three stages proceed in parallel to maximize resource utilization.

[0122] The PCIe bus is a full-duplex communication channel, which means: Simultaneous bidirectional transmission: Data can be transmitted simultaneously in both the host → GPU (uplink) and GPU → host (downlink) directions; Independent passageways: Each direction has its own passageway, and they do not interfere with each other; Bandwidth aggregation: For example, PCIe 3.0x16 has a unidirectional bandwidth of about 16GB / s, and the total bandwidth can reach 32GB / s when transmitting in both directions simultaneously.

[0123] The working diagrams of the copy queue and the compute queue are shown in Table 2.

[0124] Table 2: Schematic diagram of copy queue and compute queue operation.

[0125]

[0126] exist At time T3, a stable state is reached. PCIe uplink continues to be busy, copying the next batch of DCT coefficient blocks; PCIe downlink continues to be busy, reading back the original image results decoded from the previous batch from GPU memory to host memory; GPU cores continue to be busy, calculating the IDCT transformation and color space conversion of the previous batch.

[0127] The three processes run in parallel without blocking each other, and all three resources operate at full capacity. As long as there are enough batches of DCT coefficient blocks, data is always being transmitted in both the uplink and downlink directions of the PCIe bus. While the GPU is computing, both directions of the bus are busy.

[0128] The integer IDCT transform uses fixed-point integer arithmetic. After dequantization is completed on the CPU, the resulting DCT coefficient block is stored in the circular buffer in 16-bit signed integer format.

[0129] This invention fully utilizes the full-duplex capability of the PCIe bus and achieves triple parallelism of uplink data transmission, GPU computing, and downlink data readback through a dual-queue pipeline design, enabling the bandwidth of the PCIe bus in both directions and the GPU computing resources to be fully utilized simultaneously.

[0130] The purpose of step S3 is to simulate an efficient asynchronous scheduler at the software level using two sequential queues and event synchronization, even if the hardware does not support out-of-order delivery. This allows the latency of PCIe transmission to be completely hidden in the GPU's computation process, so that the throughput of the decoding process is no longer limited by the round-trip time of the PCIe bus, but only by the GPU's computing speed and the upper limit of the PCIe bus bandwidth.

[0131] In summary, the JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs according to embodiments of the present invention has the following beneficial effects: 1. Existing technologies, when accelerating decoding on general-purpose GPU platforms without dedicated JPEG hardware, often result in excessive PCIe transmission overhead due to frequent small-batch data copying, leading to a performance inversion where "hardware acceleration is actually slower." This invention constructs a circular buffer in host memory and designs a dual-trigger mechanism based on a quantity threshold (≥128 blocks) and a time threshold (≤5ms) to aggregate scattered DCT coefficient blocks into large data blocks and submit them to GPU memory in batches. This allows the PCIe bus to transmit over 256KB of data at a time, effectively distributing transmission overhead and improving bandwidth utilization. Real-world testing shows that in high-resolution image batch decoding scenarios, this invention improves decoding throughput by 3-5 times compared to traditional block-by-block transmission schemes, completely resolving the acceleration failure problem of JPEG decoding on general-purpose GPU platforms. It addresses the core pain point of "performance inversion" and significantly improves decoding throughput. 2. This invention employs a double-buffered pipeline for batch copying of DCT coefficient blocks, integer IDCT transformation, color space conversion calculations, and original image readback, achieving saturated utilization of CPU and GPU resources. Addressing the hardware limitation of domestic general-purpose GPUs not supporting out-of-order execution, this invention proposes a dual-command queue collaborative scheduling mechanism: creating two independent in-order OpenCL command queues, dedicated to data copying and computation readback respectively, and constructing a "copy-first" mechanism through event synchronization. Batch, calculation Batch, reread The design utilizes a "batch" pipeline architecture. It fully leverages the full-duplex nature of the PCIe bus, enabling uplink data transmission, GPU core computation, and downlink result readback to proceed in parallel. This improves the utilization of GPU computing units and simultaneously saturates the bidirectional bandwidth of the PCIe bus, achieving maximum utilization of both computing and communication resources. 3. This invention addresses the hardware and software compatibility challenges of domestic platforms by constructing a multi-layered runtime fallback mechanism: Before decoding begins, small images are filtered using an image size threshold (640×480 pixels). If an image is detected as too small, the system seamlessly switches to the CPU SIMD decoding path (such as libjpeg-turbo), ensuring uninterrupted decoding functionality. During the initialization phase, dual detection via OpenCL version lookup and context creation accurately identifies GPU unavailability anomalies. This design gives the solution the robustness of industrial-grade products, achieving a perfect balance between high-efficiency acceleration and reliable operation on domestic CPU+GPU heterogeneous platforms; the comprehensive robust design ensures stable availability in complex environments.

[0132] Those skilled in the art will understand that all or part of the processes of the methods described in the above embodiments can be implemented by a computer program instructing related hardware, and the program can be stored in a computer-readable storage medium. The computer-readable storage medium may be a disk, optical disk, read-only memory, or random access memory, etc.

[0133] The above description is only a preferred embodiment of the present invention, but the scope of protection of the present invention is not limited thereto. Any changes or substitutions that can be easily conceived by those skilled in the art within the scope of the technology disclosed in the present invention should be included within the scope of protection of the present invention.

Claims

1. A batch processing and double-buffered pipeline scheduling method for JPEG decoding on general-purpose GPUs, characterized in that, Includes the following steps: Step S1: Obtain the bitstream data of the JPEG image to be decoded. The CPU performs entropy decoding and dequantization on the bitstream data to obtain multiple DCT coefficient blocks. If the size of the JPEG image to be decoded is not less than a preset size threshold, the DCT coefficient blocks are cached in a circular buffer located in the host memory. Step S2: Continuously monitor the status of the circular buffer. When the preset batch triggering conditions are met and the GPU is detected to be available, copy the DCT coefficient blocks in the circular buffer to the GPU memory in batches. Step S3: Based on the DCT coefficient block in the GPU memory, the GPU performs integer IDCT transformation and color space conversion calculations to obtain the original image corresponding to the JPEG image to be decoded, and then reads the original image back to the host memory. The process employs a double-buffered pipeline for batch copying of DCT coefficient blocks, integer IDCT transformation, color space conversion calculation, and reading back of the original image.

2. The JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs according to claim 1, characterized in that, A double-buffered pipeline is used for batch copying of DCT coefficient blocks, integer IDCT transformation, color space conversion calculation, and readback of the original image, including: Create a first OpenCL command queue and a second OpenCL command queue; the first OpenCL command queue is used for batch copying from the host memory's circular buffer to the GPU's video memory, and the second OpenCL command queue is used for performing integer IDCT transformation and color space conversion calculations and reading back the original image; The first OpenCL command queue is used to perform a batch copy operation of the first batch of DCT coefficient blocks from the circular buffer in the host memory to the GPU memory; after the first batch of copying is completed, the second OpenCL command queue performs integer IDCT transformation and color space conversion calculation operations for the first batch of DCT coefficient blocks. For batch serial number , Execute the first OpenCL command queue Batch copy operation of DCT coefficient blocks from the host memory's circular buffer to the GPU's video memory; In the After the batch copy is completed, the second OpenCL command queue is used to execute the first... Integer IDCT transform and color space conversion calculation operations for batch DCT coefficient blocks; In the After the batch calculation is completed, the second OpenCL command queue is used to execute the first... The corresponding raw image data is read back from GPU memory to host memory. For the first Batch copy operation, for the first Batch calculation operations, for the first Batch readback operations are executed in parallel.

3. The JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs according to claim 1, characterized in that, The preset batch triggering conditions include any one or a combination of the following: Condition 1: The number of DCT coefficient blocks cached in the circular buffer reaches or exceeds the preset number of blocks; Condition 2: The waiting time since the last copy of DCT coefficient block data from the circular buffer to the GPU has reached or exceeded the preset time.

4. The JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs according to claim 1, characterized in that, The detection process includes ensuring that the size of the JPEG image to be decoded is not less than a preset size threshold, including: The CPU is used to obtain the SOF of the JPEG image to be decoded, and the width and height of the JPEG image to be decoded are obtained based on the SOF. If the width and height of the JPEG image to be decoded are both less than the preset size threshold, then the size of the JPEG image to be decoded is determined to be less than the preset size threshold; the CPU is directly used to perform integer IDCT transformation and color space conversion calculation operations on the DCT coefficient block. Otherwise, if the size of the JPEG image to be decoded is determined to be no greater than a preset size threshold, the DCT coefficient block is cached in a circular buffer located in the host memory.

5. The JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs according to claim 1, characterized in that, The detection of GPU availability includes: Detect the GPU's OpenCL version to determine if it supports OpenCL 1.2 or higher; Attempt to create a GPU context; If the GPU's OpenCL version supports OpenCL 1.2 or higher, and the GPU context is successfully created, the GPU is considered usable; otherwise, the GPU is considered unusable.

6. The JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs according to claim 2, characterized in that, Both the first OpenCL command queue and the second OpenCL command queue are in-order command queues.

7. The JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs according to claim 1, characterized in that, The circular buffer uses a lock-free queue mechanism to achieve thread safety.

8. The JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs according to claim 7, characterized in that, The integer IDCT transform uses fixed-point integer arithmetic. After dequantization is completed on the CPU, the resulting DCT coefficient block is stored in the circular buffer in 16-bit signed integer format.

9. The JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs according to any one of claims 1-8, characterized in that, The GPU is a general-purpose GPU and does not contain a dedicated JPEG hardware encoding / decoding unit; The CPU and GPU communicate in full-duplex mode via the PCIe bus.

10. The JPEG decoding batch processing and double-buffered pipeline scheduling method for general-purpose GPUs according to any one of claims 1-8, characterized in that, Obtain the bitstream data of the JPEG image to be decoded, including: The JPEG image to be decoded is read into the host memory; Using JPEG decoding library functions, the marked segments of the JPEG image to be decoded are parsed, and the scan start mark SOS is located. From the SOS marker to the end-of-image marker (EOI), continuous binary data is extracted as the bitstream data of the JPEG image to be decoded.

Images