GPU flow scheduling performance measurement method and device

Abstract

The invention belongs to the field of GPU computing performance testing, and provides a GPU flow scheduling performance measurement method and device.The method comprises the steps that in a pre-established basic execution environment, N null-kernel functions are submitted to a reference flow in sequence, and the reference flow executes the received null-kernel functions and measures the reference execution duration; alternately submitting the N null kernel functions to the first test flow and the second test flow, introducing a cross-flow synchronization event, and measuring the total execution duration; and according to the reference execution duration and the total execution duration, calculating to obtain synchronization time consumption between the streams. According to the method, the influence of a tracking tool and the like on the performance of the measured object is eliminated, the method has obvious advantages in accuracy and stability, high-precision and low-interference GPU scheduling performance measurement is realized, and a quantitative basis is provided for performance optimization of a system during GPU driving and operation.

Description

Technical Field

[0001] This application belongs to the field of GPU computing performance testing and driver layer performance analysis, specifically, it relates to a GPU stream scheduling performance measurement method and apparatus. Background Technology

[0002] A GPU (Graphics Processing Unit) is a graphics processing chip capable of graphics rendering, mathematical calculations, geometric operations, and more. In GPU parallel computing, GPU streams are widely used to implement asynchronous task execution and parallel scheduling. Performance analysis during the scheduling process typically employs the following methods: (1) Event Timestamp Measurement Method: The time difference between two events is obtained through cudaEventRecord() and cudaEventElapsedTime(), which are used to calculate the kernel function execution time or data transfer time. cudaEventRecord() is used for event recording and time measurement, measuring the time consumed by GPU operations (such as kernel execution and data transfer), and can also serve as a marker for stream synchronization. cudaEventRecord() inserts an event into the specified stream, marking the GPU's execution progress at that moment. cudaEventElapsedTime() is used to calculate the time difference between two events, helping to analyze and optimize GPU kernel performance, and stores the result in a specified variable, usually in milliseconds.

[0003] (2) Driver-level tracing: Using performance analysis / testing tools such as NVIDIA Nsight, CUPTI (CUDA Profiling Tools Interface), and NVTX (NVIDIA Tools Extension), various timestamps are recorded at the driver level to analyze scheduling and synchronization time. The CUPTI tool can perform performance testing on the program, obtain performance indicators or understand the performance status, and provides some function interfaces for developers to call to implement performance testing tasks. The NVTX tool library is used to insert markers in the application, making it easier to more clearly display the program's execution flow and the time consumption of each part in the performance analysis tool.

[0004] (3) Macro statistical method: Run multiple rounds of benchmark tests and calculate the average execution time through standardized evaluation methods.

[0005] While the methods described above can measure kernel function execution time, they cannot effectively measure the synchronization and scheduling time between GPU streams. Especially in scenarios with multiple concurrent streams, event synchronization and driver layer task switching only incur microsecond-level overhead, and the time resolution and measurement granularity of existing methods are insufficient to accurately reflect this microsecond-level overhead. Summary of the Invention

[0006] This application provides a method and apparatus for measuring GPU stream scheduling performance, which can calculate microsecond-level synchronization or scheduling overhead and achieve high-precision, low-interference GPU scheduling performance measurement.

[0007] In a first aspect, embodiments of this application provide a method for measuring GPU stream scheduling performance, including: In a pre-created basic execution environment N empty kernel functions are submitted to the reference stream in sequence. The reference stream executes the received empty kernel functions and measures the execution time. N empty kernel functions are alternately submitted to the first test stream and the second test stream, and cross-stream synchronization events are introduced to measure the total execution time. Based on the reference execution time and the total execution time, the synchronization time between streams is calculated.

[0008] The pre-created basic execution environment includes: creating two independent streams, which are designated as the first test stream and the second test stream, respectively; the reference stream is any one of the default stream, the first test stream, and the second test stream.

[0009] The pre-created basic execution environment includes: creating two events, designated as the first event and the second event, respectively; the introduction of cross-stream synchronization events includes: establishing a dependency relationship between N empty kernel functions alternately submitted to the first test stream and the second test stream through the first event in the first test stream and the second event in the second test stream.

[0010] Wherein, the N empty kernel functions submitted to the reference stream have no dependencies, and the measurement of the reference execution time includes: Start the timer when the first empty kernel function is submitted; Execute the stream synchronization instruction to ensure that all empty kernel functions in the reference stream have completed execution; When all empty kernel functions in the reference stream have finished executing, the timer stops and the reference execution time is recorded.

