A memory write method, device, storage medium and program product

Abstract

The application discloses a memory writing method and device, a storage medium and a program product, relates to the technical field of memory access, and comprises the following steps: determining the memory step length of an initial image tensor corresponding to a target image, determining a target vectorization width according to the memory step length and the memory data width of a target graphics processor, determining the target data dimension of the target graphics processor, determining a target thread grid corresponding to the initial image tensor according to the target data dimension, performing image tensor slicing again by using the target thread grid, and obtaining a target image tensor written into the memory. The memory layout (step length) of the input tensor can be dynamically analyzed, so that the optimal vectorization width is adaptively determined, the memory access efficiency is maximized, slicing is performed according to the data dimension when the target graphics processor outputs data, the adaptability and the efficiency are combined, the GPU display memory bandwidth is fully utilized, and the memory bandwidth utilization rate is improved.

Description

Technical Field

[0001] The present application relates to the field of memory access technology, and in particular to a memory writing method, device, storage medium and program product. Background Art

[0002] Image tensor operations are frequently used when processing image data on GPUs (Graphics Processing Units), such as those used for facial recognition. As model sizes continue to expand and data dimensions become increasingly complex, higher demands are placed on the computational efficiency of image tensor operations. Image tensor slicing, a key operation for data preprocessing and feature extraction, has a direct impact on overall computing performance. Traditional image tensor slicing operations often suffer from issues such as discontinuous memory access and poor data locality, making it difficult to fully utilize the hardware's computing potential on parallel computing devices like GPUs. Summary of the Invention

[0003] The present application provides a memory writing method, device, storage medium and program product to solve the problem of poor writing effect when writing tensors into memory after performing tensor processing operations on facial image-related tensors.

[0004] This application provides a memory writing method, including:

[0005] Acquire a target image, determine an initial image tensor corresponding to the target image, and determine a memory step size of the initial image tensor; wherein the initial image tensor is a discontinuous image tensor;

[0006] Determining a target vectorization width corresponding to the initial image tensor according to the memory stride and the memory data width of the target graphics processor;

[0007] Determining a target data dimension of the target graphics processor, and determining a target thread grid corresponding to the initial image tensor according to the target data dimension; the target data dimension is a data dimension when the target graphics processor outputs data;

[0008] The initial image tensor is sliced ​​using the target thread grid to obtain a target image tensor, and the target image tensor is written into a target memory.

[0009] The present application also provides an electronic device, comprising: a memory for storing a computer program; and a processor for implementing the steps of any one of the above-mentioned memory writing methods when executing the computer program.

[0010] The present application also provides a computer-readable storage medium, in which a computer program is stored. When the computer program is executed by a processor, the steps of any of the above-mentioned memory writing methods are implemented.

[0011] The present application also provides a computer program product, including a computer program, which implements the steps of any of the above-mentioned memory writing methods when executed by a processor.

[0012] This application first determines the memory stride of the initial image tensor corresponding to the target image, determines the target vectorization width based on the memory stride and the memory data width of the target graphics processor, and determines the target data dimension of the target graphics processor. The target thread grid corresponding to the initial image tensor is determined based on the target data dimension, and then the target thread grid is used to slice the image tensor to obtain the target image tensor and write it to the memory. Through this application, the memory layout (stride) of the input tensor can be dynamically analyzed to adaptively determine the optimal vectorization width to maximize memory access efficiency, and slice according to the data dimension when the target graphics processor outputs data. The performance bottleneck in tensor operations is solved from the dimension of memory access optimization (dynamic vector width). It is suitable for scenarios when operating facial image tensors in irregular memory layouts. It can fully utilize the GPU memory bandwidth when performing facial image tensor operations and improve memory bandwidth utilization. BRIEF DESCRIPTION OF THE DRAWINGS

[0013] In order to more clearly illustrate the embodiments of the present application, the following is a brief introduction to the drawings required for use in the embodiments. Obviously, the drawings described below are only some embodiments of the present application. For ordinary technicians in this field, other drawings can be obtained based on these drawings without any creative work.

[0014] Figure 1 A flow chart of a memory writing method provided in an embodiment of the present application;

[0015] Figure 2 A schematic diagram of an adaptive vectorization decision mechanism provided in an embodiment of the present application;

[0016] Figure 3 A flow chart of a specific memory writing method provided in an embodiment of the present application;

[0017] Figure 4 A tensor slicing flowchart provided for an embodiment of the present application. DETAILED DESCRIPTION

[0018] The following will be combined with the accompanying drawings in the embodiments of this application to clearly and completely describe the technical solutions in the embodiments of this application. Obviously, the embodiments described are only part of the embodiments of this application, not all of them. Based on the embodiments in this application, all other embodiments obtained by ordinary technicians in this field without making creative efforts are within the scope of protection of this application.

[0019] It should be noted that, in the description of this application, the terms "comprises," "includes," or any other variations thereof are intended to cover non-exclusive inclusion, such that a process, method, article, or device comprising a series of elements includes not only those elements, but also other elements not explicitly listed, or elements inherent to such process, method, article, or device. The terms "first," "second," etc., in this application are used to distinguish similar objects, and are not used to describe a particular order or sequence.

