A GPU+CPU architecture radar signal processing implementation method based on grid stride loop

Abstract

The present invention discloses a GPU+CPU architecture radar signal processing implementation method based on a grid stride loop. The method comprises: connecting a CPU to a GPU via a PCIE data transmission bus for data exchange and complex logical operations; leveraging the GPU's high-performance computing capabilities, high concurrency, and high access speed to place data requiring a large number of threads for parallel computation on the GPU, thereby increasing computational speed; establishing a cache mechanism using GPU on-chip shared memory to store echo data frequently accessed by threads; and reusing threads in a grid stride loop to access shared memory data based on fine-grained thread indexing, minimizing data access time and thread creation costs, improving computational speed and efficiency, and rapidly implementing radar signal processing. This method has low development costs, simple hardware platform construction, and thread reuse, reducing data access time and thread activation and destruction costs, thereby increasing data computation speed and rapidly implementing radar signal processing.

Description

Technical Field

[0001] The present invention relates to the field of radar technology, and in particular to a method for rapidly processing radar signal data. The method is a method for implementing radar signal processing in a GPU+CPU architecture based on a grid stride loop. Technical Background

[0002] With the continuous development of information technology, such as drones, the sky environment is becoming increasingly complex. Hundreds or even thousands of drones may be present in the sky, along with a large number of false targets. Furthermore, various natural factors and human interference are increasing the computational complexity of radar signal processing systems, placing increasing demands on real-time data processing. Therefore, flexible, high-performance, and rapid radar signal processing based on software-based, universal capabilities is a key direction for future radar signal processing systems. Early radars used field-programmable gate arrays (FPGAs) and dedicated digital signal processors (DSPs) to build systems and perform radar signal processing tasks. However, these large-scale computations placed high demands on hardware, resulting in a complex hardware platform, limited flexibility, and poor maintainability. With the advancement of technology, Nvidia proposed the GPU-based CUDA (Compute Unified Device Architecture, or CUDA) architecture, which gave rise to a heterogeneous GPU + CPU architecture. The CPU handles complex logical operations, while data requiring a large number of threads for parallel computation is stored in the GPU's global memory. When the GPU performs parallel computations, kernel functions dispatch data cached in global memory to complete the computation. As radar detection ranges increase and the number of channels increases, the amount of data increases. When the GPU calculates large amounts of echo data, each echo data point is accessed multiple times. Frequent accesses to the GPU's global memory data cause significant latency, resulting in slower computation speeds. Summary of the Invention

[0003] The present invention aims to address the shortcomings of existing technologies and provide a GPU+CPU architecture radar signal processing implementation method based on grid stride loops. This method can improve computing speed and efficiency and quickly implement radar signal processing.

[0004] The technical solution for achieving the purpose of the present invention is:

[0005] A GPU+CPU architecture radar signal processing implementation method based on grid stride loop includes the following steps:

[0006] 1) Initialize radar parameters on the CPU host, open up memory and video memory space, and receive radar echo signal S echo (m, n), m = 1, 2, ... M, n = 1, 2, ... N, M represents the total number of radar echoes, and N represents the total number of sampling points for each echo;

[0007] 2) The CPU transmits the radar echo signal S through the PCIE data transmission bus echo (m, n) data is copied from the CPU host to the GPU device and stored in the GPU global memory;

[0008] 3) Radar echo signal S echo The (m, n) data is pulse compressed using a one-dimensional grid-stride cycle to obtain the pulse compression result S p (m, n), including:

[0009] 3-1) The CPU host side flips the original linear frequency modulation signal left and right, and performs conjugate complex operation on the flipped result to obtain the pulse compression coefficient S c (k), k = 1, 2, ... K, K represents the total number of pulse compression coefficients and distributes the pulse compression result S p (m, n) video memory space and thread space are opened to start the pulse compression processing kernel function;

[0010] 3-2) The echo signal S echo (m, n) data is vectorized to obtain the vectorized echo signal S echo (t), t = 1, 2, ... M*N, M*N represents the total number of radar echo samples;

[0011] 3-3) The GPU device opens a shared memory space in the pulse compression processing kernel function according to the number of threads in the block, and sequentially sends the radar echo signal S echo (t) Data is cached from GPU global memory to shared memory;