[0011] The total execution time of the measurement includes: When the first empty kernel function is submitted, a timer is started; a stream synchronization instruction is executed to ensure that the empty kernel functions in the first test stream and the second test stream have completed execution; when the empty kernel functions in the first test stream and the second test stream have completed execution, the timer is stopped, and the total execution time is obtained.

[0012] The synchronization time between streams is calculated based on the reference execution time and the total execution time, including: calculating the average synchronization time SyncCost using the following formula:

[0013] Among them, T dual T represents the total execution time. single For reference, the execution time is [not specified].

[0014] This also includes: calculating the throughput metric TP based on the total execution time and the reference execution time. TP = (N - 1) × 10^6 / (T dual - T single ) The throughput metric TP is used to describe the average task processing capacity of the system under synchronous conditions.

[0015] Secondly, this application provides a GPU stream scheduling performance measurement device, comprising: The first measurement unit is used to sequentially submit N empty kernel functions to the reference stream in a pre-created basic execution environment, and the reference stream executes the received empty kernel functions to measure the reference execution time. The second measurement unit is used to alternately submit N empty kernel functions to the first test stream and the second test stream in a pre-created basic execution environment, introduce cross-stream synchronization events, and measure the total execution time. The calculation unit is used to calculate the synchronization time between streams based on the reference execution time and the total execution time.

[0016] Thirdly, this application provides a computer-readable storage medium having a computer program stored thereon that, when executed by a processor, implements the steps of any of the methods described above.

[0017] Fourthly, this application provides an electronic device, including a memory, a processor, and a computer program stored in the memory and executable on the processor, wherein the processor executes the program to implement the steps of any of the methods described above.

[0018] The GPU stream scheduling performance measurement method and apparatus of this application have the following beneficial effects: In this application, N empty kernel functions are sequentially submitted to a reference stream within a pre-created basic execution environment. The execution time of tasks in the reference stream is measured, i.e., the reference execution time. The N empty kernel functions are then alternately submitted to two test streams, introducing cross-stream synchronization events to establish dependencies between the empty kernel functions. The total execution time of tasks in the two test streams is measured. Finally, based on the reference execution time and the total execution time, the synchronization time between streams is calculated. This eliminates the impact of tracing tools on the performance of the tested object, exhibiting significant advantages in accuracy and stability. It achieves high-precision, low-interference GPU scheduling performance measurement, providing a quantitative basis for performance optimization of GPU drivers and runtime systems. Attached Figure Description

[0019] Figure 1 This is a schematic flowchart of the GPU stream scheduling performance measurement method in an embodiment of this application. Figure 1 ; Figure 2 This is a schematic flowchart of the GPU stream scheduling performance measurement method in an embodiment of this application. Figure 2 ; Figure 3 This is a schematic flowchart of the GPU stream scheduling performance measurement method in an embodiment of this application. Figure 3 ; Figure 4 This is a schematic diagram illustrating the principle of the GPU stream scheduling performance measurement method in this application embodiment; Figure 5 This is a schematic diagram of the structure of the GPU stream scheduling performance measurement device according to an embodiment of this application. Detailed Implementation

[0020] The present application will be further described below with reference to the accompanying drawings and embodiments.

[0021] In the following description, the terms "first" and "second" are used for descriptive purposes only and should not be construed as indicating or implying relative importance. The following description provides multiple embodiments of the invention, which can be substituted or combined with each other. Therefore, this application can also be considered to include all possible combinations of the same and / or different embodiments described. Thus, if one embodiment includes features A, B, and C, and another embodiment includes features B and D, then this application should also be considered to include embodiments containing one or more other possible combinations of features A, B, C, and D, even if such embodiments are not explicitly described in the following text.

[0022] like Figure 1As shown, the GPU stream scheduling performance measurement method of this application includes: S101, in a pre-created basic execution environment, N empty kernel functions are sequentially submitted to a reference stream, the reference stream executes the received empty kernel functions, and the reference execution time is measured; S103, in the pre-created basic execution environment, N empty kernel functions are alternately submitted to a first test stream and a second test stream, introducing cross-stream synchronization events, and the total execution time is measured; S105, based on the reference execution time and the total execution time, the synchronization time between streams is calculated. Each step is described below.

[0023] S101, In the pre-created basic execution environment, N empty kernel functions are submitted to the reference stream in sequence. The reference stream executes the received empty kernel functions and measures the reference execution time.

