Heterogeneous computing-based memory management method, system and equipment and medium

Abstract

The invention discloses a memory management method, system and device based on heterogeneous computing and a medium. In the method, firstly, specific variables in a CUDA kernel function are declared as unified memory variables; respectively allocating a host memory to the specific variables at the CPU side, and allocating an equipment memory corresponding to the host memory to the GPU side; the specific variables are globally visible at the GPU side and the CPU side, respectively. Before the kernel function is called, data in a host memory is copied to a corresponding device memory; after the kernel function calling is finished, the calculation result in the device memory is copied back to the corresponding host memory, and the host memory is kept unchanged during the kernel function calling period and is not in linkage with the device memory. The data is copied only before and after the kernel function is called, so that unnecessary data transmission is reduced, and the overall calculation speed is effectively increased; meanwhile, through global visibility of the specific variable, programming complexity increase caused by dependent variable transmission limitation is avoided, and convenience and performance are both considered.

Description

Technical Field

[0001] The embodiments of this application relate to, but are not limited to, the fields of computer architecture and parallel computing technology, and in particular to memory management methods, systems, devices and media based on heterogeneous computing. Background Technology

[0002] In the field of CUDA-based GPU parallel computing, unified memory technology is widely used to simplify programming processes and reduce the complexity of heterogeneous memory management. This technology provides a shared global virtual address space between the CPU and GPU, eliminating the need for developers to manually maintain the address mapping between CPU host memory and GPU device memory. Variables can be directly accessed by both the CPU and GPU, reducing the amount of code required for parameter passing.

[0003] However, the convenience of unified memory technology comes at the cost of system resource overhead. To ensure data consistency between the corresponding storage units on the CPU and GPU sides, the system needs to monitor and synchronize data at both ends in real time: when the CPU or GPU accesses a unified memory variable, a data copy operation is automatically triggered in the background to maintain the same value. This real-time synchronization mechanism leads to a significant increase in data access latency, especially in high-frequency access or large data volume scenarios. Frequent cross-device data transfers can significantly offset the speed advantage brought by GPU parallel computing, and may even result in computational efficiency lower than pure CPU computing, which contradicts the core goal of CUDA technology to pursue efficient parallel computing. Summary of the Invention

[0004] The following is an overview of the subject matter described in detail herein. This overview is not intended to limit the scope of the claims.

[0005] This application provides a memory management method, system, device, and medium based on heterogeneous computing, which can effectively improve the overall computing speed and achieve a balance between convenience and efficiency in CUDA heterogeneous computing.

[0006] In a first aspect, embodiments of this application provide a memory management method based on heterogeneous computing, comprising: declaring certain specific variables in a CUDA kernel function as unified memory variables using managed memory keywords; allocating host memory on the CPU side for the specific variables declared as unified memory variables, and allocating device memory corresponding to the host memory on the GPU side for the specific variables, wherein the allocated memory is globally visible on both the GPU side and the CPU side; copying data in the host memory to the corresponding device memory before calling the kernel function; and copying the calculation results in the device memory back to the corresponding host memory after the kernel function call ends; wherein the host memory remains unchanged during the kernel function call and is not linked to the device memory.

[0007] In conjunction with the first aspect, in one embodiment of this application, allocating host memory on the CPU side for a specific variable that has been declared as the unified memory variable includes: allocating corresponding CPU-side host memory for an array pointer variable in the specific variable, wherein the capacity of the host memory matches the length of the array pointed to by the array pointer variable.

[0008] In conjunction with the first aspect, in one embodiment of this application, allocating device memory on the GPU side corresponding to the host memory for the specific variable includes: allocating GPU-side device memory for the array pointer variable in the specific variable, wherein the capacity of the device memory matches the length of the array pointed to by the array pointer variable.

[0009] In conjunction with the first aspect, in one embodiment of this application, copying the data in the host memory to the corresponding device memory includes: copying the initial data stored in the CPU-side host memory corresponding to the specific variable to the GPU-side device memory corresponding to the specific variable.

[0010] In conjunction with the first aspect, in one embodiment of this application, copying the calculation result in the device memory back to the corresponding host memory includes: copying the calculation result stored in the GPU-side device memory corresponding to the specific variable back to the CPU-side host memory corresponding to the specific variable.

[0011] In conjunction with the first aspect, in one embodiment of this application, the method further includes: calling the kernel function, wherein the kernel function directly accesses the device memory and does not require passing the specific variable as a parameter to the kernel function.

[0012] Secondly, embodiments of this application provide a memory management system for CUDA heterogeneous parallel computing, comprising: a variable declaration unit, used to declare certain specific variables in a CUDA kernel function as unified memory variables using managed memory keywords; a memory allocation unit, used to allocate host memory on the CPU side for the specific variables declared as the unified memory variables, and allocate device memory corresponding to the host memory on the GPU side for the specific variables, wherein the allocated memory is globally visible on both the GPU side and the CPU side; a data preprocessing unit, used to copy data from the host memory to the corresponding device memory before calling the kernel function; and a result return unit, used to copy the calculation results from the device memory back to the corresponding host memory after the kernel function call ends; wherein, during the kernel function execution, the host memory remains unchanged and is not linked to the device memory.

[0013] In conjunction with the second aspect, in one embodiment of this application, the memory allocation unit includes: a device memory allocation module, configured to allocate GPU-side device memory to the array pointer variable in the specific variable, wherein the capacity of the device memory matches the length of the array pointed to by the array pointer variable; and a host memory allocation module, configured to allocate corresponding CPU-side host memory to the array pointer variable in the specific variable, wherein the host memory matches the length of the array pointed to by the array pointer variable.

[0014] On the other hand, embodiments of this application provide an electronic device including at least one processor, at least one memory, and computer program instructions stored in the memory, which, when executed by the processor, implement the method described above.

[0015] On the other hand, embodiments of this application provide a computer-readable storage medium storing computer program instructions that, when executed by a processor, implement the method described above.

[0016] This application provides a memory management method, system, electronic device, and computer-readable storage medium based on heterogeneous computing. In this method, firstly, specific variables in a CUDA kernel function are declared as uniform memory variables using managed memory keywords. Then, host memory is allocated for the specific variables on the CPU side, and device memory corresponding to the host memory is allocated for the specific variables on the GPU side. The allocated memory is globally visible on both the GPU and CPU sides. Before calling the kernel function, data in the host memory is copied to the corresponding device memory; after the kernel function call ends, the computation results in the device memory are copied back to the corresponding host memory. The host memory remains unchanged during the kernel function call and is not linked to the device memory. This application's embodiments replace the real-time synchronization mechanism of unified memory by copying CPU-side host memory data to GPU-side device memory only before the CUDA kernel function call and returning the device memory calculation results to host memory after the kernel function call. This reduces unnecessary data transfers and makes the data access and calculation process more aligned with the high-efficiency characteristics of GPU parallel computing, effectively improving the overall calculation speed. At the same time, unified memory variables declared through the managed memory keyword have global visibility on both the CPU and GPU sides and can be directly accessed by kernel functions. This continues the convenience of unified memory in simplifying programming with multiple variables and complex data structures, avoiding the increase in programming complexity caused by variable passing restrictions, thus achieving a balance between convenience and efficiency. Attached Figure Description