[0020] Traditional image tensor slicing operations often face issues such as discontinuous memory access and poor data locality, making it difficult to fully exploit the hardware computing potential on parallel computing devices such as GPUs. Furthermore, current image tensor computation frameworks typically use fixed vectorization strategies that are unable to adapt to the characteristics of different image tensor layouts. This new approach analyzes the memory layout of the input tensor and adaptively determines the optimal vectorization width to maximize memory access efficiency. This approach is suitable for high-performance computing scenarios with irregular memory layouts, fully utilizing GPU memory bandwidth and improving memory bandwidth utilization.

[0021] In order to enable those skilled in the art to better understand the present application, the present application is further described in detail below with reference to the accompanying drawings and specific implementation methods.

[0022] The embodiment of the present application provides a memory writing method, and the method is described in detail in conjunction with the execution process of the memory writing method, such as Figure 1 As shown, the memory writing method includes:

[0023] Step S11: Acquire a target image, determine an initial image tensor corresponding to the target image, and determine a memory step size of the initial image tensor; wherein the initial image tensor is a discontinuous image tensor.

[0024] This embodiment is mainly used in the tensor processing process of the GPU of the AI ​​(Artificial Intelligence) computing framework, and writes the processed tensor into the corresponding memory based on CUDA (Compute Unified Device Architecture). It can be understood that the technical solution in this embodiment is not only applicable to the AI ​​field, but can also be applied to other fields that need to process multi-dimensional data and complex data formats, such as image processing, signal processing, etc., to expand and optimize the efficiency of data processing and calculation in these fields. After obtaining the target image, this embodiment first needs to determine the non-continuous initial image tensor corresponding to the target image and determine the memory stride of the initial image tensor, so as to dynamically analyze the memory layout (i.e., memory stride) of the input tensor through an adaptive memory access optimization mechanism.

[0025] Step S12: Determine a target vectorization width corresponding to the initial image tensor based on the memory stride and the memory data width of the target graphics processor.

[0026] In this embodiment, the target vectorization width corresponding to the initial image tensor can be determined based on the memory step size and the memory data width of the target graphics processor. Specifically, the memory step size of the initial image tensor can be determined first, and the minimum positive step size of the initial image tensor can be determined based on the memory step size. Then, the memory data width of the target graphics processor can be determined, and the target vectorization width corresponding to the initial image tensor can be determined based on the minimum step size and the memory data width. And when determining the target vectorization width corresponding to the initial image tensor based on the minimum step size and the memory data width, the byte step size corresponding to the initial image tensor can be determined based on the minimum step size and the number of bytes of the initial image tensor, and then the target vectorization width corresponding to the initial image tensor can be determined based on the byte step size and the memory data width. In this way, the dynamic vectorization decision algorithm based on the tensor memory layout characteristics can accurately determine the optimal vectorization width.

[0027] That is to say, the adaptive memory access optimization mechanism in this embodiment can accurately locate the minimum positive stride (min_stride) of the data in the memory by dynamically analyzing the memory layout of the input tensor, that is, the core indicator reflecting the most compact continuity of the memory obtained after traversing all dimensional strides, thereby adaptively determining the optimal vectorization width. Specifically, the byte stride corresponding to the minimum positive stride (byte_stride=min_stride×elem_size (number of bytes)) can be calculated, and the vectorization strategy can be dynamically selected based on the memory transaction width (i.e., memory data width) of the Ampere architecture.

[0028] Specifically, in the process of dynamically selecting the vectorization strategy, the byte step can be divided by the memory data width; if the memory data width is divisible by the byte step, the target vectorization width is determined to be four; if the memory data width is not divisible by the byte step, and the memory data width is divisible by twice the byte step, the target vectorization width is determined to be two; if the memory data width is not divisible by the byte step, and the memory data width is not divisible by twice the byte step, the target vectorization width is determined to be one. In a specific embodiment, if Figure 2 As shown, if the memory data width is 128, then if 128 bits is divisible by byte_stride, 4-element vectorization is used; if it is divisible by byte_stride × 2, 2-element vectorization is used; otherwise, it degenerates to scalar access. This helps maximize memory access efficiency, optimize memory merged access, reduce the number of transactions, and adapt to arbitrary memory continuity characteristics without assuming data layout. By accurately identifying memory bottlenecks based on the minimum positive stride, reducing transaction fragmentation and cache pollution, and closely integrating with hardware features to achieve full memory merged access, it ultimately significantly improves data loading and storage efficiency at the hardware level. This is particularly suitable for deep learning and scientific computing scenarios that are sensitive to memory bandwidth.

[0029] A specific implementation process is as follows:

[0030] / / Host-side vector width decision function:

[0031] int decide_vector_width(int* strides, int dims, int elem_size) {

[0032] int min_stride = INT_MAX;

[0033] for(int d=0; d <dims; ++d){

[0034] if(strides[d]>0)

[0035] min_stride = min(min_stride, strides[d]);

[0036] };

[0037] / / Calculate byte stride alignment:

[0038] const int bytes_per_access = 128; / / Ampere architecture memory transaction width

[0039] int byte_stride = min_stride * elem_size;

[0040] if(bytes_per_access % byte_stride == 0) return 4;

[0041] else if(bytes_per_access % byte_stride*2 ==0) return 2;

[0042] else return 1;

[0043] }.