[0024] First, we will introduce the pre-created basic execution environment. The pre-created basic execution environment is the GPU's execution resources. Building the basic execution environment enables accurate observation of task scheduling behavior between GPU streams. Specifically, it includes creating two GPU streams and two event objects.

[0025] A GPU stream represents a series of asynchronous operations executed on the GPU. A stream corresponds to an execution queue, where operations are performed in the order determined by the host code. Streaming allows operations to be queued and ensures they are executed only after all previous operations have completed. The use of streams can significantly improve the performance of GPU programs, especially in scenarios requiring large amounts of data transfer and computation.

[0026] In some embodiments, for example, calling cudaStreamCreate() or the corresponding API (Application Programming Interface) creates two independent streams: Stream A and Stream B. Stream A and Stream B are used to carry different sequences of tasks. cudaStreamCreate() is an API in CUDA (Compute Unified Device Architecture) programming, used to create streams. CUDA is a general-purpose parallel computing architecture developed by NVIDIA, allowing programmers to leverage the parallel computing capabilities of NVIDIA GPUs to accelerate various computationally intensive applications.

[0027] In this step, the two GPU streams created are, for example, the first test stream and the second test stream, and the reference stream is any one of the default stream, the first test stream, and the second test stream. The purpose of creating two GPU streams is to distinguish between execution paths of "continuous execution within the same stream" and "synchronous execution across stream events" in subsequent experiments, thereby forming a comparable differential comparison.

[0028] In a GPU, events are used to record the queue state of a stream. An event marks a point in the stream's execution process, allowing checks whether an operation in the stream has reached that point. In some embodiments, the CUDA interface `cudaEventCreate()` is called to create a first and second event; alternatively, Event A and Event B can be created. The event mechanism is a fundamental synchronization primitive provided by the CUDA / MUSA runtime, enabling the establishment of task dependencies between different streams. In this application's scheme, events are used to precisely control the scheduling order of tasks between streams to simulate synchronization or waiting behavior within the GPU.

[0029] In some embodiments, this step further includes: allocating video memory space, specifically allocating a video memory buffer for flag variables via cudaMalloc(), for use in sharing state or delay control between kernels. Although this embodiment primarily uses empty kernel functions, this step is retained to ensure consistency with real-world workload scenarios.

[0030] N empty kernel functions are submitted sequentially to the reference stream. The reference stream executes the received empty kernel functions, and the process of measuring the execution time of the reference stream is called the single-stream testing phase, which is described below. The single-stream testing phase measures the execution time of continuously submitting N dependency-free tasks (empty kernel functions) under a single stream to obtain the baseline execution time T. single This benchmark execution time reflects the minimum command submission and execution latency of the GPU under no synchronization interference.

[0031] Kernel functions are code that runs on the GPU device, i.e., functions executed on the GPU. They are used to process large-scale data in parallel. Kernel functions are executed by multiple threads on the GPU, and the identifier of the current thread can be obtained in the kernel function.

[0032] like Figure 2As shown, in step S1011, N empty kernel functions are sequentially submitted to the same reference stream. This reference stream can be any one of the default stream, the first test stream, or the second test stream, and the N empty kernel functions have no dependency relationship. In step S1012, when the first empty kernel function is submitted, a timer is started. The timer can be, for example, a high-precision perf_counter timer, or other timers; this application does not limit the timer. In this application, the timer starts when the first task is submitted, and then the first empty kernel function (empty kernel task) is sent out. The time from sending the first empty kernel function to its transmission to the hardware is usually very short, on the order of microseconds (µs). The reference stream, the first test stream, and the second test stream will all have this time, which is eventually canceled out. Even if it cannot be guaranteed that the two are completely consistent, when the number of empty kernel functions N is very large, this part of the time can be ignored relative to the difference between the total execution time in the first test stream and the second test stream and the execution time in the reference stream.

[0033] After the tasks are submitted, a stream synchronization instruction is executed to ensure that all empty kernel functions in the reference stream have completed execution. The stream synchronization instruction is, for example, `cudaStreamSynchronize()`, which waits until the queue in the reference stream is empty, ensuring that all tasks have completed execution. In step S1013, when all empty kernel functions in the reference stream have completed execution, the timer stops, and the reference execution duration T is recorded. single .