[0017] Figure 1 This is a schematic diagram of the CPU and its memory access in related technologies; Figure 2This is a schematic diagram of the GPU and its memory access in related technologies; Figure 3 This is a schematic diagram of a heterogeneous platform composed of a CPU and a GPU provided in an embodiment of this application; Figure 4 This is a schematic diagram of a unified memory architecture for a heterogeneous platform provided in an embodiment of this application; Figure 5 This is a flowchart of a memory management method based on heterogeneous computing provided in an embodiment of this application; Figure 6 This is a block diagram of a memory management system based on heterogeneous computing provided in an embodiment of this application. Detailed Implementation

[0018] To make the objectives, technical solutions, and advantages of this application clearer, the following detailed description is provided in conjunction with the accompanying drawings and embodiments. It should be understood that the specific embodiments described herein are merely illustrative and not intended to limit the scope of this application.

[0019] It should be noted that although the flowchart shows a logical order, in some cases, the steps shown or described may be performed in a different order than that shown in the flowchart. The terms "first," "second," etc., used in the specification, claims, and the aforementioned drawings are used to distinguish similar objects and are not necessarily used to describe a specific order or sequence. It should be understood that the structures, proportions, sizes, etc., depicted in the drawings are only used to complement the content disclosed in the specification for those skilled in the art to understand and read, and are not intended to limit the implementation conditions of this application. Therefore, they have no substantial technical significance. Any modifications to the structure, changes in proportions, or adjustments to size, without affecting the effects and purposes achieved by this application, should still fall within the scope of the technical content disclosed in this application. Similarly, the terms such as "upper," "lower," "left," "right," "middle," and "one" used in this specification are only for clarity of description and are not used to limit the scope of implementation of this application. Changes or adjustments in their relative relationships, without substantially altering the technical content, should also be considered within the scope of implementation of this application.

[0020] Unless otherwise defined, all technical and scientific terms used herein have the same meaning as commonly understood by one of ordinary skill in the art to which this application belongs. The terminology used herein is for the purpose of describing embodiments of this application only and is not intended to limit this application.

[0021] In the field of CUDA-based GPU parallel computing, unified memory technology is widely used to simplify programming processes and reduce the complexity of heterogeneous memory management. This technology provides a shared global virtual address space between the CPU and GPU, eliminating the need for developers to manually maintain the address mapping between CPU host memory and GPU device memory. Variables can be directly accessed by both the CPU and GPU, reducing the amount of code required for parameter passing.

[0022] However, the convenience of unified memory technology comes at the cost of system resource overhead. To ensure data consistency between the corresponding storage units on the CPU and GPU sides, the system needs to monitor and synchronize data at both ends in real time: when the CPU or GPU accesses a unified memory variable, a data copy operation is automatically triggered in the background to maintain the same value. This real-time synchronization mechanism leads to a significant increase in data access latency, especially in high-frequency access or large data volume scenarios. Frequent cross-device data transfers can significantly offset the speed advantage brought by GPU parallel computing, and may even result in computational efficiency lower than pure CPU computing, which contradicts the core goal of CUDA technology to pursue efficient parallel computing.

[0023] In view of this, embodiments of this application provide a memory management method, a memory management system, an electronic device, and a computer-readable storage medium based on heterogeneous computing. In this method, firstly, certain specific variables in a CUDA kernel function are declared as unified memory variables using managed memory keywords; then, host memory is allocated for the specific variables on the CPU side, and device memory corresponding to the host memory is allocated for the specific variables on the GPU side, wherein the allocated memory is globally visible on both the GPU and CPU sides. Before calling the kernel function, data in the host memory is copied to the corresponding device memory; after the kernel function call ends, the computation results in the device memory are copied back to the corresponding host memory. The host memory remains unchanged during the kernel function call and is not linked to the device memory. This application's embodiments replace the real-time synchronization mechanism of unified memory by copying CPU-side host memory data to GPU-side device memory only before CUDA kernel function calls and returning the device memory calculation results to host memory after kernel function calls. This reduces unnecessary data transfers and makes the data access and calculation process more aligned with the high-efficiency characteristics of GPU parallel computing, effectively improving the overall calculation speed. At the same time, unified memory variables declared through managed memory keywords have global visibility on both the CPU and GPU sides and can be directly accessed by kernel functions. This continues the convenience of unified memory in simplifying programming with multiple variables and complex data structures, avoiding the increase in programming complexity caused by variable passing restrictions, thus achieving a balance between convenience and performance.

[0024] It is understandable that in CUDA-based heterogeneous computing scenarios, the access characteristics of variables are closely related to the hardware architecture. For example... Figure 1As shown, from the perspective of variable visibility, global variables in C language, because their lifecycle covers the entire program and they have global visibility, can be directly accessed by the CPU at any time without needing to be passed through function parameters. This characteristic is similar to the "CPU and GPU global visibility" goal pursued by unified memory technology, sharing a similar appeal in simplifying programming convenience. However, CPU access to memory requires via the BUS bus, and due to the physical characteristics of external memory, the access speed is relatively slow. Figure 2 As shown, the GPU's hardware architecture determines its advantage in parallel computing: the GPU processor and video memory are directly integrated on the board, allowing access to video memory without a bus, resulting in extremely fast access speeds; simultaneously, GPUs have tens of thousands of processor cores, far exceeding the number of CPU cores, making them naturally suitable for large-scale parallel computing. However, the integrated video memory also brings limitations: the capacity is fixed and limited (e.g., the A100 and H100 only have 80GB of video memory), making efficient use of video memory and reducing unnecessary data interactions key to improving GPU computing efficiency. Figure 3 As shown, in a heterogeneous platform composed of a CPU and a GPU, data is transferred between the two via a bus, and bus transfer itself incurs significant overhead. As mentioned in the background section, excessive data transfer severely negates the parallel advantages of the GPU. Figure 4 As shown, the core of unified memory technology for heterogeneous platforms is to build a shared virtual address space between CPU memory and GPU memory, enabling a logical correspondence between some of their storage units. This technology can bring great convenience to programming: for example, CUDA stipulates that variables in the global unified memory can be directly accessed by functions, and these variables can be used without parameter passing (virtual-real combination) during function calls. However, this convenience comes at a performance cost: to maintain the consistency of unified memory data in CPU memory and GPU memory, the system needs to frequently exchange data via the bus during function calls. The high-frequency access to data by the large number of processors (tens of thousands) on the GPU further exacerbates this exchange, ultimately severely reducing the execution speed of function calls.

[0025] The embodiments of this application will be further described below with reference to the accompanying drawings.

[0026] Reference Figure 5 , Figure 5 This is a flowchart of a memory management method based on heterogeneous computing provided in an embodiment of this application. The process may specifically include, but is not limited to, steps 510 to 540.

[0027] Step 510: Declare certain specific variables in the CUDA kernel functions as uniform memory variables using managed memory keywords; Step 520: Allocate host memory on the CPU side for the specific variable that has been declared as a uniform memory variable, and allocate device memory corresponding to the host memory on the GPU side for the specific variable. The allocated memory is globally visible on both the GPU side and the CPU side. Step 530: Before calling the kernel function, copy the data from the host memory to the corresponding device memory; Step 540: After the kernel function call ends, copy the calculation results from the device memory back to the corresponding host memory.