[0044] Based on the above function, this embodiment can dynamically determine the vectorization width of memory access by analyzing the strides of the input tensor. First, traverse all dimensions to find the minimum positive stride (min_stride), and then calculate the byte stride (byte_stride = min_stride * elem_size). According to the 128-byte memory transaction width of the Ampere architecture, if the byte stride can be divided by 128, 4-element vectorization is selected (most efficient); if it can be divided by 64 (ie, byte_stride*2), 2-element vectorization is selected; otherwise, it degenerates to scalar access (width 1). In this way, the optimal vectorization width can be dynamically determined by analyzing the memory layout (stride) of the input tensor to maximize memory access efficiency.

[0045] Step S13: determining a target data dimension of the target graphics processor, and determining a target thread grid corresponding to the initial image tensor according to the target data dimension; the target data dimension is the data dimension when the target graphics processor outputs data.

[0046] In this embodiment, the target data dimension of the target GPU can be determined, and the target thread grid corresponding to the initial image tensor can be determined based on the target data dimension. The target data dimension is the data dimension of the target GPU output data, i.e., the data dimension corresponding to the output tensor. In this embodiment, after slicing, the input tensor is read and processed from non-contiguous memory via vectorized load (vectorized_load) in the kernel function (slice_kernel). The slicing strategy ensures that the kernel function processes the input data in a planned parallel manner and ultimately writes the results to the output tensor (via vectorized_store).

[0047] Step S14: Slice the initial image tensor using the target thread grid to obtain a target image tensor, and write the target image tensor into the target memory.

[0048] In this embodiment, the determined target thread grid can be used to slice the initial image tensor to obtain the target image tensor, and the target image tensor can be written into the target memory. Before slicing, a kernel function corresponding to each target vectorization width can be constructed so that the kernel function corresponding to the target vectorization width can be directly called to slice the initial image tensor. That is, in this embodiment, the corresponding kernel function can be pre-constructed based on each vectorization width, and then the kernel function template that matches the vectorization width can be automatically selected for operation when slicing the tensor. That is, the kernel function slice_kernel in this embodiment is defined as a template function, such as slice_kernel<T,VEC_WIDTH> , where VEC_WIDTH is determined at compile time, but different versions can be dynamically selected at runtime based on the stride of the input tensor (for example, via a switch statement in the launch_slice function). This allows pre-compiled kernels with different vectorization widths (1, 2, or 4) to be automatically selected at runtime based on the dynamic decision. This combines compile-time optimization with runtime flexibility, effectively improving computational efficiency.

[0049] Furthermore, in this embodiment, before slicing the initial image tensor, the initial pointer can be converted into a target pointer of the corresponding data type according to the current target vectorization width of the initial image tensor, and the initial image tensor can be read using the pointer so as to slice the initial image tensor using the target thread grid. The specific example is shown in the following template function:

[0050] / / Vectorized loading template (FP32 example):

[0051] template<int VEC_WIDTH>

[0052] __device__ void vectorized_load(const float* addr, float* reg) {

[0053] static_assert(VEC_WIDTH<=4, "Unsupported vector width");

[0054] if constexpr (VEC_WIDTH ==4) {

[0055] const float4 vec = *reinterpret_cast<const float4*> (addr);

[0056] reg[0]=vec.x; reg[1]=vec.y; reg[2]=vec.z; reg[3]=vec.w;

[0057] }

[0058] / / Other widths are implemented similarly...

[0059] };

[0060] Based on the above template function, different loading methods can be selected based on the compile-time constant VEC_WIDTH (1, 2, or 4). For example, when VEC_WIDTH = 4, the pointer is directly converted to a float4 type, read, and then unpacked into the register array reg. Similarly, the implementations with widths of 2 and 1 load a float2 or a single element, respectively. A static_assert ensures that the width does not exceed 4. This templated design avoids runtime branching while helping to keep the code concise.

[0061] This embodiment first determines the memory stride of the initial image tensor corresponding to the target image, determines the target vectorization width based on the memory stride and the memory data width of the target graphics processor, and determines the target data dimension of the target graphics processor. The target thread grid corresponding to the initial image tensor is then determined based on the target data dimension. The target thread grid is then used to slice the image tensor, and the target image tensor is written to memory. Through the above technical solution, the memory layout (stride) of the input tensor can be dynamically analyzed to adaptively determine the optimal vectorization width. By introducing a runtime dynamic decision-making mechanism, the optimal vectorization width is automatically selected based on the actual tensor memory layout to maximize memory access efficiency. Slicing is performed based on the data dimension when the target graphics processor outputs data, solving performance bottlenecks in tensor operations from two dimensions: memory access optimization (dynamic vector width) and parallel efficiency improvement (tensor slicing). This approach is both adaptable and efficient, suitable for high-performance computing scenarios with irregular memory layouts, fully utilizing GPU memory bandwidth and improving memory bandwidth utilization.

[0062] Based on the previous embodiment, we can see that this application can analyze the memory layout (step size) of the input tensor to adaptively determine the optimal vectorization width. Next, this embodiment will explain in detail the process of tensor slicing and memory writing. Figure 3 As shown, the embodiment of the present application provides a specific memory writing method, including:

[0063] Step S21 : determining a target data dimension of a target graphics processor, and determining a target thread grid corresponding to the initial image tensor according to the target data dimension; the target data dimension is the data dimension when the target graphics processor outputs data.

[0064] In this embodiment, when determining the target thread grid corresponding to the initial image tensor, the dimension size and thread block capacity corresponding to the target dimension in the target data dimension can be determined first, and the number of threads can be determined based on the dimension size and thread block capacity. A thread block can be constructed based on the number of threads, and then a two-dimensional thread grid can be constructed based on the thread block, and the two-dimensional thread grid can be used as the target thread grid. The target dimension is the last dimension of the target data dimension. Specifically, when constructing a two-dimensional thread grid based on the thread block, the number of threads can be determined as the first dimension of the thread grid, and the data batch corresponding to the initial image tensor can be determined as the second dimension of the thread grid. Then, based on the target number of threads and the data batch, a two-dimensional thread grid can be constructed based on the thread block. In other words, the key parameters for starting the tensor slicing operation in this embodiment include: vectorization width: dynamically determined by the decide_vector_width function; number of thread blocks: blocks = (last_dim + elements_per_block - 1) / elements_per_block, where elements_per_block = 256 × vec_width; grid configuration: two-dimensional grid (blocks, out_shape[0]). This approach uses slicing based on the last dimension of the target data. The number of thread blocks is calculated based on the size of the last dimension of the output tensor (256 threads per block x vectorization width), ensuring continuous address access within the thread block while maximizing parallelism through a two-dimensional grid. Furthermore, when organizing thread blocks based on a two-dimensional grid, the first dimension corresponds to the sliced ​​processing blocks, while the second dimension maintains the original batch dimension. Each thread block consists of 256 threads, which, combined with the vectorization width, enables efficient data parallel processing. Through compile-time generation and runtime dynamic dispatch of multiple kernel versions, a write-once, multiple-optimization approach is achieved.

[0065] Step S22: Slice the initial image tensor using the target thread grid to obtain a target image tensor, and write the target image tensor into the target memory.

[0066] In this embodiment, the correspondence between the initial image tensor and the threads in the two-dimensional thread grid can be determined, and the logical coordinates of the initial image tensor can be determined based on the correspondence. Then, the initial image tensor is read according to the logical coordinates, and the initial image tensor is sliced ​​using the target thread grid to obtain the target image tensor. When reading the initial image tensor according to the logical coordinates, the corresponding input address can be determined based on the non-contiguous memory step size and logical coordinates of the initial image tensor, and the corresponding initial image tensor can be read according to the input address. The specific process is as follows:

[0067] / / --------------------- Device kernel---------------------

[0068] template<typename T, int VEC_WIDTH>