[0034] In this step, during single-stream execution, all empty kernel functions (tasks) are submitted sequentially to the same command queue. The GPU hardware does not need to perform additional context switching or dependency resolution. Therefore, the reference execution time T is relatively short. single This is essentially equivalent to the cumulative value of the minimum empty kernel function scheduling time over N iterations. This step provides a reference performance for GPUs in asynchronous scenarios, which can be used for differential comparison with measurement results under multi-stream scheduling.

[0035] S103, in a pre-created basic execution environment, N empty kernel functions are alternately submitted to the first test stream and the second test stream, introducing cross-stream synchronization events and measuring the total execution time. In this step, the first and second test streams are, for example, the two streams created in step S101. An event records the state of the queues within the stream. An event marks a point in the stream's execution process; it can be treated as an operation inserted into the numerous operations within the stream. It checks whether the currently executing operations in the stream have reached that point; when they do, the tasks in the stream execute to the point recorded by the event. In this step, the events are, for example, created by `cudaEventCreate()`, the first and second events. Through the first event in the first test stream and the second event in the second test stream, dependencies are established between the N empty kernel functions alternately submitted to the first and second test streams. This step is the dual-stream testing phase. By introducing cross-stream synchronization events, it simulates multi-task collaboration or dependency scenarios and measures the total execution time T when synchronization dependencies are included. dual .

[0036] like Figure 3 and Figure 4 As shown, Figure 4 In the diagram, ① represents an empty kernel function, ② represents the scheduling time of the empty kernel function, and ③ represents the sum of the scheduling time and synchronization time of the empty kernel function. In step S1031, the first empty kernel function is submitted to the first test stream, the second empty kernel function is submitted to the second test stream, the third empty kernel function is submitted to the first test stream, the fourth empty kernel function is submitted to the second test stream, the fifth empty kernel function is submitted to the first test stream, and so on, until all N empty kernel functions are submitted. When the first empty kernel function is submitted, a timer is started.

[0037] In step S1032, a first event occurs in the first test stream, and a second event occurs in the second test stream. Through these first and second events, dependencies are established between the second and first empty kernel functions, the third and fourth empty kernel functions, and so on, establishing dependencies between the first and second test streams. After these dependencies are established, the execution of the second empty kernel function must wait for the completion of the first empty kernel function, the third empty kernel function must wait for the completion of the second empty kernel function, the fourth empty kernel function must wait for the completion of the third empty kernel function, the fifth empty kernel function must wait for the completion of the fourth empty kernel function, and so on, thus constructing a cross-stream dependency chain measurement structure.

[0038] In some embodiments, for each empty kernel function (except the first one), a dependency is established using cudaStreamWaitEvent(): the i-th commit of the first test stream waits for the second event on the second test stream; the (i+1)-th commit of the second test stream waits for the first event on the first test stream; and so on, N-1 synchronization relationships are established in an alternating manner. cudaStreamWaitEvent is an API in CUDA used to wait for an event in another stream to complete in one stream.

[0039] In step S1033, the first and second test streams are synchronized to ensure that all empty kernel functions in both streams complete execution. This is because the completion of an empty kernel function (task) does not guarantee its actual completion; therefore, synchronizing the first and second test streams ensures that the empty kernel functions (tasks) are truly completed. For example, `cudaStreamSynchronize()` is executed to synchronize the two test streams, waiting for all empty kernel functions (tasks) to complete. When all empty kernel functions in the first and second test streams have completed execution, the timer is stopped, and the total execution time T is obtained. dual .

[0040] Internally, inter-stream synchronization triggers mechanisms such as event dependency resolution, command sequence merging, and resource scheduling switching. This step, through a carefully designed alternating dependency structure, forces the GPU to perform multiple scheduling and synchronization operations between tasks, thus truly exposing the additional time overhead of synchronization. This step obtains the actual total execution time T, including the effects of synchronization. dual This is used to calculate the average loss caused by a single synchronization.

[0041] S107, calculate the synchronization time between streams based on the reference execution time and the total execution time.

[0042] This step includes calculating the average synchronization time (time overhead) using the following formula:

[0043] Among them, T dual T represents the total execution time. single For reference execution time, N is the number of empty kernel functions, and SyncCost is the average time cost introduced by one inter-stream synchronization or GPU scheduling. SyncCost can be used as an indicator to compare the performance of different GPU architectures, drivers or scheduling strategies.

[0044] In some embodiments, this application further includes: performing multiple tests to obtain multiple average synchronization time costs, calculating the mean, standard deviation, and maximum / minimum values ​​to obtain stable performance statistics. For example, using the Celero performance testing framework to perform multiple experiments: automatically repeating tests; collecting samples; calculating the mean, standard deviation, and maximum / minimum values, thereby obtaining stable and repeatable performance statistics. Celero is a C++ micro-benchmarking framework that enables the writing and running of performance tests in C++ projects.