[0028] Steps 510 to 540 will be described in detail below.

[0029] Understandably, in CUDA programming, different types of memory arrays or pointers correspond to different storage locations and management methods: CPU-side memory arrays or pointers (such as `double x[n];` or `double* x = new double[n];`) have their storage units located in the CPU's memory modules and are directly managed by the host; GPU-side video memory arrays or pointers (such as `__device__ double* x;`, which needs to be allocated via `cudaMalloc(&x, n * sizeof(double))`) have their storage units located in the GPU's onboard video memory and are only accessible to the GPU processor. Unified Memory, however, is implemented differently. For example, variables declared using `__device__ __managed__ double* x;` and allocated via `cudaMallocManaged(&x, n * sizeof(double))` can simultaneously allocate corresponding storage units in both CPU memory and GPU video memory, and the system automatically maintains the numerical consistency of the storage units on both sides, meaning that variables on the CPU and GPU sides are always highly synchronized. This feature allows direct data interaction between the CPU and GPU, enabling functions to access variables without parameter passing (a combination of virtual and physical parameters), thus simplifying the programming process. However, this convenience comes at a significant cost: to maintain the same values ​​in the CPU and GPU storage units at all times, the system needs to frequently copy data between the CPU and GPU via the bus. This continuous synchronization operation can result in data access speeds that are even slower than when using the CPU alone.

[0030] In a feasible embodiment, in step 510, certain specific variables that need to be accessed in CUDA kernel functions (parallel functions running on the GPU) are declared as uniform memory variables using a specific keyword combination. These variables can be recognized and accessed by both the CPU and GPU simultaneously, without requiring additional address mapping configuration. The managed memory keyword refers to the keyword combination used to declare "uniform memory variables," and its core function is to give variables global visibility to both the CPU and GPU. For example, the __device__ and __managed__ keyword combination in CUDA (the former indicates that the variable can be accessed by the GPU, and the latter indicates that the variable supports shared address space between the CPU and GPU). Variables declared using managed memory keywords can be directly accessed by both parties and can be called by kernel functions without being passed as function parameters.