[0012] 3-4) Using a one-dimensional grid stride Grid-Stride loop, with the number of threads in a grid as the loop step, the one-dimensional grid threads are reused to form a thread fine-grained index, which corresponds one-to-one to the echo signal S after index vectorization echo (t), the thread index expression is:

[0013] G idx =S echo (idx+i*gridDim.x*blockDim.x),

[0014] idx=threadIdx.x+blockIdx.x*blockDim.x,

[0015] Among them, G idx Represents a grid thread, idx is the thread index number within the grid, i = 0, 1, ... I, I represents the total number of grid stride cycles, threadIdx.x represents the thread number within the thread block, and its value is threadIdx.x∈[0, blockDim.x), blockIdx.x represents the index of the thread block, and its value is blockIdx.x∈[0, gridDim.x), gridDim.x represents the dimension of the grid in the thread space, blockDim.x represents the dimension of the block in the thread space, and the values ​​of gridDim.x and blockDim.x are determined by the opened thread space;

[0016] 3-5) The radar echo signal S cached in the shared memory is accessed by the thread fine-grained index echo (t), calculate the radar echo signal S echo (t) for pulse compression coefficient S c (k) is discretely convolved to obtain the vectorized pulse compression result S p (t), the operation expression is:

[0017]

[0018] 3-6) The pulse compression result S p (t) Move from shared memory to global memory to provide source data for radar moving target display processing;

[0019] 4) Use shared memory to process the pulse compression result S p (t) Perform MTI processing and obtain the MTI processing result S MTI (m, n), specifically:

[0020] 4-1) The pulse compression result S p (t) The pulse compression result S is a two-dimensional matrix in the form of a cache to facilitate data cancellation, with the total number of M radar echoes as a row and each echo with N sampling points as a column. p (m, n);

[0021] 4-2) CPU allocation target display result S MTI (m, n) video memory space and thread space are opened, and the MTI processing kernel function is started;

[0022] 4-3) In the MTI processing kernel function, the GPU opens a shared memory space according to the number of threads in the block, and uses the control thread index method to sequentially convert the pulse compression result S in the form of a two-dimensional matrix into p (m, n) data is cached from GPU global memory to shared memory;

[0023] 4-4) Using a one-dimensional grid stride loop, with the number of threads in the grid as the loop step, thread reuse constitutes a thread fine-grained index, and one-to-one correspondence indexes the pulse compression result S in the form of a two-dimensional matrix p (m, n), the thread index expression is:

[0024] G idx =S p (idx,idy*gridDim.x*blockDim.x),

[0025] idx=threadIdx.x+blockIdx.x*blockDim.x,

[0026] idy=threadIdx.y+blockIdx.y*blockDim.y,

[0027] Among them G idx is the grid thread, idx is the thread index number in the grid, which also represents the x-axis index in the two-dimensional matrix, and idy is the y-axis index in the two-dimensional matrix, which also represents the number of grid stride cycles.

[0028] threadIdx.y, blockIdx.y, blockDim.x and threadIdx.x, blockIdx.x, blockDim.x have the same meaning, representing the thread number in the y and x directions, the thread block index number, and the thread dimension within the block respectively;

[0029] 4-5) The pulse compression result S cached in the shared memory is accessed by the thread fine-grained index p (m, n) and perform row subtraction to cancel the operation, and obtain the MTI processing result S MTI (m, n), the subtraction cancellation expression is: S MTI (m, n) = S p (m, n+1)-S p (m, n);

[0030] 4-6) MTI processing result S MTI (m, n) is moved from shared memory to GPU global memory;

[0031] 5) The pulse compression result S cached in a two-dimensional matrix form p (m, n) performs MTD processing and obtains the MTD processing result S MTD (m, n), the specific steps are as follows:

[0032] 5-1) CPU allocates MTD processing result S MTD (m, n) video memory space;

[0033] 5-2) Taking the length of a pulse signal as the sliding window size, perform M-point windowed FFT on M pulse pressure data in the same range Doppler dimension, and obtain the result S after FFT. MTD_FFT (m, n);

[0034] 5-3) The result S after FFT MTD_FFT (m, n) performs matrix transposition to obtain the MTD processing result S MTD (m, n);