[0045] In some embodiments, this application further includes: calculating the throughput metric TP based on the total execution time and the reference execution time. TP = (N - 1) × 10^6 / (T dual - T single ) The throughput metric TP (unit: s^-1) is used to describe the average task processing capacity of a system under synchronous conditions.

[0046] In this application, N empty kernel functions are sequentially submitted to a reference stream within a pre-created basic execution environment. The execution time of tasks in the reference stream is measured, i.e., the reference execution time. The N empty kernel functions are then alternately submitted to two test streams, introducing cross-stream synchronization events to establish dependencies between the empty kernel functions. The total execution time of tasks in the two test streams is measured. Finally, based on the reference execution time and the total execution time, the synchronization time between streams is calculated. This eliminates the impact of tracing tools on the performance of the tested object, exhibiting significant advantages in accuracy and stability. It achieves high-precision, low-interference GPU scheduling performance measurement, providing a quantitative basis for performance optimization of GPU drivers and runtime systems.

[0047] Tracing mechanisms can interfere with GPU execution timing, causing disruptions to the system and affecting the true performance of the object under test. Furthermore, microsecond-level events are often easily masked by sampling noise. This application constructs two sets of task sequences and compares the difference in their total execution time to indirectly estimate microsecond-level synchronization or scheduling overhead. This achieves accurate measurement of microsecond-level inter-stream synchronization time in GPUs without the aid of tracing tools.

[0048] This application uses an empty kernel function for measurement, which avoids computational load interference. The empty kernel function greatly reduces measurement interference. The measurement process does not change the original scheduling behavior of the GPU, eliminates the "observer effect", and eliminates measurement noise of traditional timing methods. This application is simple to implement, requiring only the standard CUDA API and no additional drivers or underlying modifications. Furthermore, it can be used for performance comparisons of different GPU models, driver versions, or scheduling strategies. It is applicable to heterogeneous GPU architectures such as CUDA and MUSA, possessing cross-platform portability and strong versatility, and can quantitatively compare the synchronous scheduling efficiency between different GPUs and driver versions. The table below compares the effects of existing technologies with the method of this application:

[0049] like Figure 5 As shown, this application also provides a GPU stream scheduling performance measurement device, including: a first measurement unit 201, used to sequentially submit N empty kernel functions to a reference stream in a pre-created basic execution environment, wherein the reference stream executes the received empty kernel functions, and measures the reference execution time; a second measurement unit 202, used to alternately submit N empty kernel functions to a first test stream and a second test stream in a pre-created basic execution environment, introduce cross-stream synchronization events, and measure the total execution time; and a calculation unit 203, used to calculate the synchronization time between streams based on the reference execution time and the total execution time.

[0050] In this application, the embodiments of the GPU stream scheduling performance measurement device and the GPU stream scheduling performance measurement method are basically similar. For relevant details, please refer to the description of the GPU stream scheduling performance measurement method.

[0051] This application also provides a computer-readable storage medium storing a computer program thereon, which, when executed by a processor, implements the steps of the GPU stream scheduling performance measurement method described above. The computer-readable storage medium may include, but is not limited to, any type of disk, including floppy disks, optical disks, DVDs, CD-ROMs, microdrives, as well as magneto-optical disks, ROMs, RAMs, EPROMs, EEPROMs, DRAMs, VRAMs, flash memory devices, magnetic cards or optical cards, nanosystems (including molecular memory ICs), or any type of medium or device suitable for storing instructions and / or data.

[0052] This application also provides an electronic device, including a memory, a processor, and a computer program stored in the memory and executable on the processor, wherein the processor executes the program to implement the steps of any of the methods described above.

[0053] The electronic devices described in this application include, for example, laptop computers, desktop computers, workstations, personal digital assistants, servers, blade servers, mainframe computers, and also smart terminals such as mobile phones and tablets. The electronic devices of this application include a processor, memory, input devices, and output devices. The processor, memory, input devices, and output devices can be connected via a bus or other means. The memory stores a computer program that can run on the processor, and the processor executes the program to implement the steps of the GPU stream scheduling performance measurement method described above. Input devices include, for example, touchscreens, keypads, mice, trackpads, touchpads, and joysticks. Output devices include display devices, which may include, but are not limited to, liquid crystal displays (LCDs), light-emitting diode (LED) displays, plasma displays, and touchscreens.