[0069] __global__ void slice_kernel(const T* __restrict__ input, T* __restrict__ output,

[0070] const int* __restrict__ out_shape,

[0071] const int* __restrict__ out_strides,

[0072] const int* __restrict__ in_strides,

[0073] int dims) {

[0074] / / Calculate the global linear ID (by output tensor shape):

[0075] int tid = blockIdx.x * blockDim.x + threadIdx.x;

[0076] if (tid>= out_shape[0] * out_shape[1]) return;

[0077] / / Convert linear ID to logical coordinates:

[0078] int coord[2];

[0079] int remaining = tid;

[0080] for (int d = dims-1; d>= 0; --d) {

[0081] coord[d] = remaining % out_shape[d];

[0082] remaining / = out_shape[d];

[0083] };

[0084] / / Calculate input address (non-continuous mapping):

[0085] int in_offset = 0;

[0086] for (int d = 0; d <dims; ++d) {

[0087] in_offset += coord[d] * in_strides[d]; / / Input stride may be discontinuous

[0088] };

[0089] / / Vectorized load (non-contiguous addresses):

[0090] T reg[VEC_WIDTH];

[0091] / / Secondary sharding: ensuring access to consecutive addresses within a thread block

[0092] float reg[VEC_WIDTH];

[0093] vectorized_load<VEC_WIDTH> (input + in_offset, reg);

[0094] / / Example calculation: abs calculation:

[0095] #pragma unroll

[0096] for(int i=0; i <VEC_WIDTH; ++i){

[0097] reg[i] = fabs(reg[i], 0);

[0098] };

[0099] / / Calculate the output address (continuous writing):

[0100] int out_offset = coord[0] * out_strides[0]+ coord[1] * out_strides[1];

[0101] / / Vectorized storage:

[0102] vectorized_store<VEC_WIDTH> (output + out_offset, reg).

[0103] The above code includes:

[0104] Template parameters: T is the data type (such as float / int), VEC_WIDTH is the vectorization width (determines whether 1 / 2 / 4 elements are processed at a time);

[0105] __global__ modifier: identifies the function as a GPU kernel function, which is called by the host and executed in parallel on the GPU;

[0106] __restrict__ keyword: tells the compiler that the memory areas pointed to by pointers do not overlap, allowing the compiler to optimize memory access paths (such as avoiding cache conflicts);

[0107] tid = blockIdx.x * blockDim.x + threadIdx.x: In CUDA, blockIdx.x is the x-coordinate of the current block in the grid, blockDim.x is the number of threads per block (e.g., 256), and tid is a globally unique thread ID.

[0108] if (tid>= out_shape[0] * out_shape[1]) return: Ensure that the thread ID does not exceed the total number of elements in the first two dimensions of the output tensor (such as rows × columns of a two-dimensional tensor) to avoid out-of-bounds access;

[0109] One-dimensional ID conversion: Example: If out_shape = [2,3] (2 rows and 3 columns), tid = 5, then:

[0110] d=1 (column dimension): coord[1]=5%3=2, remaining=5 / 3=1;

[0111] d=0 (row dimension): coord[0]=1%2=1, the final coordinate is (1, 2) (row 1, column 2);

[0112] in_offset += coord[d] * in_strides[d]: The memory layout of the input tensor may be non-contiguous (such as a sliced ​​tensor). in_strides[d] represents the stride of the d-th dimension of the input tensor. For example, if the input stride is [8, 2] and the coordinates are (1, 2), then in_offset = 1×8+2×2=12, which means that the data is read from the base address offset by 12 elements.

[0113] vectorized_load<VEC_WIDTH> : Read data in batches from non-contiguous addresses according to VEC_WIDTH (for example, read 4 elements at a time when VEC_WIDTH=4), optimizing memory bandwidth utilization and reducing the number of transactions;

[0114] #pragma unroll:

[0115] Forces the compiler to unroll loops to avoid loop judgment overhead. For example, when VEC_WIDTH=4, 4 lines of independent fabs calculation code are generated, improving execution efficiency.

[0116] out_offset calculation: Use the output step size out_strides (usually a continuous layout) to ensure that the output address is continuous, making it easier for the GPU to merge write operations from multiple threads;

[0117] vectorized_store<VEC_WIDTH> : Write results in batches according to the vectorization width, echoing the load operation to further improve memory write efficiency.

[0118] Based on the above technical solution, this embodiment dynamically adapts VEC_WIDTH based on the input tensor stride, avoiding memory access misalignment caused by fixed vectorization. Furthermore, through in_strides, it supports irregular memory layouts, directly calculating the target address and reducing intermediate conversion overhead. Furthermore, it uses a continuous stride during output to ensure memory write merging and maximize GPU memory bandwidth utilization. This effectively combines dynamic vectorization with non-contiguous memory access. Through the template parameter VEC_WIDTH, the kernel supports runtime adaptive vectorization widths (1 / 2 / 4), optimizing memory loading efficiency based on the stride characteristics (in_strides) of the input tensor. Furthermore, by converting the global thread ID (tid) into multi-dimensional logical coordinates (coord) and calculating the input address (in_offset) based on non-contiguous strides, memory access is directly linked to data calculation, reducing intermediate overhead. Furthermore, by tightly integrating in_offset calculation with data reads, each memory access is ensured to be an efficient vectorized operation, avoiding the performance loss caused by separating address calculation from data reads. In a specific embodiment, the input address and output address correspond to the non-contiguous memory layout and continuous memory layout in tensor operations, respectively. The input address (in_offset) corresponds to the memory layout. The non-contiguous memory address reflects the irregular offset of the data in the original storage. The input address is used to read data from the non-contiguous memory. The output address (out_offset) corresponds to the continuous memory address of the output tensor, ensuring memory continuity when writing data. By writing the processed data to continuous memory, the GPU's memory merge write performance is maximized. In this way, by non-contiguous reading of the input address and continuous writing of the output address, efficient conversion of non-contiguous tensors to continuous tensors is achieved.

[0119] Accordingly, when writing the target image tensor to the target memory, the continuous step size of the target image tensor can be determined, and the data output address corresponding to the target image tensor can be determined based on the continuous step size. Then, the target image tensor can be written to the continuous target memory according to the data output address. Specifically, based on the target vectorization width and the data output address, the target image tensor can be cyclically written to the target memory using loop unrolling instructions. In this way, combined with the tensor input process, this embodiment implements a dynamically vectorized tensor slicing operation kernel that supports non-continuous memory access and efficient calculation of multi-dimensional tensors. In a specific embodiment, logical coordinates (coord) can be calculated using the thread ID (tid) to convert global linear indices into coordinates of a multidimensional tensor. The input address (in_offset) is calculated based on the non-contiguous stride (in_strides) of the input tensor, supporting irregular memory layouts. Data is then read from the non-contiguous address into a register array (reg) using vectorized loads (vectorized_load), and the loop is unrolled using #pragma unroll to perform absolute value calculations (fabs). Finally, the output address (out_offset) is calculated based on the continuous stride (out_strides) of the output tensor, and the result is written to contiguous memory using vectorized stores (vectorized_store). In this way, this embodiment dynamically adapts the vectorization width (1 / 2 / 4) through the template parameter VEC_WIDTH, combined with __restrict__ pointer and memory access optimizations to maximize GPU parallel computing efficiency. Furthermore, #pragma unroll loop unrolling and vectorized stores (vectorized_store) ensure output address continuity while maintaining computational efficiency to maximize memory merge write performance. The optimal task partitioning scheme can be calculated based on the output tensor dimension and hardware parameters. Through special optimizations for non-contiguous access, including address calculation and register usage optimization, a balance is achieved between hardware adaptability, memory access efficiency, and code versatility. Furthermore, after the loop is expanded using #pragma unroll, all threads execute the same instruction sequence, fully utilizing the CUDA core, facilitating the GPU's instruction scheduler to issue instructions in batches, and improving CUDA core utilization. This is particularly suitable for the vectorized computing scenarios in this embodiment (such as processing 4 elements when VEC_WIDTH=4), making register access and computing operations more compact and reducing loop control overhead.

[0120] Based on the above technical solution, Figure 4 As shown, based on the following code, this embodiment achieves high efficiency and universality of CUDA tensor operations through dynamic decision-making + hierarchical parallelism + collaborative design of multi-version kernels.

[0121] / / --------------------- Startup code---------------------

[0122] template<typename T>

[0123] void launch_slice(const T* d_input, T* d_output,

[0124] const std::vector <int>&out_shape,

[0125] const std::vector <int>&in_strides,

[0126] const std::vector <int>&out_strides) {

[0127] / / Dynamically determine vector width:

[0128] const int vec_width = decide_vector_width <t>(in_strides, sizeof(T));

[0129] / / Sharding strategy (sharding by the last dimension of the output tensor):

[0130] const int last_dim = out_shape.back();

[0131] const int threads_per_block = 256;

[0132] const int elements_per_block = threads_per_block * vec_width;

[0133] const int blocks = (last_dim + elements_per_block - 1) / elements_per_block;

[0134] dim3 grid(blocks, out_shape[0]); / / The first dimension is fully preserved, and the second dimension is sliced

[0135] / / Dynamic dispatch:

[0136] switch(vec_width) {

[0137] case 4:

[0138] slice_kernel<T,4> <<<grid, threads_per_block> >>(

[0139] d_input, d_output,

[0140] out_shape.data(), out_strides.data(),

[0141] in_strides.data(), out_shape.size() );

[0143] break;

[0144] case 2:

[0145] slice_kernel<T,2> <<<grid, threads_per_block> >>(...);

[0146] break;

[0147] default:

[0148] slice_kernel<T,1> <<<grid, threads_per_block> >>(...);

[0149] }

[0150] }.

[0151] These include: template<typename T> : Define generic templates so that functions can handle any data type (such as float, int) and achieve code reuse;

[0152] void launch_slice(...): Launches the GPU kernel to perform tensor slicing operations. Parameters include input tensors / output tensors and stride.

[0153] decide_vector_width <t>(in_strides, sizeof(T)): Calls the host-side function to calculate the optimal vectorization width: traverses the input stride to find the minimum positive stride (min_stride), calculates the byte stride (byte_stride = min_stride × data type size); and makes a decision based on the 128-byte transaction width of the Ampere architecture;

[0154] last_dim = out_shape.back(): Get the last dimension of the output tensor (such as the number of columns of a two-dimensional tensor) as the base dimension of the sharding;

[0155] threads_per_block = 256: Each thread block contains 256 threads. This number of threads is an efficient configuration for CUDA hardware scheduling, ensuring full utilization of resources within the thread block);