[0035] 6) Using the data S cached in the GPU global memory MTD (m, n) performs constant false alarm detection according to the one-dimensional grid stride Grid-Stride cycle method, and screens the target to obtain the detection result S det (m, n), including:

[0036] 6-1) GPU initializes the parameters of the constant false alarm detection module and calculates the moving target detection result S MTD (m, n) performs a modulo operation to obtain the moving target detection result S after modulo operation. abs (m, n);

[0037] 6-2) The CPU allocates memory space for CFAR detection and creates thread space, and starts the CFAR detection processing kernel function;

[0038] 6-3) In the constant false alarm detection processing kernel function, the GPU opens a shared memory space according to the number of threads in the block, and uses the control thread index method to sequentially load the modulo moving target detection result S abs (m, n) is cached from GPU global memory to shared memory;

[0039] 6-4) Using a one-dimensional grid stride loop, with the number of threads in the grid as the loop step, it loops through the modulo-based moving target detection results S in the shared memory. abs (m, n) data is reduced and summed row by row;

[0040] 6-5) Use a sliding window to sum the reference unit data at the left and right ends of the reduced unit to be tested, and calculate the average value to obtain the noise background signal S noise (m, n);

[0041] 6-6) Based on the noise background data S noise (m, n) solves the decision threshold S T (m, n), compare S abs (m, n) and the decision threshold S T The size of (m, n): If S abs The value of (m, n) is less than S T (m, n), then directly determine S det (m, n) is a non-target detection point, and the detection result S is set det (m, n) = 0; if S abs The value of (m, n) is greater than or equal to S T (m, n), then determine S det (m, n) is the target echo detection point, and the detection result S is set det (m, n) = 1;

[0042] 6-7) The test result S det (m, n) is moved from the shared memory to the global memory, and then the detection results are copied to the CPU using the PCIE data transmission bus. Finally, the memory and video memory space are released to complete the radar signal processing.

[0043] The purpose of this technical solution is to solve the problems of high hardware requirements, complex hardware platform, poor maintainability, etc. when traditional solutions use FPGA and DSP to build radar signal processing systems. In addition, the long access time and slow computing speed when using global memory to cache large-scale data in the CPU+GPU architecture are also solved. This technical solution provides a GPU+CPU architecture based on grid stride loop for rapid radar signal processing.

[0044] This technical solution includes pulse compression, MTI, MTD, and constant false alarm detection processing. Under the GPU+CPU architecture, the access time of different storage types in the GPU memory structure is compared. The high-performance computing capability, high concurrency of the GPU and the high access speed of the shared memory shard memory are utilized to transfer a large amount of echo signal data received by the radar terminal from the CPU to the GPU global memory via the PCIE data transmission bus. When undergoing radar signal processing such as pulse compression, moving target display, moving target detection, and constant false alarm detection, a grid stride loop method is adopted, threads are reused, and shared memory data is accessed according to fine-grained thread indexing. This reduces data access time and thread time overhead during large-scale data operations, quickly realizes radar signal processing, and meets the real-time requirements of the radar.

[0045] This technical solution has the following advantages:

[0046] 1. Software-based and modular data processing with high flexibility and maintainability: On the GPU device side, pulse compression, moving target display, moving target detection, and constant false alarm detection in the radar signal system are all written using independent kernel functions, which are convenient and maintainable;

[0047] 2. Fast calculation speed and less time consumption: GPU has high parallelism, multi-threading, powerful arithmetic computing capabilities and high-bandwidth, low-latency shared memory cache mechanism, which improves the speed of data calculation, reduces calculation time, and can better meet the real-time requirements of radar;

[0048] 3. Grid stride loop, thread reuse, and maximized memory merging: This solves the problem of requiring far more threads than can be scheduled within a block on an SM when performing large-scale data operations, reducing the cost of thread activation and destruction.

[0049] 4. Low development cost, simple hardware platform construction, and reduced operator skill requirements for radar signal processing: This technical solution uses GPU+CPU as the hardware foundation, solving the problems of traditional FPGA and DSP hardware platforms for data processing, such as expensive hardware equipment, cumbersome platform environment construction, difficult hardware design and debugging, and low software program reusability.