[0054] In the several embodiments provided in this application, it should be understood that the disclosed apparatus and methods can be implemented in other ways. The apparatus embodiments described above are merely illustrative. For example, the division of units is only a logical functional division, and in actual implementation, there may be other division methods, such as: multiple units or components may be combined, or integrated into another system, or some features may be ignored or not executed. In addition, the coupling, direct coupling, or communication connection between the various components shown or discussed may be through some interfaces, indirect coupling or communication connection of devices or units, and may be electrical, mechanical, or other forms.

[0055] In the various embodiments of the present invention, all functional units can be integrated into one processing unit, or each unit can be a separate unit, or two or more units can be integrated into one unit; the integrated unit can be implemented in hardware or in the form of hardware plus software functional units.

[0056] The above description is merely a preferred embodiment of the present invention and is not intended to limit the invention. Various modifications and variations can be made to the present invention by those skilled in the art. Any modifications, equivalent substitutions, improvements, etc., made within the spirit and principles of the present invention should be included within the scope of protection of the present invention.

Claims

1. A method for measuring GPU stream scheduling performance, characterized in that, include: In a pre-created basic execution environment N empty kernel functions are submitted to the reference stream in sequence. The reference stream executes the received empty kernel functions and measures the execution time. N empty kernel functions are alternately submitted to the first test stream and the second test stream, and cross-stream synchronization events are introduced to measure the total execution time. Based on the reference execution time and the total execution time, the synchronization time between streams is calculated.

2. The GPU stream scheduling performance measurement method according to claim 1, characterized in that, The pre-created basic execution environment includes: creating two independent streams, which are respectively used as the first test stream and the second test stream; the reference stream is any one of the default stream, the first test stream, and the second test stream.

3. The GPU stream scheduling performance measurement method according to claim 1 or 2, characterized in that, The pre-created basic execution environment includes: creating two events, designated as the first event and the second event, respectively. The introduction of cross-stream synchronization events includes: establishing a dependency between N empty kernel functions alternately submitted to the first test stream and the second test stream through the first event in the first test stream and the second event in the second test stream.

4. The GPU stream scheduling performance measurement method according to claim 1 or 2, characterized in that, The N empty kernel functions submitted to the reference stream are independent of each other, and the measurement of the reference execution time includes: Start the timer when the first empty kernel function is submitted; Execute the stream synchronization instruction to ensure that all empty kernel functions in the reference stream have completed execution; When all empty kernel functions in the reference stream have finished executing, the timer stops and the reference execution time is recorded.

5. The GPU stream scheduling performance measurement method according to claim 3, characterized in that, The total execution time of the measurement includes: When the first empty kernel function is submitted, a timer is started; a stream synchronization instruction is executed to ensure that the empty kernel functions in the first test stream and the second test stream have completed execution; when the empty kernel functions in the first test stream and the second test stream have completed execution, the timer is stopped, and the total execution time is obtained.

6. The GPU stream scheduling performance measurement method according to claim 1 or 2, characterized in that, Based on the reference execution time and total execution time, the synchronization time between streams is calculated, including: calculating the average synchronization time SyncCost using the following formula: Among them, T dual T represents the total execution time. single For reference, the execution time is [not specified].

7. The GPU stream scheduling performance measurement method according to claim 1 or 2, characterized in that, Also includes: Calculate the throughput metric TP based on the total execution time and the reference execution time: TP = (N - 1) × 10^6 / (T dual - T single ) The throughput metric TP is used to describe the average task processing capacity of the system under synchronous conditions.

8. A GPU stream scheduling performance measurement device, characterized in that, include: The first measurement unit is used to sequentially submit N empty kernel functions to the reference stream in a pre-created basic execution environment, and the reference stream executes the received empty kernel functions to measure the reference execution time. The second measurement unit is used to alternately submit N empty kernel functions to the first test stream and the second test stream in a pre-created basic execution environment, introduce cross-stream synchronization events, and measure the total execution time. The calculation unit is used to calculate the synchronization time between streams based on the reference execution time and the total execution time.

9. A computer-readable storage medium having a computer program stored thereon, characterized in that, When executed by a processor, the program implements the steps of the method described in any one of claims 1-7.

10. An electronic device comprising a memory, a processor, and a computer program stored in the memory and executable on the processor, characterized in that, When the processor executes the program, it implements the steps of the method according to any one of claims 1-7.

Images