[0031] In a feasible embodiment, in step 520, independent physical memory is allocated on both the CPU and GPU sides for the specific variables that have been declared as unified memory variables, and the storage structures (such as array length and data type) are consistent to ensure accurate correspondence in subsequent data copying. Since the specific variables have been declared as unified memory variables in step 510, their addresses on the CPU and GPU sides are visible to both. Here, host memory refers to the memory on the CPU side (usually the storage space in a computer's RAM), directly managed by the CPU, used to store initial data or final results processed by the CPU; device memory refers to the video memory on the GPU side, typically integrated on the GPU board, directly managed by the GPU, and used to store intermediate data or computational objects required by kernel functions during runtime.

[0032] In a feasible embodiment, device memory can be allocated for these specific variables on the GPU side using device memory allocation functions, while on the CPU side, host memory corresponding to the device memory can be allocated for the specific variables using host memory allocation functions. Here, host memory allocation functions refer to functions that allocate host memory on the CPU side, such as the `new()` function in C++ or the `malloc()` function in C, used to allocate storage space for specific variables on the CPU side; device memory allocation functions refer to functions that allocate device memory on the GPU side, such as the `cudaMalloc()` function in CUDA, used to allocate storage space for specific variables on the GPU side that matches the host memory.

[0033] In a feasible embodiment, in step 530, since the GPU can only directly access data in the device memory, the initialized raw data in the CPU-side host memory needs to be transferred to the corresponding device memory on the GPU side to ensure that the kernel function can obtain the correct initial value during runtime. A copy operation is performed using a data copy function (such as the cudaMemcpy() function in CUDA), with the copy direction being from the CPU-side host memory to the GPU-side device memory. The amount of data copied matches the memory capacity allocated in step 520 (e.g., array length × number of bytes per element). Here, the data copy function refers to a function used to transfer data between the CPU host memory and the GPU device memory, supporting bidirectional transfer between the host and device, such as the cudaMemcpy() function in CUDA.

[0034] In one feasible embodiment, when calling the kernel function, the kernel function can directly access the device memory corresponding to a specific variable without passing the specific variable as a parameter to the kernel function. Since the specific variable has global visibility to both the CPU and GPU after being declared in step 510, the kernel function does not need to receive variables through a parameter list (avoiding parameter quantity limitations) and can directly access the data in the device memory through the variable name, simplifying the kernel function calling logic. This is especially suitable for scenarios with multi-variable arrays or dynamic variables (such as variables whose access objects cannot be determined during the programming phase).

[0035] It is worth noting that during the execution of the kernel function, the host memory data allocated to these specific variables on the CPU side remains fixed and is not synchronized in real time with the device memory allocated to these variables on the GPU side; only the device memory data on the GPU side changes as the kernel function performs its calculations.

[0036] In one feasible embodiment, in step 540, the result of the kernel function's calculation on the GPU is stored in the device memory and can be transferred back to the host memory on the CPU side so that the CPU can perform subsequent processing on the result (such as output, secondary calculation, etc.).

[0037] Through the processing flow of steps 510 to 540, globally visible variables are first declared, then memory is allocated separately, and then data is copied as needed. This not only retains the programming convenience of certain specific variables being globally visible and requiring no parameter passing, but also reduces unnecessary data transfer overhead by replacing the traditional unified memory "real-time synchronization" by "copying data only before and after kernel function calls". This achieves a balance between convenience and computational efficiency.

[0038] In a feasible embodiment, the process of allocating host memory on the CPU side for specific variables is as follows: CPU-side host memory is allocated for array pointer variables within the specific variables, with the capacity of the host memory matching the length of the array pointed to by the array pointer variable. This process primarily involves matching suitable storage space for the array pointer variables contained within these specific variables. Here, an array pointer variable refers to a pointer within a specific variable that points to a specific array (possessing global visibility to both the CPU and GPU). Its function is to associate array data in CPU-side host memory and GPU-side device memory through address mapping. When allocating host memory, the memory capacity can be determined based on the length of the array pointed to by the array pointer variable. For example, if the array pointer variable points to an array containing 1000 double-type elements, then the host memory capacity must be 1000 × sizeof(double).

[0039] In a feasible embodiment, the process of allocating device memory corresponding to host memory for a specific variable on the GPU side is as follows: GPU-side device memory is allocated for the array pointer variable within the specific variable, and the capacity of the device memory matches the length of the array pointed to by the array pointer variable. This process primarily allocates matching video memory space for the array pointer variable within the specific variable, ensuring that the GPU can efficiently access array data. The array pointer variable points to the allocated device memory array on the GPU side. When allocating device memory, the capacity can be determined based on the length of the array pointed to by the pointer. For example, if the array pointer variable points to an array containing 500 int type elements, then the capacity of the device memory needs to be 500 × sizeof(int).

[0040] In a feasible embodiment, the process of copying data from the host memory to the corresponding device memory is as follows: the initial data stored in the CPU-side host memory corresponding to the specific variable is copied to the GPU-side device memory corresponding to the specific variable. The CPU-side host memory corresponding to the specific variable refers to the storage space allocated on the CPU side for the specific variable (such as an array pointer variable) in step 520, which stores the original data required for kernel function calculation; the GPU-side device memory corresponding to the specific variable is the video memory space allocated on the GPU side in step 520 for the same specific variable, which matches the capacity of the host memory and is the storage area directly accessed by the kernel function. Through a data copy function (such as cudaMemcpy), the initial data stored in the CPU-side host memory is completely transferred to the corresponding device memory on the GPU side. The amount of data transferred is strictly matched with the length of the array pointed to by the array pointer variable (such as a double type array of length n, the amount of data transferred is n×sizeof(double)), ensuring that the data in the GPU-side device memory is completely consistent with the initial data on the CPU side, providing a foundation for the accurate execution of parallel computation by the subsequent kernel function. Meanwhile, because certain variables have global visibility, the kernel function on the GPU side can directly access the copied data through the variable name without passing any additional parameters. In a feasible embodiment, the process of copying the computation results from the device memory back to the corresponding host memory is as follows: the computation results stored in the GPU-side device memory corresponding to the specific variable are copied back to the CPU-side host memory corresponding to the specific variable. The GPU-side device memory corresponding to the specific variable stores the result data after parallel computation of the kernel function; the CPU-side host memory corresponding to the specific variable is the storage space that the CPU can directly access. After the copy is completed, the CPU can perform subsequent operations based on the results in the host memory, which not only ensures the integrity of the computation process, but also ensures the accuracy of data transmission by relying on the correspondence of specific variables declared as unified memory variables.

[0041] The memory management method of this embodiment is illustrated below through a specific example.

[0042] #include "cuda_runtime.h" #include "device_launch_parameters.h" #include <cstdio> #include <fstream> #include <cstring> #include <cstdlib> #include <cmath> #include <iostream> #include <valarray> #include<stdio.h> using namespace std; int restart = 0; int const nx0 = 385; int const ny0 = 385; int const nz0 = 385; / / s, 1, 2, ..., 32, 33, s int const nx = nx0 + 2; int const ny = ny0 + 2; int const nz = nz0 + 2; / / s, 1, 2, ..., 32, 33, s const int tmax = 100000000; int constnn0 = nx0 * ny0 * nz0; int node[nx][ny][nz]; int threadsPerBlock; int blocksPerGrid; int* hn0 = new int[nn0] { 0}; int* hn1 = new int[nn0] { 0}; int* hn2 = new int[nn0] { 0}; int* hn3 = new int[nn0] { 0}; int* hn4 = new int[nn0] { 0}; int* hn5 = new int[nn0] { 0}; int* hn6 = new int[nn0] { 0}; int* hn7 = new int[nn0] { 0}; int* hn8 = new int[nn0] { 0}; int* hn9 = new int[nn0] { 0}; int* hn10 = new int[nn0] { 0}; int* hn11 = new int[nn0] { 0}; int* hn12 = new int[nn0] { 0}; int* hn13 = new int[nn0] { 0}; int* hn14 = new int[nn0] { 0}; int* hn15 = new int[nn0] { 0}; int* hn16 = new int[nn0] { 0}; int* hn17 = new int[nn0] { 0}; int* hn18 = new int[nn0] { 0}; int* hn19 = new int[nn0] { 0}; int* hn20 = new int[nn0] { 0}; int* hn21 = new int[nn0] { 0}; int* hn22 = new int[nn0] { 0}; int* hn23 = new int[nn0] { 0}; int* hn24 = new int[nn0] { 0}; int* hn25 = new int[nn0] { 0}; int* hn26 = new int[nn0] { 0}; int* hx0 = new int[nn0] { 0}; int* hx1 = new int[nn0] { 0}; int* hx2 = new int[nn0] { 0}; double* hf0 = new double[nn0]; double* hf1 = new double[nn0]; double* hf2 = new double[nn0]; double* hf3 = new double[nn0]; double* hf4 = new double[nn0]; double* hf5 = new double[nn0]; double* hf6 = new double[nn0]; double* hf7 = new double[nn0]; double* hf8 = new double[nn0]; double* hf9 = new double[nn0]; double* hf10 = new double[nn0]; double* hf11 = new double[nn0]; double* hf12 = new double[nn0]; double* hf13 = new double[nn0]; double* hf14 = new double[nn0]; double* hf15 = new double[nn0]; double* hf16 = new double[nn0]; double* hf17 = new double[nn0]; double* hf18 = new double[nn0]; double* hf19 = new double[nn0]; double* hf20 = new double[nn0]; double* hf21 = new double[nn0]; double* hf22 = new double[nn0]; double* hf23 = new double[nn0]; double* hf24 = new double[nn0]; double* hf25 = new double[nn0]; double* hf26 = new double[nn0]; double* hden = new double[nn0]; double* hu0 = new double[nn0]; double* hu1 = new double[nn0]; double* hu2 = new double[nn0]; __device__ __managed__double u00 =0.0; __device__ __managed__double C = 1. / 3.; __device__ __managed__ double acclx0; __device__ __managed__ double accly0; __device__ __managed__ double acclz0; __device__ __managed__doubletao; __device__ __managed__doubletao1; __device__ __managed__doubletao2; __device__ __managed__doubletao3; __device__ __managed__int ec

[27] [3]; __device__ __managed__double wt

[27] ; __device__ __managed__int* x0; __device__ __managed__int* x1; __device__ __managed__int* x2; __device__ __managed__double* den; __device__ __managed__double* u0; __device__ __managed__double* u1; __device__ __managed__double* u2; __device__ __managed__int* n0; __device__ __managed__int* n1; __device__ __managed__int* n2; __device__ __managed__int* n3; __device__ __managed__int* n4; __device__ __managed__int* n5; __device__ __managed__int* n6; __device__ __managed__int* n7; __device__ __managed__int* n8; __device__ __managed__int* n9; __device__ __managed__int* n10; __device__ __managed__int* n11; __device__ __managed__int* n12; __device__ __managed__int* n13; __device__ __managed__int* n14; __device__ __managed__int* n15; __device__ __managed__int* n16; __device__ __managed__int* n17; __device__ __managed__int* n18; __device__ __managed__int* n19; __device__ __managed__int* n20; __device__ __managed__int* n21; __device__ __managed__int* n22; __device__ __managed__int* n23; __device__ __managed__int* n24; __device__ __managed__int* n25; __device__ __managed__int* n26; __device__ __managed__double* fc0; __device__ __managed__double* fc1; __device__ __managed__double* fc2; __device__ __managed__double* fc3; __device__ __managed__double* fc4; __device__ __managed__double* fc5; __device__ __managed__double* fc6; __device__ __managed__double* fc7; __device__ __managed__double* fc8; __device__ __managed__double* fc9; __device__ __managed__double* fc10; __device__ __managed__double* fc11; __device__ __managed__double* fc12; __device__ __managed__double* fc13; __device__ __managed__double* fc14; __device__ __managed__double* fc15; __device__ __managed__double* fc16; __device__ __managed__double* fc17; __device__ __managed__double* fc18; __device__ __managed__double* fc19; __device__ __managed__double* fc20; __device__ __managed__double* fc21; __device__ __managed__double* fc22; __device__ __managed__double* fc23; __device__ __managed__double* fc24; __device__ __managed__double* fc25; __device__ __managed__double* fc26; __device__ __managed__double* f0; __device__ __managed__double* f1; __device__ __managed__double* f2; __device__ __managed__double* f3; __device__ __managed__double* f4; __device__ __managed__double* f5; __device__ __managed__double* f6; __device__ __managed__double* f7; __device__ __managed__double* f8; __device__ __managed__double* f9; __device__ __managed__double* f10; __device__ __managed__double* f11; __device__ __managed__double* f12; __device__ __managed__double* f13; __device__ __managed__double* f14; __device__ __managed__double* f15; __device__ __managed__double* f16; __device__ __managed__double* f17; __device__ __managed__double* f18; __device__ __managed__double* f19; __device__ __managed__double* f20; __device__ __managed__double* f21; __device__ __managed__double* f22; __device__ __managed__double* f23; __device__ __managed__double* f24; __device__ __managed__double* f25; __device__ __managed__double* f26; __global__ void sstreamKernel() { int i = blockDim.x * blockIdx.x + threadIdx.x; if (i<nn0) { if (n0[i]>= 0) f0[i] = fc0[n0[i]]; if (n2[i]>= 0) f1[i] = fc1[n2[i]]; else f1[i] = fc2[n0[i]]; if (n1[i]>= 0) f2[i] = fc2[n1[i]]; else f2[i] = fc1[n0[i]]; if (n4[i]>= 0) f3[i] = fc3[n4[i]]; else f3[i] = fc4[n0[i]]; if (n3[i]>= 0) f4[i] = fc4[n3[i]]; else f4[i] = fc3[n0[i]]; if (n6[i]>= 0) f5[i] = fc5[n6[i]]; else f5[i] = fc6[n0[i]]; if (n5[i]>= 0) f6[i] = fc6[n5[i]]; else f6[i] = fc5[n0[i]]; if (n8[i]>= 0) f7[i] = fc7[n8[i]]; else f7[i] = fc8[n0[i]]; if (n7[i]>= 0) f8[i] = fc8[n7[i]]; else f8[i] = fc7[n0[i]]; if (n10[i]>= 0) f9[i] = fc9[n10[i]]; else f9[i] = fc10[n0[i]]; if (n9[i]>= 0) f10[i] = fc10[n9[i]]; else f10[i] = fc9[n0[i]]; if (n12[i]>= 0) f11[i] = fc11[n12[i]]; else f11[i] = fc12[n0[i]]; if (n11[i]>= 0) f12[i] = fc12[n11[i]]; else f12[i] = fc11[n0[i]]; if (n14[i]>= 0) f13[i] = fc13[n14[i]]; else f13[i] = fc14[n0[i]]; if (n13[i]>= 0) f14[i] = fc14[n13[i]]; else f14[i] = fc13[n0[i]]; if (n16[i]>= 0) f15[i] = fc15[n16[i]]; else f15[i] = fc16[n0[i]]; if (n15[i]>= 0) f16[i] = fc16[n15[i]]; else f16[i] = fc15[n0[i]]; if (n18[i]>= 0) f17[i] = fc17[n18[i]]; else f17[i] = fc18[n0[i]]; if (n17[i]>= 0) f18[i] = fc18[n17[i]]; else f18[i] = fc17[n0[i]]; if (n20[i]>= 0) f19[i] = fc19[n20[i]]; else f19[i] = fc20[n0[i]]; if (n19[i]>= 0) f20[i] = fc20[n19[i]]; else f20[i] = fc19[n0[i]]; if (n22[i]>= 0) f21[i] = fc21[n22[i]]; else f21[i] = fc22[n0[i]]; if (n21[i]>= 0) f22[i] = fc22[n21[i]]; else f22[i] = fc21[n0[i]]; if (n24[i]>= 0) f23[i] = fc23[n24[i]]; else f23[i] = fc24[n0[i]]; if (n23[i]>= 0) f24[i] = fc24[n23[i]]; else f24[i] = fc23[n0[i]]; if (n26[i]>= 0) f25[i] = fc25[n26[i]]; else f25[i] = fc26[n0[i]]; if (n25[i]>= 0) f26[i] = fc26[n25[i]]; else f26[i] = fc25[n0[i]]; }; }; __global__ void adiabaticBoundaryConditionKernel() { int i = blockDim.x * blockIdx.x + threadIdx.x; if (i<nn0) { if ((n6[i]<0)&&(n1[i]>= 0)&&(n2[i]>= 0)&&(n3[i]>= 0)&&(n4[i]>= 0)) { f5[i] = f6[i]; f11[i] = f12[i]+ ((f2[i] + f8[i]+ f9[i]) - (f1[i]+ f7[i] + f10[i])) / 2.0; f13[i] = f14[i]- ((f2[i] + f8[i]+ f9[i]) - (f1[i]+ f7[i] + f10[i])) / 2.0; f15[i] = f16[i]+ ((f4[i] + f8[i]+ f10[i]) - (f3[i]+ f7[i] + f9[i])) / 2.0; f17[i] = f18[i]- ((f4[i] + f8[i]+ f10[i]) - (f3[i]+ f7[i] + f9[i])) / 2.0; f19[i] = f20[i]; f21[i] = f22[i]; f23[i] = f24[i]; f25[i] = f26[i]; f0[i] = fc0[i]+ (fc6[i] + fc12[i]+ fc14[i] + fc16[i]+ fc18[i] + fc20[i]+ fc22[i] + fc24[i]+ fc26[i]) - (f5[i]+ f11[i] + f13[i]+ f15[i] + f17[i]+f19[i] + f21[i]+ f23[i] + f25[i]); }; if ((n5[i]<0)&&(n1[i]>= 0)&&(n2[i]>= 0)&&(n4[i]>= 0)&&(n3[i]>= 0)) { f6[i] = f5[i]; f14[i] = f13[i]+ ((f2[i] + f9[i]+ f8[i]) - (f1[i]+ f10[i] + f7[i])) / 2.0; f12[i] = f11[i]- ((f2[i] + f9[i]+ f8[i]) - (f1[i]+ f10[i] + f7[i])) / 2.0; f16[i] = f15[i]+ ((f3[i] + f9[i]+ f7[i]) - (f4[i]+ f10[i] + f8[i])) / 2.0; f18[i] = f17[i]- ((f3[i] + f9[i]+ f7[i]) - (f4[i]+ f10[i] + f8[i])) / 2.0; f22[i] = f21[i]; f20[i] = f19[i]; f26[i] = f25[i]; f24[i] = f23[i]; f0[i] = fc0[i]+ (fc5[i] + fc13[i]+ fc11[i] + fc15[i]+ fc17[i] + fc21[i]+ fc19[i] + fc25[i]+ fc23[i]) - (f6[i]+ f14[i] + f12[i]+ f16[i] + f18[i]+f22[i] + f20[i]+ f26[i] + f24[i]); }; if ((n4[i]<0)&&(n1[i]>= 0)&&(n2[i]>= 0)&&(n6[i]>= 0)&&(n5[i]>= 0)) { f3[i] = f4[i]; f7[i] = f8[i]+ ((f2[i] + f13[i]+ f12[i]) - (f1[i]+ f14[i] + f11[i])) / 2.0; f9[i] = f10[i]- ((f2[i] + f13[i]+ f12[i]) - (f1[i]+ f14[i] + f11[i])) / 2.0; f18[i] = f17[i]+ ((f5[i] + f13[i]+ f11[i]) - (f6[i]+ f14[i] + f12[i])) / 2.0; f15[i] = f16[i]- ((f5[i] + f13[i]+ f11[i]) - (f6[i]+ f14[i] + f12[i])) / 2.0; f24[i] = f23[i]; f26[i] = f25[i]; f21[i] = f22[i]; f19[i] = f20[i]; f0[i] = fc0[i]+ (fc4[i] + fc8[i]+ fc10[i] + fc17[i]+ fc16[i] + fc23[i]+ fc25[i] + fc22[i]+ fc20[i]) - (f3[i]+ f7[i] + f9[i]+ f18[i] + f15[i]+f24[i] + f26[i]+ f21[i] + f19[i]); }; if ((n3[i]<0)&&(n1[i]>= 0)&&(n2[i]>= 0)&&(n5[i]>= 0)&&(n6[i]>= 0)) { f4[i] = f3[i]; f10[i] = f9[i]+ ((f2[i] + f12[i]+ f13[i]) - (f1[i]+ f11[i] + f14[i])) / 2.0; f8[i] = f7[i]- ((f2[i] + f12[i]+ f13[i]) - (f1[i]+ f11[i] + f14[i])) / 2.0; f17[i] = f18[i]+ ((f6[i] + f12[i]+ f14[i]) - (f5[i]+ f11[i] + f13[i])) / 2.0; f16[i] = f15[i]- ((f6[i] + f12[i]+ f14[i]) - (f5[i]+ f11[i] + f13[i])) / 2.0; f25[i] = f26[i]; f23[i] = f24[i]; f20[i] = f19[i]; f22[i] = f21[i]; f0[i] = fc0[i]+ (fc3[i] + fc9[i]+ fc7[i] + fc18[i]+ fc15[i] + fc26[i]+ fc24[i] + fc19[i]+ fc21[i]) - (f4[i]+ f10[i] + f8[i]+ f17[i] + f16[i]+ f25[i] + f23[i]+ f20[i] + f22[i]); }; }; }; int main() { tao = 2.0; tao1 = 1 / tao; tao2 = (1 - tao1); tao3 = 1; acclx0 = 0.0001; accly0 = 0.0; acclz0 = 0.0; threadsPerBlock = 1024; blocksPerGrid = (nn0 + threadsPerBlock - 1) / threadsPerBlock; wt[0] = 8.0 / 27.0; wt[1] = 2.0 / 27.0; wt[2] = 2.0 / 27.0; wt[3] = 2.0 / 27.0; wt[4] = 2.0 / 27.0; wt[5] = 2.0 / 27.0; wt[6] = 2.0 / 27.0; wt[7] = 1.0 / 54.0; wt[8] = 1.0 / 54.0; wt[9] = 1.0 / 54.0; wt

[10] = 1.0 / 54.0; wt

[11] = 1.0 / 54.0; wt

[12] = 1.0 / 54.0; wt

[13] = 1.0 / 54.0; wt

[14] = 1.0 / 54.0; wt

[15] = 1.0 / 54.0; wt

[16] = 1.0 / 54.0; wt

[17] = 1.0 / 54.0; wt

[18] = 1.0 / 54.0; wt

[19] = 1.0 / 216.0; wt

[20] = 1.0 / 216.0; wt

[21] = 1.0 / 216.0; wt

[22] = 1.0 / 216.0; wt

[23] = 1.0 / 216.0; wt

[24] = 1.0 / 216.0; wt

[25] = 1.0 / 216.0; wt

[26] = 1.0 / 216.0; ec[0][0]= 0;ec[0][1]= 0;ec[0][2]= 0; ec[1][0]= 1;ec[1][1]= 0;ec[1][2]= 0; ec[2][0]= -1;ec[2][1]= 0;ec[2][2]= 0; ec[3][0]= 0;ec[3][1]= 1;ec[3][2]= 0; ec[4][0]= 0;ec[4][1]= -1;ec[4][2]= 0; ec[5][0]= 0;ec[5][1]= 0;ec[5][2]= 1; ec[6][0]= 0;ec[6][1]= 0;ec[6][2]= -1; ec[7][0]= 1;ec[7][1]= 1;ec[7][2]= 0; ec[8][0]= -1;ec[8][1]= -1;ec[8][2]= 0; ec[9][0]= -1;ec[9][1]= 1;ec[9][2]= 0; ec

[10] [0]= 1;ec

[10] [1]= -1;ec

[10] [2]= 0; ec

[11] [0]= 1;ec

[11] [1]= 0; ec

[11] [2]= 1; ec

[12] [0]= -1; ec

[12] [1]= 0; ec

[12] [2]= -1; ec

[13] [0]= -1;ec

[13] [1]= 0;ec

[13] [2]= 1; ec

[14] [0]= 1;ec

[13] [1] = 0;ec

[14] [2] = -1; ec

[15] [0]= 0;ec

[15] [1]= 1;ec

[15] [2]= 1; ec

[16] [0]= 0;ec

[16] [1]= -1;ec

[16] [2]= -1; ec

[17] [0]= 0;ec

[17] [1]= -1;ec

[17] [2]= 1; ec

[18] [0]= 0;ec

[18] [1]= 1;ec

[18] [2]= -1; ec

[19] [0]= 1;ec

[19] [1]= 1;ec

[19] [2]= 1; ec

[20] [0]= -1;ec

[20] [1]= -1;ec

[20] [2]= -1; ec

[21] [0]= -1;ec

[21] [1]= 1;ec

[21] [2]= 1; ec

[22] [0]= 1;ec

[22] [1]= -1;ec

[22] [2]= -1; ec

[23] [0]= -1;ec

[23] [1]= -1;ec

[23] [2]= 1; ec

[24] [0]= 1;ec

[24] [1]= 1;ec

[24] [2]= -1; ec

[25] [0]= 1;ec

[25] [1]= -1;ec

[25] [2]= 1; ec

[26] [0]= -1;ec

[26] [1]= 1;ec

[26] [2]= -1; cudaStatus = cudaSetDevice(0); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaSetDevice failed!Do you have a CUDA-capable GPUinstalled?"); return 1; } cudaMalloc(&n0, nn0 * sizeof(int)); cudaMalloc(&n1, nn0 * sizeof(int)); cudaMalloc(&n2, nn0 * sizeof(int)); cudaMalloc(&n3, nn0 * sizeof(int)); cudaMalloc(&n4, nn0 * sizeof(int)); cudaMalloc(&n5, nn0 * sizeof(int)); cudaMalloc(&n6, nn0 * sizeof(int)); cudaMalloc(&n7, nn0 * sizeof(int)); cudaMalloc(&n8, nn0 * sizeof(int)); cudaMalloc(&n9, nn0 * sizeof(int)); cudaMalloc(&n10, nn0 * sizeof(int)); cudaMalloc(&n11, nn0 * sizeof(int)); cudaMalloc(&n12, nn0 * sizeof(int)); cudaMalloc(&n13, nn0 * sizeof(int)); cudaMalloc(&n14, nn0 * sizeof(int)); cudaMalloc(&n15, nn0 * sizeof(int)); cudaMalloc(&n16, nn0 * sizeof(int)); cudaMalloc(&n17, nn0 * sizeof(int)); cudaMalloc(&n18, nn0 * sizeof(int)); cudaMalloc(&n19, nn0 * sizeof(int)); cudaMalloc(&n20, nn0 * sizeof(int)); cudaMalloc(&n21, nn0 * sizeof(int)); cudaMalloc(&n22, nn0 * sizeof(int)); cudaMalloc(&n23, nn0 * sizeof(int)); cudaMalloc(&n24, nn0 * sizeof(int)); cudaMalloc(&n25, nn0 * sizeof(int)); cudaMalloc(&n26, nn0 * sizeof(int)); cudaMalloc(&f0, nn0 * sizeof(double)); cudaMalloc(&f1, nn0 * sizeof(double)); cudaMalloc(&f2, nn0 * sizeof(double)); cudaMalloc(&f3, nn0 * sizeof(double)); cudaMalloc(&f4, nn0 * sizeof(double)); cudaMalloc(&f5, nn0 * sizeof(double)); cudaMalloc(&f6, nn0 * sizeof(double)); cudaMalloc(&f7, nn0 * sizeof(double)); cudaMalloc(&f8, nn0 * sizeof(double)); cudaMalloc(&f9, nn0 * sizeof(double)); cudaMalloc(&f10, nn0 * sizeof(double)); cudaMalloc(&f11, nn0 * sizeof(double)); cudaMalloc(&f12, nn0 * sizeof(double)); cudaMalloc(&f13, nn0 * sizeof(double)); cudaMalloc(&f14, nn0 * sizeof(double)); cudaMalloc(&f15, nn0 * sizeof(double)); cudaMalloc(&f16, nn0 * sizeof(double)); cudaMalloc(&f17, nn0 * sizeof(double)); cudaMalloc(&f18, nn0 * sizeof(double)); cudaMalloc(&f19, nn0 * sizeof(double)); cudaMalloc(&f20, nn0 * sizeof(double)); cudaMalloc(&f21, nn0 * sizeof(double)); cudaMalloc(&f22, nn0 * sizeof(double)); cudaMalloc(&f23, nn0 * sizeof(double)); cudaMalloc(&f24, nn0 * sizeof(double)); cudaMalloc(&f25, nn0 * sizeof(double)); cudaMalloc(&f26, nn0 * sizeof(double)); cudaMalloc(&fc0, nn0 * sizeof(double)); cudaMalloc(&fc1, nn0 * sizeof(double)); cudaMalloc(&fc2, nn0 * sizeof(double)); cudaMalloc(&fc3, nn0 * sizeof(double)); cudaMalloc(&fc4, nn0 * sizeof(double)); cudaMalloc(&fc5, nn0 * sizeof(double)); cudaMalloc(&fc6, nn0 * sizeof(double)); cudaMalloc(&fc7, nn0 * sizeof(double)); cudaMalloc(&fc8, nn0 * sizeof(double)); cudaMalloc(&fc9, nn0 * sizeof(double)); cudaMalloc(&fc10, nn0 * sizeof(double)); cudaMalloc(&fc11, nn0 * sizeof(double)); cudaMalloc(&fc12, nn0 * sizeof(double)); cudaMalloc(&fc13, nn0 * sizeof(double)); cudaMalloc(&fc14, nn0 * sizeof(double)); cudaMalloc(&fc15, nn0 * sizeof(double)); cudaMalloc(&fc16, nn0 * sizeof(double)); cudaMalloc(&fc17, nn0 * sizeof(double)); cudaMalloc(&fc18, nn0 * sizeof(double)); cudaMalloc(&fc19, nn0 * sizeof(double)); cudaMalloc(&fc20, nn0 * sizeof(double)); cudaMalloc(&fc21, nn0 * sizeof(double)); cudaMalloc(&fc22, nn0 * sizeof(double)); cudaMalloc(&fc23, nn0 * sizeof(double)); cudaMalloc(&fc24, nn0 * sizeof(double)); cudaMalloc(&fc25, nn0 * sizeof(double)); cudaMalloc(&fc26, nn0 * sizeof(double)); cudaMalloc(&u0, nn0 * sizeof(double)); cudaMalloc(&u1, nn0 * sizeof(double)); cudaMalloc(&u2, nn0 * sizeof(double)); cudaMalloc(&den, nn0 * sizeof(double)); cudaMalloc(&x0, nn0 * sizeof(int)); cudaMalloc(&x1, nn0 * sizeof(int)); cudaMalloc(&x2, nn0 * sizeof(int)); cudaDeviceSynchronize(); cudaStatus = cudaGetLastError(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc devic failed!"); return 1; } SoilStructure(); cudaMemcpy(x0, hx0, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(x1, hx1, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(x2, hx2, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n0, hn0, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n1, hn1, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n2, hn2, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n3, hn3, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n4, hn4, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n5, hn5, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n6, hn6, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n7, hn7, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n8, hn8, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n9, hn9, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n10, hn10, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n11, hn11, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n12, hn12, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n13, hn13, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n14, hn14, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n15, hn15, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n16, hn16, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n17, hn17, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n18, hn18, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n19, hn19, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n20, hn20, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n21, hn21, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n22, hn22, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n23, hn23, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n24, hn24, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n25, hn25, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaMemcpy(n26, hn26, (size_t)(nn0 * sizeof(int)),cudaMemcpyHostToDevice); cudaDeviceSynchronize(); cudaStatus = cudaGetLastError(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "copy state failed!"); return 1; } lbm_ns_3d(); return 0; } The CUDA parallel computing scenario shown in the above examples is based on the "flow field" as a specific case. The "flow field" refers to the physical field that characterizes the motion state of a fluid (liquid, gas, etc.), which is the set of motion parameters (such as velocity, particle distribution, pressure, etc.) at all locations when the fluid moves in space.

[0043] In this example, `__global__ void sstreamKernel()` and `__global__ void adiabaticBoundaryConditionKernel()` are two core CUDA kernel functions, involving a large number of variables and dynamic access logic. From the perspective of variable composition, the example contains multiple types of key arrays: the node numbers of the flow field are stored in a three-dimensional array `node[][][]`; the topology of the flow field is described by 27 integer arrays `n0[]`, `n1[]`, ..., `n26[]`; the real-time direction particle count of the flow field is stored in 27 double-precision arrays `f0[]`, `f1[]`, ..., `f26[]`; and the collision direction particle count of the flow field is stored in 27 double-precision arrays `fc0[]`, `fc1[]`, ..., `fc26[]`. The number of these arrays is enormous (27 each for the `n`, `f`, and `fc` series alone), far exceeding the number of variables that the CUDA kernel function parameter list can hold. This would be impossible to achieve through traditional parameter passing (a combination of virtual and real parameters). More importantly, there are a lot of dynamic variable accesses in kernel functions. For example, in sstreamKernel, there are logics such as f1[i] = fc1[n2[i]] and f2[i] = fc2[n1[i]], and in adiabaticBoundaryConditionKernel, there are conditional judgments and assignments based on indices such as n6[i] and n1[i]. The access objects (such as the specific elements of fc1) need to be dynamically determined at runtime through indices such as n2[i]. This cannot be predicted during the programming stage, and therefore cannot be passed as parameters.

[0044] To address this issue, in this example, all array pointers that need to be accessed within the kernel function (such as n0 to n26, f0 to f26, fc0 to fc26, etc.) are declared as unified memory variables using the `__device__` and `__managed__` keywords. This declaration provides global visibility to both the CPU and GPU, allowing the kernel function to access them directly by their names without passing them as parameters, perfectly adapting to the access requirements of multiple and dynamic variables. However, unlike traditional unified memory, this example employs a separation strategy during memory allocation: on the GPU side, device memory is allocated for these unified memory pointers using `cudaMalloc(&n0, nn0 * sizeof(int))`, rather than `cudaMallocManaged`; on the CPU side, corresponding host memory is allocated using `new` (such as `int* hn0 = new int[nn0] {0}`), forming a "host memory - device memory" correspondence.

[0045] Data interaction occurs only twice, before and after the kernel function call: before the call, initial data (such as hn0, hx0) is copied from the CPU-side host memory to the GPU-side device memory (such as n0, x0) via cudaMemcpy, in the direction cudaMemcpyHostToDevice; after the kernel function finishes execution, the computation results are copied back from the device memory to the host memory via cudaMemcpyDeviceToHost. This approach avoids the real-time synchronous copying required by traditional unified memory to maintain data consistency, retaining only the necessary two data transfers. While ensuring that the kernel function can directly access global variables, it significantly reduces data interaction overhead, thus maintaining the programming convenience of unified memory while aligning with the core goal of GPU parallel computing: high efficiency.

[0046] It should be noted that although the embodiments of this application differ to some extent from the CUDA specification, actual verification (as shown in the above example) shows that this solution does not cause any functional abnormalities, but can effectively improve the execution speed of function calls.

[0047] In addition, embodiments of this application also provide a memory management system for CUDA heterogeneous parallel computing, such as... Figure 6 As shown, the system 600 includes: Variable declaration unit 610 is used to declare certain specific variables in CUDA kernel functions as uniform memory variables using managed memory keywords; Memory allocation unit 620 is used to allocate host memory on the CPU side for a specific variable that has been declared as a uniform memory variable, and to allocate device memory corresponding to the host memory on the GPU side for the specific variable. The allocated memory is globally visible on both the GPU side and the CPU side. The data preprocessing unit 630 is used to copy data from the host memory to the corresponding device memory before calling the kernel function; The result return unit 640 is used to copy the calculation results in the device memory back to the corresponding host memory after the kernel function call ends. During the execution of the kernel function, the host memory remains unchanged and is not linked to the device memory.

[0048] In one feasible embodiment, the memory allocation unit includes: The device memory allocation module is used to allocate GPU-side device memory for array pointer variables in the unified memory variable. The capacity of the device memory matches the length of the array pointed to by the array pointer variable. The host memory allocation module is used to allocate corresponding CPU-side host memory for array pointer variables, and the length of the host memory matches the length of the array pointed to by the array pointer variable.

[0049] It should be noted that the execution logic of each unit of the memory management system 600 corresponds one-to-one with the steps of the memory management method described above. For details, please refer to the relevant descriptions above, which will not be repeated here.

[0050] This application also discloses an electronic device, which includes at least one processor, at least one memory, and computer program instructions stored in the memory. When the computer program instructions are executed by the processor, the memory management method based on heterogeneous computing, as described above, is implemented.

[0051] This application also discloses a computer-readable storage medium storing a processor-executable program, which, when executed by a processor, is used to perform the memory management method based on heterogeneous computing as described above.

[0052] The above description of the disclosed embodiments enables those skilled in the art to make or use this application. Various modifications to these embodiments will be readily apparent to those skilled in the art, and the general principles defined herein may be implemented in other embodiments without departing from the spirit or scope of this application. Therefore, this application is not to be limited to the embodiments shown herein, but is to be accorded the widest scope consistent with the principles and novel features disclosed herein.< / valarray> < / iostream> < / cmath> < / cstdlib> < / cstring> < / fstream> < / cstdio>

Claims

1. A memory management method based on heterogeneous computing, characterized by, The method comprises the following steps: Some specific variables in a CUDA kernel function are declared as unified memory variables by means of a managed memory keyword; Host memory is allocated for the specific variables which have been declared as the unified memory variables on the CPU side, and device memory corresponding to the host memory is allocated for the specific variables on the GPU side, wherein the allocated memory is globally visible on the GPU side and the CPU side respectively; Data in the host memory is copied to the corresponding device memory before the kernel function is called; The calculation result in the device memory is copied back to the corresponding host memory after the kernel function call ends; During the kernel function call, the host memory remains unchanged and is not linked with the device memory.

2. The method of claim 1, wherein, The step of allocating host memory for the specific variables which have been declared as the unified memory variables on the CPU side comprises the following steps: CPU-side host memory corresponding to the array pointer variable in the specific variable is allocated, and the capacity of the host memory matches the length of the array pointed to by the array pointer variable.

3. The method of claim 1, wherein, The step of allocating device memory corresponding to the host memory for the specific variables on the GPU side comprises the following steps: GPU-side device memory corresponding to the array pointer variable in the specific variable is allocated, and the capacity of the device memory matches the length of the array pointed to by the array pointer variable.

4. The method of claim 1, wherein, The step of copying data in the host memory to the corresponding device memory comprises the following steps: The initial data stored in the CPU-side host memory corresponding to the specific variable is copied to the GPU-side device memory corresponding to the specific variable.

5. The method of claim 1, wherein, The step of copying the calculation result in the device memory back to the corresponding host memory comprises the following steps: The calculation result stored in the GPU-side device memory corresponding to the specific variable is copied back to the CPU-side host memory corresponding to the specific variable.

6. The method of claim 1, wherein, The method further comprises the following step:

7. A memory management system for CUDA heterogeneous parallel computing, characterized in that, The kernel function is called, wherein the kernel function directly accesses the device memory, and the specific variable does not need to be passed into the kernel function as a parameter. The method comprises the following steps: A variable declaration unit is configured to declare some specific variables in a CUDA kernel function as unified memory variables by means of a managed memory keyword; A memory allocation unit is configured to allocate host memory for the specific variables which have been declared as the unified memory variables on the CPU side, and to allocate device memory corresponding to the host memory for the specific variables on the GPU side, wherein the allocated memory is globally visible on the GPU side and the CPU side respectively; A data preprocessing unit is configured to copy data in the host memory to the corresponding device memory before the kernel function is called; A result returning unit is configured to copy the calculation result in the device memory back to the corresponding host memory after the kernel function call ends; 8. The system of claim 7, wherein, During the kernel function call, the host memory remains unchanged and is not linked with the device memory. The memory allocation unit comprises the following steps: A device memory allocation module is configured to allocate GPU-side device memory for the array pointer variable in the specific variable, and the capacity of the device memory matches the length of the array pointed to by the array pointer variable. A host memory allocation module is configured to allocate a corresponding CPU-side host memory for an array pointer variable in the specific variable, the host memory matching a length of an array pointed to by the array pointer variable.

9. An electronic device, comprising: Comprising: at least one processor, at least one memory, and computer program instructions stored in the memory that, when executed by the processor, implement the method of any of claims 1-6.

10. A computer-readable storage medium having stored thereon computer program instructions, wherein, When the computer program instructions are executed by a processor, the method of any of claims 1-6 is implemented. When the computer program instructions are executed by a processor, the method of any of claims 1-6 is implemented.

Images