[0050] 5. High target detection accuracy: The unified computing device architecture (CUDA) technology used in this technical solution contains library functions related to complex matrix operations such as FFT, IFFT, and multiplication, and has relatively powerful floating-point calculation capabilities.

[0051] This method can improve computing speed and efficiency and can quickly implement radar signal processing. BRIEF DESCRIPTION OF THE DRAWINGS

[0052] Figure 1 A schematic diagram of a method flow chart of an embodiment;

[0053] Figure 2 Schematic diagram of a one-dimensional grid strided cycle index radar echo signal in an embodiment;

[0054] Figure 3 Schematic diagram of pulse compression data in the form of a one-dimensional grid strided cyclic indexed two-dimensional matrix in an embodiment;

[0055] Figure 4 Flowchart of constant false alarm detection processing in the embodiment;

[0056] Figure 5 A diagram comparing the data processing time of the method of this example and the method of global memory cache in the embodiment;

[0057] Figure 6 This is a linear comparison chart of the data processing time using the method of this example and the method using global memory cache in the embodiment. DETAILED DESCRIPTION

[0058] The present invention will be further described below with reference to the accompanying drawings and embodiments, but the present invention is not limited thereto.

[0059] Example:

[0060] Reference Figure 1 A radar signal processing implementation method for GPU+CPU architecture based on grid stride loop includes the following steps:

[0061] 1) Initialize radar parameters on the CPU host, open up memory and video memory space, and receive radar echo signal S echo (m, n), m = 1, 2, ... M, n = 1, 2, ... N, M represents the total number of radar echoes, and N represents the total number of sampling points for each echo;

[0062] 2) The CPU transmits the radar echo signal S through the PCIE data transmission bus echo (m, n) data is copied from the CPU host to the GPU device and stored in the GPU global memory;

[0063] 3) Radar echo signal S echo The (m, n) data is pulse compressed using a one-dimensional grid-stride cycle to obtain the pulse compression result S p (m, n), including:

[0064] 3-1) The CPU host side flips the original linear frequency modulation signal left and right, performs conjugate complex operation on the flipped result, and obtains the pulse compression coefficient S by conjugate complex operation. c (k), k = 1, 2, ... K, K represents the total number of pulse compression coefficients and distributes the pulse compression result S p (m, n) video memory space and thread space are opened to start the pulse compression processing kernel function;

[0065] 3-2) The echo signal S echo (m, n) data is vectorized to obtain the vectorized echo signal S echo (t), t = 1, 2, ... M*N, M*N represents the total number of radar echo samples;

[0066] 3-3) The GPU device opens a shared memory space in the pulse compression processing kernel function according to the number of threads in the block, and sequentially sends the radar echo signal S echo(t) Data is cached from GPU global memory to shared memory;

[0067] 3-4) Using a one-dimensional grid stride Grid-Stride loop, with the number of threads in a grid as the loop step, the one-dimensional grid threads are reused to form a thread fine-grained index, which corresponds one-to-one to the echo signal S after index vectorization echo (t), the grid stride cycle principle is as follows Figure 2 As shown, the thread index expression is:

[0068] G idx =S echo (idx+i*gridDim.x*blockDim.x),

[0069] idx=threadIdx.x+blockIdx.x*blockDim.x,

[0070] Among them, G idx Represents a grid thread, idx is the thread index number within the grid, i = 0, 1, ... I, I represents the total number of grid stride cycles, threadIdx.x represents the thread number within the thread block, and its value is threadIdx.x∈[0, blockDim.x), blockIdx.x represents the index of the thread block, and its value is blockIdx.x∈[0, gridDim.x), gridDim.x represents the dimension of the grid in the thread space, blockDim.x represents the dimension of the block in the thread space, and the values ​​of gridDim.x and blockDim.x are determined by the opened thread space;

[0071] 3-5) The radar echo signal S cached in the shared memory is accessed by the thread fine-grained index echo (t), calculate the radar echo signal S echo (t) for pulse compression coefficient S c (k) is discretely convolved to obtain the vectorized pulse compression result S p (t), the operation expression is:

[0072]

[0073] 3-6) The pulse compression result S p (t) Move from shared memory to global memory to provide source data for radar moving target display processing;

[0074] 4) Use shared memory to process the pulse compression result Sp (t) Perform MTI processing and obtain the MTI processing result S MTI (m, n), specifically:

[0075] 4-1) The pulse compression result S p (t) The pulse compression result S is a two-dimensional matrix in the form of a cache to facilitate data cancellation, with the total number of M radar echoes as a row and each echo with N sampling points as a column. p (m, n);

[0076] 4-2) CPU allocation target display result S MTI (m, n) video memory space and thread space are opened, and the MTI processing kernel function is started;

[0077] 4-3) In the MTI processing kernel function, the GPU opens a shared memory space according to the number of threads in the block, and uses the control thread index method to sequentially convert the pulse compression result S in the form of a two-dimensional matrix into p (m, n) data is cached from GPU global memory to shared memory;

[0078] 4-4) Using a one-dimensional grid stride loop, with the number of threads in the grid as the loop step, thread reuse constitutes a thread fine-grained index, and one-to-one correspondence indexes the pulse compression result S in the form of a two-dimensional matrix p (m, n), such as Figure 3 As shown, the thread index expression is:

[0079] G idx =S p (idx,idy*gridDim.x*blockDim.x),

[0080] idx=threadIdx.x+blockIdx.x*blockDim.x,

[0081] idy=threadIdx.y+blockIdx.y*blockDim.y,

[0082] Among them G idx is the grid thread, idx is the thread index number in the grid, which also represents the x-axis index in the two-dimensional matrix, and idy is the y-axis index in the two-dimensional matrix, which also represents the number of grid stride cycles.

[0083] threadIdx.y, blockIdx.y, blockDim.x and threadIdx.x, blockIdx.x, blockDim.x have the same meaning, representing the thread number in the y and x directions, the thread block index number, and the thread dimension within the block respectively;

[0084] 4-5) The pulse compression result S cached in the shared memory is accessed by the thread fine-grained index p (m, n) and perform row subtraction to cancel the operation, and obtain the MTI processing result S MTI (m, n), the subtraction cancellation expression is: S MTI (m, n) = S p (m, n+1)-S p (m, n);

[0085] 4-6) MTI processing result S MTI (m, n) is moved from shared memory to GPU global memory;

[0086] 5) The pulse compression result S cached in a two-dimensional matrix form p (m, n) performs MTD processing and obtains the MTD processing result S MTD (m, n), the specific steps are as follows:

[0087] 5-1) CPU allocates MTD processing result S MTD (m, n) video memory space;

[0088] 5-2) Taking the length of a pulse signal as the sliding window size, perform M-point windowed FFT on M pulse pressure data in the same range Doppler dimension, and obtain the result S after FFT. MTD_FFT (m, n);

[0089] 5-3) The result S after FFT MTD_FFT (m, n) performs matrix transposition to obtain the MTD processing result S MTD (m, n);

[0090] 6) Using the data S cached in the GPU global memory MTD (m, n) performs constant false alarm detection according to the one-dimensional grid stride Grid-Stride cycle method, such as Figure 4 As shown, the target is screened to obtain the test result S det (m, n), including:

[0091] 6-1) GPU initializes the parameters of the constant false alarm detection module and calculates the moving target detection result S MTD (m, n) performs a modulo operation to obtain the moving target detection result S after modulo operation.abs (m, n);

[0092] 6-2) The CPU allocates memory space for CFAR detection and creates thread space, and starts the CFAR detection processing kernel function;

[0093] 6-3) In the constant false alarm detection processing kernel function, the GPU opens a shared memory space according to the number of threads in the block, and uses the control thread index method to sequentially load the modulo moving target detection result S abs (m, n) is cached from GPU global memory to shared memory;

[0094] 6-4) Using a one-dimensional grid stride loop, with the number of threads in the grid as the loop step, it loops through the modulo-based moving target detection results S in the shared memory. abs (m, n) data is reduced and summed row by row;

[0095] 6-5) Use a sliding window to sum the reference unit data at the left and right ends of the reduced unit to be tested, and calculate the average value to obtain the noise background signal S noise (m, n);

[0096] 6-6) Based on the noise background data S noise (m, n) solves the decision threshold S T (m, n), compare S abs (m, n) and the decision threshold S T The size of (m, n): If S abs The value of (m,) is less than S T (m, n), then directly determine S det (m, n) is a non-target detection point, and the detection result S is set det (m, n) = 0; if S abs The value of (m, n) is greater than or equal to S T (m, n), then determine S det (mm, n) is the target echo detection point, and the detection result S is set det (m, n) = 1;