[0156] elements_per_block = 256 * vec_width: The number of elements processed per block = the number of threads × vectorization width. For example, when vec_width=4, 1024 elements are processed per block.

[0157] blocks = (last_dim + elements_per_block - 1) / elements_per_block: round up to the nearest integer to calculate the number of thread blocks. For example, when last_dim = 1000 and elements_per_block = 1024, blocks = 1; when last_dim = 1025, blocks = 2 to ensure that all elements are processed.

[0158] dim3 grid(blocks, out_shape[0]): Build a two-dimensional thread grid: the first dimension blocks: the number of blocks sliced ​​by the last dimension; the second dimension out_shape[0]: retain the first dimension of the output tensor to avoid cross-batch data interference;

[0159] switch(vec_width): selects the corresponding kernel version according to the vectorization width to achieve runtime adaptation: case 4: calls slice_kernel<T,4> , using 4-element vectorization, suitable for scenarios with high memory continuity; case 2: calling slice_kernel<T,2> , 2-element vectorization, balanced memory access efficiency; default: call slice_kernel<T,1> , scalar access, suitable for highly non-contiguous memory scenarios;

[0160] <<<grid, threads_per_block> >>: CUDA kernel launch syntax: grid: thread grid dimension, determines the layout of the global thread block; threads_per_block: number of threads per block, determines the parallelism within the block;

[0161] slice_kernel parameters: pass input / output pointers, shape and stride arrays so that the kernel can calculate addresses and process data correctly.

[0162] Based on the above code, this embodiment implements an adaptively optimized GPU tensor slicing operation launcher, whose core function is to dynamically select the optimal execution strategy to complete efficient data copy operations based on the memory layout characteristics of the input tensor. The launcher automatically determines the vectorization width and divides the thread blocks by analyzing the dimension information and span parameters of the input / output tensors, and finally calls the specialized CUDA kernel function to achieve high-performance memory access. Specifically including: dynamic vector width decision: analyze the input span (in_strides) and data type size through the decide_vector_width function, and intelligently select a vectorization width of 1 / 2 / 4 to maximize the use of the GPU's SIMD instruction set; intelligent slicing strategy: calculate the optimal number of thread blocks based on the last dimension size of the output tensor and the thread block capacity (256 threads × vectorization width); dynamic kernel function dispatch: automatically select the kernel function template (slice_kernel) that matches the vectorization width through the switch-case structure<T,4 / 2 / 1> In this way, by adopting a layered sharding strategy, we maximize parallelism while maintaining data locality, and through multi-version kernel dynamic dispatch technology, we achieve compile-time optimization and runtime flexibility. In particular, we have deeply optimized non-continuous tensor operations, which can significantly improve the execution efficiency of various tensor slicing operations.