[0097] 6-7) The test result S det (m, n) is moved from the shared memory to the global memory, and then the detection results are copied to the CPU using the PCIE data transmission bus. Finally, the memory and video memory space are released to complete the radar signal processing.

[0098] 1. Experimental Verification Conditions: The performance of this method was analyzed through simulation experiments. This example used Visual Studio 2017 and CUDA 11.4 programming on a Windows system as the software architecture. The hardware platform was an Intel Core i5-10400 CPU and an NVIDA Telsa A100 GPU. Different GPU cache mechanisms were employed to test the method under the condition of 16 radar pulses with 480 sampling points per pulse.

[0099] 2. Experimental results: Figure 5 The following is a comparison of the radar signal pulse compression, moving target display, moving target detection, and constant false alarm detection results using this method and using the GPU global memory to cache 16 radar pulses and 480 sampling points of echo data per pulse. Figure 6 The following are line graphs of two GPU cache types. As can be seen from the two figures, this method uses a grid stride loop approach, thread reuse, and shared memory data storage. Compared with the currently common radar signal processing method using global memory cache processing, this method is faster and less time-consuming, and can better meet the real-time requirements of radar signal processing.

Claims

1. A GPU+CPU architecture radar signal processing implementation method based on grid stride loop, characterized in that: The steps include: 1) Initialize radar parameters on the CPU host, open up memory and video memory space, and receive radar echo signal S echo (m, n), m = 1, 2, ... M, n = 1, 2, ... N, M represents the total number of radar echoes, and N represents the total number of sampling points for each echo; 2) The CPU transmits the radar echo signal S through the PCIE data transmission bus echo (m, n) data is copied from the CPU host to the GPU device and stored in the GPU global memory; 3) Radar echo signal S echo The (m, n) data is pulse compressed using a one-dimensional grid-stride cycle to obtain the pulse compression result S p (m, n), including: 3-1) The CPU host side flips the original linear frequency modulation signal left and right, and performs conjugate complex operation on the flipped result to obtain the pulse compression coefficient S c (k), k = 1, 2, ... K, K represents the total number of pulse compression coefficients and distributes the pulse compression result S p (m, n) video memory space and thread space are opened to start the pulse compression processing kernel function; 3-2) The echo signal S echo (m, n) data is vectorized to obtain the vectorized echo signal S echo (t), t = 1, 2, ... M*N, M*N represents the total number of radar echo samples; 3-3) The GPU device side opens a shared memory space in the pulse compression processing kernel function according to the number of threads in the block, and sequentially transmits the radar echo signal S echo (t) Data is cached from GPU global memory to shared memory; 3-4) Using a one-dimensional grid stride Grid-Stride loop, with the number of threads in a grid as the loop step, the one-dimensional grid threads are reused to form a thread fine-grained index, which corresponds one-to-one to the echo signal S after index vectorization echo (t), the thread index expression is: G idx =S echo (idx+i*gridDim.x*blockDim.x), idx=threadIdx.x+blockIdx.x*blockDim.x, Among them, G idx Represents a grid thread, idx is the thread index number within the grid, i = 0, 1, ... I, I represents the total number of grid stride loops, threadIdx.x represents the thread number within the thread block, and its value is threadIdx.x∈[0,blockDim.x), blockIdx.x represents the index of the thread block, and its value is blockIdx.x∈[0,gridDim.x), gridDim.x represents the dimension of the grid in the thread space, blockDim.x represents the dimension of the block in the thread space, and the values ​​of gridDim.x and blockDim.x are determined by the opened thread space; 3-5) The radar echo signal S cached in the shared memory is accessed by the thread fine-grained index echo (t), calculate the radar echo signal S echo (t) for pulse compression coefficient S c (k) is discretely convolved to obtain the vectorized pulse compression result S p (t), the operation expression is: 3-6) The pulse compression result S p (t) Move from shared memory to global memory; 4) Use shared memory to process the pulse compression result S p (t) Perform MTI processing and obtain the MTI processing result S MTI (m, n), specifically: 4-1) The pulse compression result S p (t) The pulse compression result S is a two-dimensional matrix in the form of a cache to facilitate data cancellation, with the total number of M radar echoes as a row and each echo with N sampling points as a column. p (m, n); 4-2) CPU allocation target display result S MTI (m, n) video memory space and thread space are opened, and the MTI processing kernel function is started; 4-3) In the MTI processing kernel function, the GPU opens a shared memory space according to the number of threads in the block, and uses the control thread index method to sequentially convert the pulse compression result S in the form of a two-dimensional matrix into p (m, n) data is cached from GPU global memory to shared memory; 4-4) Using a one-dimensional grid stride loop, with the number of threads in the grid as the loop step, thread reuse constitutes a thread fine-grained index, and one-to-one correspondence indexes the pulse compression result S in the form of a two-dimensional matrix p (m, n), the thread index expression is: G idx =S p (idx, idy*gridDim.x*blockDim.x), idx=threadIdx.x+blockIdx.x*blockDim.x, idy=threadIdx.y+blockIdx.y*blockDim.y, Among them G idx is the grid thread, idx is the thread index number in the grid, which also represents the x-axis index in the two-dimensional matrix, and idy is the y-axis index in the two-dimensional matrix, which also represents the number of grid stride cycles. threadIdx.y, blockIdx.y, blockDim.y and threadIdx.x, blockIdx.x, blockDim.x have the same meaning, representing the thread number in the y and x directions, the thread block index number, and the thread dimension within the block respectively; 4-5) The pulse compression result S cached in the shared memory is accessed by the thread fine-grained index p (m, n) and perform row subtraction to cancel the operation, and obtain the MTI processing result S MTI (m, n), the subtraction cancellation expression is: S MTI (m, n) = S p (m, n+1)-S p (m, n); 4-6) MTI processing result S MTI (m, n) is moved from shared memory to GPU global memory; 5) The pulse compression result S cached in a two-dimensional matrix form p (m, n) performs MTD processing and obtains the MTD processing result S MTD (m, n), the specific steps are as follows: 5-1) CPU allocates MTD processing result S MTD (m, n) video memory space; 5-2) Taking the length of a pulse signal as the sliding window size, perform M-point windowed FFT on M pulse pressure data in the same range Doppler dimension, and obtain the result S after FFT. MTD_FFT (m, n); 5-3) The result S after FFT MTD_FFT (m, n) performs matrix transposition to obtain the MTD processing result S MTD (m, n); 6) Using the data S cached in the GPU global memory MTD (m, n) performs constant false alarm detection according to the one-dimensional grid stride Grid-Stride cycle method, and screens the target to obtain the detection result S det (m, n), including: 6-1) GPU initializes the parameters of the constant false alarm detection module and calculates the moving target detection result S MTD (m, n) performs a modulo operation to obtain the moving target detection result S after modulo operation. abs (m, n); 6-2) The CPU allocates memory space for CFAR detection and creates thread space, and starts the CFAR detection processing kernel function; 6-3) In the constant false alarm detection processing kernel function, the GPU opens a shared memory space according to the number of threads in the block, and uses the control thread index method to sequentially load the modulo moving target detection result S abs (m, n) is cached from GPU global memory to shared memory; 6-4) Using a one-dimensional grid stride Grid-Stride loop, with the number of threads in the grid as the loop step, it loops through the modulo-based moving target detection results S in the shared memory. abs (m, n) data is reduced and summed row by row; 6-5) Use a sliding window to sum the reference unit data at the left and right ends of the reduced unit to be tested, and calculate the average value to obtain the noise background signal S noise (m, n); 6-6) Based on the noise background data S noise (m, n) solves the decision threshold S T (m, n), compare S abs (m, n) and the decision threshold S T The size of (m, n): If S abs The value of (m, n) is less than S T (m, n), then directly determine S det (m, n) is a non-target detection point, and the detection result S is set det (m, n) = 0; if S abs The value of (m, n) is greater than or equal to S T (m, n), then determine S det (m, n) is the target echo detection point, and the detection result S is set det (m, n) = 1; 6-7) The test result S det (m, n) is moved from the shared memory to the global memory, and then the detection results are copied to the CPU using the PCIE data transmission bus. Finally, the memory and video memory space are released to complete the radar signal processing.

Images