[0163] Through the above technical solutions, this embodiment achieves the coordinated optimization of computing mode and memory access mode while maintaining the simplicity of the interface by combining compile-time polymorphism with runtime decision-making, fully tapping the parallel potential of modern GPU architectures and providing a unified high-performance tensor operation interface for the upper layer. Combined with the previous embodiment, the performance bottleneck in tensor operations is solved from the two dimensions of memory access optimization (dynamic vector width) and parallel efficiency improvement (sharding and dynamic dispatch). The optimized vectorized access mode fully utilizes the GPU memory bandwidth, improves the memory bandwidth utilization, and based on intelligent task division, improves the parallel calculation process and improves data throughput. At the same time, the adaptability is significantly enhanced, and it can automatically adapt to different hardware architectures and tensor layouts, and can achieve near-optimal performance without manual intervention. In this way, this embodiment has both adaptability and efficiency, and can be effectively applied to high-performance computing scenarios under irregular memory layouts.

[0164] In a specific embodiment, the above-mentioned CUDA-based dynamic vectorization and intelligent slicing can be applied to facial image processing scenarios, specifically targeting core requirements such as local feature extraction and multi-region parallel processing of facial images, to solve the efficiency bottleneck caused by non-continuous memory access. It is understandable that in facial image processing, it is necessary to crop multiple non-continuous local areas (such as eyes, nose, mouth, etc.) from the original image (or feature map) and extract image features, and different features are distributed non-continuously in memory. This embodiment can first adapt the memory layout of image tensors corresponding to different features, and automatically select the optimal vectorization width. At the same time, it divides the thread blocks according to the image feature dimensions, and calls the determined optimal kernel according to the image feature layout to adapt to the dynamically changing facial areas.

[0165] For example, for the cropping coordinates corresponding to the input facial image tensor, the stride in the original image memory is calculated. Then, the decide_vector_width function in the above embodiment is called to determine the vectorization width based on the determined minimum positive stride (min_stride) and the facial image pixel data type (e.g., uint8 or float32). The tensor is then sliced ​​based on the last dimension of the facial image output tensor to be stored in memory, preserving the integrity of the batch_size. It is also understood that this embodiment can precompile corresponding kernel versions based on common data types and vectorization widths (1 / 2 / 4) for facial image features. Then, when the function is started, the corresponding kernel is directly called through a switch-case based on the memory stride and data type of the input image features. This solution can be further applied to scenarios such as facial image generation. Through dynamic vectorization and intelligent slicing, it addresses parallel computing efficiency issues in non-contiguous memory layouts, providing GPU-based low-level acceleration support for real-time facial processing systems (such as security monitoring and identity verification).

[0166] Through the description of the above implementation methods, those skilled in the art can clearly understand that the method according to the above embodiment can be implemented by means of software plus the necessary general hardware platform, and of course it can also be implemented by hardware, but in many cases the former is a better implementation method.

[0167] An embodiment of the present application further provides an electronic device, comprising a memory and a processor, wherein the memory stores a computer program, and the processor is configured to run the computer program to execute the steps in any one of the above-mentioned memory writing method embodiments.

[0168] An embodiment of the present application further provides a computer-readable storage medium, in which a computer program is stored, wherein the computer program is configured to execute the steps of any of the above-mentioned memory writing method embodiments when run.

[0169] In an exemplary embodiment, the computer-readable storage medium may include, but is not limited to, various media that can store computer programs, such as a USB flash drive, a read-only memory (ROM), a random access memory (RAM), a mobile hard disk, a magnetic disk, or an optical disk.

[0170] An embodiment of the present application further provides a computer program product, which includes a computer program. When the computer program is executed by a processor, the steps of any of the above-mentioned memory writing method embodiments are implemented.

[0171] An embodiment of the present application also provides another computer program product, including a non-volatile computer-readable storage medium, wherein the non-volatile computer-readable storage medium stores a computer program, and when the computer program is executed by a processor, the steps of any of the above-mentioned memory writing method embodiments are implemented.

[0172] Professionals may further appreciate that the units and algorithm steps of each example described in conjunction with the embodiments disclosed herein can be implemented in electronic hardware, computer software, or a combination of the two. In order to clearly illustrate the interchangeability of hardware and software, the above description has generally described the components and steps of each example according to their functions. Whether these functions are performed in hardware or software depends on the specific application and design constraints of the technical solution. Professionals and technicians may use different methods to implement the described functions for each specific application, but such implementation should not be considered beyond the scope of this application.

[0173] The above is a detailed introduction to a memory writing method, device, storage medium and program product provided by the present application. Specific examples are used herein to illustrate the principles and implementation methods of the present application. The description of the above embodiments is only used to help understand the method and core ideas of the present application. It should be pointed out that for ordinary technicians in this technical field, without departing from the principles of the present application, several improvements and modifications can be made to the present application, and these improvements and modifications also fall within the scope of protection of the claims of the present application.< / t> < / t> < / int> < / int> < / int>

Claims

1. A memory writing method, characterized in that: Applicable to target graphics processors, including: Acquire a target image, determine an initial image tensor corresponding to the target image, and determine a memory step size of the initial image tensor; wherein the initial image tensor is a discontinuous image tensor; Determining a target vectorization width corresponding to the initial image tensor according to the memory stride and the memory data width of the target graphics processor; Determining a target data dimension of the target graphics processor, and determining a target thread grid corresponding to the initial image tensor based on the target data dimension; the target data dimension is a data dimension when the target graphics processor outputs data; The initial image tensor is sliced ​​using the target thread grid to obtain a target image tensor, and the target image tensor is written into a target memory.

2. The memory writing method according to claim 1, wherein: The determining, based on the memory stride and the memory data width of the target graphics processor, a target vectorization width corresponding to the initial image tensor, includes: Determining the memory step size of the initial image tensor, and determining a minimum step size of the initial image tensor according to the memory step size; the minimum step size is a positive number; The memory data width of the target graphics processor is determined, and the target vectorization width corresponding to the initial image tensor is determined according to the minimum step size and the memory data width.

3. The memory writing method according to claim 2, characterized in that: The determining, according to the minimum step size and the memory data width, the target vectorization width corresponding to the initial image tensor includes: Determining a byte step size corresponding to the initial image tensor based on the minimum step size and the number of bytes of the initial image tensor; Based on the byte stride and the memory data width, the target vectorization width corresponding to the initial image tensor is determined.

4. The memory writing method according to claim 3, characterized in that: The determining, based on the byte stride and the memory data width, the target vectorization width corresponding to the initial image tensor includes: Dividing the byte stride by the memory data width; If the memory data width is divisible by the byte stride, determining the target vectorization width to be four; If the memory data width is not divisible by the byte stride, and the memory data width is divisible by twice the byte stride, determining the target vectorization width to be two; If the memory data width is not divisible by the byte stride and the memory data width is not divisible by twice the byte stride, the target vectorization width is determined to be one.

5. The memory writing method according to claim 1, wherein: Before slicing the initial image tensor using the target thread grid to obtain the target image tensor, the method further includes: A kernel function corresponding to each of the target vectorization widths is constructed so as to directly call the kernel function corresponding to the target vectorization width to slice the initial image tensor.

6. The memory writing method according to claim 5, characterized in that: Before slicing the initial image tensor using the target thread grid, the method further includes: Converting an initial pointer to a target pointer of a corresponding data type according to the target vectorization width of the current initial image tensor; The initial image tensor is read using the pointer to slice the initial image tensor using the target thread grid.

7. The memory writing method according to any one of claims 1 to 6, characterized in that: The determining, according to the target data dimension, a target thread grid corresponding to the initial image tensor includes: Determining a dimension size and a thread block capacity corresponding to a target dimension in the target data dimension; the target dimension is the last dimension of the target data dimension; Determining the number of threads based on the dimension size and the thread block capacity, and constructing a thread block according to the number of threads; A two-dimensional thread grid is constructed based on the thread block, and the two-dimensional thread grid is used as the target thread grid.

8. The memory writing method according to claim 7, characterized in that: The constructing of a two-dimensional thread grid based on the thread block includes: Determine the number of threads as a first dimension of a thread grid, and determine the data batch corresponding to the initial image tensor as a second dimension of the thread grid; The two-dimensional thread grid is constructed based on the thread blocks according to the target number of threads and the data batch.

9. The memory writing method according to claim 8, characterized in that: The slicing the initial image tensor by using the target thread grid to obtain a target image tensor includes: Determining a correspondence between the initial image tensor and threads in the two-dimensional thread grid, and determining logical coordinates of the initial image tensor based on the correspondence; The initial image tensor is read according to the logical coordinates, and the initial image tensor is sliced ​​using the target thread grid to obtain the target image tensor.

10. The memory writing method according to claim 9, characterized in that: The reading the initial image tensor according to the logical coordinates includes: Determining a corresponding input address according to the memory step size and the logical coordinates of the initial image tensor, wherein the memory step size is a discontinuous step size; The corresponding initial image tensor is read according to the input address.

11. The memory writing method according to claim 10, characterized in that: Writing the target image tensor into the target memory includes: Determining a continuous step length of the target image tensor, and determining a data output address corresponding to the target image tensor based on the continuous step length; The target image tensor is written into the target memory according to the data output address; the target memory is a continuous memory.

12. The memory writing method according to claim 11, characterized in that: Writing the target image tensor into the target memory according to the data output address includes: Based on the target vectorization width and the data output address, the target image tensor is cyclically written into the target memory using a loop unrolling instruction.

13. An electronic device, characterized in that: include: Memory for storing computer programs; A processor, configured to implement the steps of the memory writing method according to any one of claims 1 to 12 when executing the computer program.

14. A computer-readable storage medium, characterized in that The computer-readable storage medium stores a computer program, wherein the computer program, when executed by a processor, implements the steps of the memory writing method according to any one of claims 1 to 12.

15. A computer program product comprising a computer program, characterized in that When the computer program is executed by a processor, the steps of the memory writing method according to any one of claims 1 to 12 are implemented.

Images