Multi-scalar multiplication acceleration method based on resource pre-estimation and pre-calculation strategy

Abstract

The invention discloses a multi-scalar multiplication acceleration method based on resource estimation and a pre-calculation strategy, and provides a systematic solution for the problems of insufficient resource estimation, pre-calculation factor stiffness and low sorting efficiency of a Pippenger algorithm in zero-knowledge proof. A multi-resolution point multiplication table is dynamically generated through a hierarchical displacement pre-calculation strategy, and global memory access delay is remarkably reduced by combining memory layout optimization and asynchronous flow task scheduling of window inner barrel continuous storage. A dynamic resource adaptation mechanism is designed, thread block topology, grid division and sorting algorithm selection are adjusted in real time based on GPU hardware features, and load balancing and cache utilization rate maximization are achieved. Iterative reduction and double temporary bucket strategies are introduced, and data scale is compressed and boundary processing is optimized through multi-round reduction. In the large-scale MSM operation in the block chain and privacy computing field, the computing efficiency and the resource utilization rate can be remarkably improved, and the method is suitable for an elliptic curve cryptography acceleration task.

Description

Technical Field

[0001] The present invention relates to an acceleration method based on GPU parallel processing and hardware feature optimization for efficient calculation of multi-scalar multiplication (MSM) in zero-knowledge proof, so as to solve the low efficiency problem of the current Pippenger algorithm caused by insufficient resource estimation and insufficient adjustment of pre-calculation factors. Specifically, the present invention proposes a multi-scalar multiplication parallel optimization process combined with GPU hardware characteristics, and constructs a unified scheduling model for heterogeneous computing resources through resource estimation, dynamic task allocation and the implementation mechanism of hierarchical displacement pre-calculation factors. The method realizes intelligent prediction and adaptive adjustment of the sorting scale of the bucket reduction stage, and completes efficient decomposition and allocation of data preprocessing and computing tasks through a multi-level pipeline architecture. The multi-scalar multiplication acceleration method based on resource estimation and pre-calculation strategy of the present invention can be widely used in curve operation acceleration tasks in the field of cryptography, supports load balancing and cross-hardware collaborative execution of large-scale parallel computing, and significantly improves the computing efficiency and resource utilization of zero-knowledge proof algorithms. Background Art

[0002] With the widespread application of zero-knowledge proof (ZKP) technology in blockchain, privacy-preserving computing, and other fields, multi-scalar multiplication (MSM), one of its core computational tasks, has placed higher demands on algorithm efficiency and hardware resource utilization. The Pippenger algorithm, a classic MSM optimization algorithm, significantly reduces computational complexity by decomposing a scalar into multiple windows and performing bucket accumulation operations, making it a mainstream choice for current zero-knowledge proof implementations. However, in practical applications, especially in GPU-based parallel implementations, the Pippenger algorithm still faces numerous performance bottlenecks.

[0003] First, the implementation of the Pippenger algorithm on the GPU has the problem of insufficient resource estimation. This is manifested in the unreasonable allocation of GPU video memory and computing resources, which may cause video memory overflow or low computing efficiency when processing large-scale data. The root cause is: (1) The problem of uneven bucket allocation. The Pippenger algorithm decomposes the MSM into multiple bucket accumulation operations. Each bucket requires memory allocation, and the number and size of buckets depend on the data scale and window size. When processing large-scale data, the number of buckets may increase sharply, causing video memory pressure and a decrease in computational parallelism. (2) Improper thread configuration. The GPU's thread block and grid configuration needs to be optimized based on the data scale and the number of buckets. If the thread allocation is too sparse or dense, it may lead to low thread utilization, which in turn affects the overall computing performance.

[0004] Secondly, the pre-computation factor is not dynamically adjusted for multi-scalar operations of varying scales, limiting the algorithm's adaptability. The pre-computation factor is a key parameter in the Pippenger algorithm, determining the number of buckets and computational complexity. However, current implementations typically use a fixed pre-computation factor that cannot be dynamically adjusted based on data size. This can lead to resource waste on small-scale data and potentially insufficient video memory or decreased computational efficiency on large-scale data.

[0005] Furthermore, insufficient sorting efficiency further exacerbates the performance bottleneck of the Pippenger algorithm on GPUs. The bucket reduction stage requires sorting the input data to ensure that data in the same bucket can be efficiently merged. However, traditional sorting methods exhibit significant limitations when faced with varying data sizes. For example, a single sorting algorithm may perform well on small-scale data, but its efficiency may decrease significantly on large-scale data. Conversely, a sorting algorithm suitable for large-scale data may incur additional overhead on small-scale data. This inefficiency directly impacts the degree of parallelization and overall computational performance of the bucket reduction stage.

[0006] In summary, the existing Pippenger algorithm implementation on GPUs faces three major problems: (1) Insufficient resource estimation leads to irrational memory allocation and low computational parallelism; (2) The pre-computation factor is not dynamically adjusted, making it unable to adapt to the needs of multi-scalar operations of varying scales; and (3) Inefficient sorting limits the parallelization potential of the bucket reduction stage. These problems severely restrict the performance of the Pippenger algorithm in heterogeneous computing environments, and a comprehensive optimization solution is urgently needed to improve its computational efficiency and resource utilization. Summary of the Invention

[0007] This paper proposes a multi-scalar multiplication acceleration method based on resource estimation and pre-computation strategy. It is an efficient bucket sorting and reduction optimization method for GPU-accelerated multi-scalar multiplication (MSM) operations. It covers multiple aspects such as memory layout optimization, computational scheduling optimization, dynamic resource adaptation, pipeline parallel processing and hardware feature adaptation, aiming to significantly improve the execution efficiency and resource utilization of large-scale data processing tasks.

[0008] First, during the bucket sorting phase, we optimized the core data structure. By introducing the bucket_sizes array to store the capacity of each bucket and combining it with the thrust::sort_by_key function, we use the bucket capacity as the key and the bucket index array single_bucket_indices as the value to perform descending sorting, thereby achieving an orderly arrangement of bucket indices by capacity. The sorted results provide a basis for subsequent resource scheduling and significantly improve local access efficiency. To further optimize memory access patterns, we designed a memory layout for continuous storage of buckets within a window. Specifically, buckets within the same window are tightly arranged in physical memory using the w*2c+bucket_id method, effectively utilizing the cache prefetch mechanism and reducing global memory access latency.

[0009] To address the bottleneck problems existing in atomic operations and dynamic memory allocation, a prefix sum method is used to pre-calculate the memory offset of each bucket to avoid resource competition during dynamic allocation and ensure that large buckets can be efficiently accessed through continuous memory blocks. At the same time, a multi-stream parallel strategy is introduced, and independent CUDA streams (such as stream_large_buckets) are set up to handle large bucket computing tasks. The calculations of large and small buckets are reasonably coordinated through the cudaStreamWaitEvent synchronization mechanism to avoid idle resources and improve overall parallelism. At the thread scheduling level, the number of thread blocks (NUM_BLOCKS) and the number of threads per block (NUM_THREADS) are dynamically adjusted according to the real-time load to maximize the occupancy rate of the streaming multiprocessor (SM) and achieve a precise match between computing resources and task load.

[0010] During the reduction phase, the present invention proposes an iterative reduction and dual temporary buckets strategy. This strategy progressively compresses the data size through multiple reduction rounds, halving the number of data bits with each round. The temporary buckets used for writing data are dynamically switched based on the parity bit to handle boundary conditions and optimize storage access. Furthermore, the combination of an asynchronous computation mechanism enables pipelined processing of the reduction process, significantly improving data processing throughput.

[0011] To further enhance overall system performance, this paper constructs a multi-level memory management mechanism. Based on object size, it employs fixed block pre-allocation, a buddy algorithm, and physical page mapping strategies to dynamically manage small-scale metadata, medium-sized buffers, and large-scale data blocks, minimizing memory fragmentation and maximizing continuity. Furthermore, the radix sort bit width is dynamically calculated based on the GPU L2 cache capacity, partitioning the data blocks to fit the cache and enhancing data locality.

[0012] In terms of computing resource management, this paper introduces a dynamic resource scheduling and hybrid sorting selection mechanism. By analyzing data features in real time, it dynamically switches between bitonic sorting, radix sorting, and library function sorting, and determines the decision threshold based on offline modeling, thereby maintaining near-optimal time complexity. Furthermore, by splitting the computational process into three phases: radix partitioning, sub-block sorting, and global merging, and employing multi-stream parallelism and event synchronization techniques, it effectively achieves overlapping computation and communication.

[0013] For ultra-large-scale bucket data processing, the present invention proposes a mechanism to isolate large bucket and small bucket computing channels, and introduces space filling curves (such as Z-Order rearrangement) to optimize memory access patterns, thereby further improving cache hit rates and reducing latency in two-dimensional locality.

[0014] During the pre-calculation phase, the present invention introduces a hierarchical displacement strategy to construct a multi-resolution pre-calculation table. This hierarchical displacement operation dynamically adjusts the displacement and generates a pre-calculation table, enabling rapid table lookup to obtain the dot product results corresponding to the scalar blocks during actual calculations, significantly shortening the computational path. The pre-calculation table utilizes a continuous storage layout, combining the optimized design of each layer's displacement with the graphics memory layout to achieve coalesced memory access, improving bandwidth utilization.

[0015] Furthermore, to adapt to different GPU hardware platforms, this paper proposes a dynamic architecture adaptation mechanism based on multidimensional feature modeling. Based on the CUDA core ratio, memory bandwidth, streaming multiprocessor capacity, and cache characteristics, a runtime decision tree is established to dynamically adjust thread block configuration, meshing, and warp-level computing strategies based on task load, ensuring optimal resource utilization under different hardware environments. Specifically, this includes technologies such as dynamic thread block topology optimization, meshing adaptation, and warp-level double-buffered pipelining, effectively improving the robustness and efficiency of cross-platform execution.

[0016] In summary, the present invention significantly improves the performance of GPU-accelerated multi-scalar multiplication (MSM) operations in large-scale data processing scenarios through a systematic design from memory layout, computational scheduling, reduction optimization, resource adaptation to pipeline parallel architecture. It has broad application prospects and high industrial value.

[0017] A multi-scalar multiplication (MSM) optimization system based on GPU acceleration, characterized by comprising:

[0018] A resource estimation module, configured to dynamically analyze GPU memory capacity, number of streaming multiprocessors, and cache characteristics to generate a resource allocation strategy for multi-scalar multiplication tasks;

[0019] Dynamic task allocation module, configured to dynamically adjust thread block size and grid division based on window width, number of buckets, and hardware resource characteristics to achieve computational load balancing;

[0020] a layered displacement pre-calculation module configured to construct a multi-resolution pre-calculation table, dynamically adjust the scalar block coefficients by the layered displacement factors, and generate pre-calculated dot product results for continuous storage;

[0021] The bucket sort optimization module is configured to adopt a memory layout with continuous storage of buckets within a window, combining prefix and offset calculation with an asynchronous stream scheduling mechanism to eliminate atomic operation contention;

[0022] The iterative reduction module is configured to compress the data size through multiple rounds of reduction and use a dual temporary bucket strategy and parity bit adaptive writing technology to handle boundary conditions;

[0023] The hardware adaptation module is configured to dynamically select the sorting algorithm and thread block topology based on the GPU architecture characteristics to optimize cache hit rate and computing resource utilization.

[0024] The mathematical expression of the layered displacement pre-calculation module is:

[0025]

[0026] Among them, P k,i is the hierarchical displacement, k and i are preset values, j is the sequence number, c is the window width, δ is the dynamically adjusted block displacement, P is the elliptic curve base point, and a hierarchical pre-calculation table is generated to support fast point multiplication table lookup operations.

[0027] The bucket sorting optimization module is further configured as follows:

[0028] Sort the bucket indexes in descending order of capacity using `thrust::sort_by_key` to generate an ordered sequence of `single_bucket_indices`;

[0029] Adopting the physical memory continuous layout strategy, the bucket index of window w is mapped to w·2 c + bucket_id, improves spatial locality;

[0030] Use independent CUDA streams to separate large and small bucket computing tasks, and use `cudaStreamWaitEvent` to achieve asynchronous stream synchronization.

[0031] The iterative reduction module is further configured to:

[0032] Define temporary buckets `temp_buckets1` and `temp_buckets2` to handle the even and odd bit reduction phases respectively;

[0033] Dynamically switch the write target bucket according to the reduction round, and control the parity data distribution through the `is_odd_c` flag;

[0034] A multi-stage pipeline architecture is adopted to decompose the reduction operation into three stages: cardinality partitioning, sub-block sorting and global merging.

[0035] The hardware adaptation module is further configured as:

[0036] Build a four-dimensional feature vector model to dynamically collect CUDA core ratio, memory bandwidth, SM capacity and cache strategy;

[0037] Dynamically switch between Bitonic sorting, radix sorting, and library function sorting through a runtime decision tree;

[0038] The Z-Order space filling curve is used to optimize the two-dimensional memory access pattern and improve the L2 cache hit rate.

[0039] It also includes a multi-level memory management module, configured as:

[0040] For data sizes smaller than 4KB, a fixed block pre-allocation strategy is used;

[0041] Applying the buddy algorithm to optimize memory continuity for 4KB-1MB buffers;

[0042] Physical page mapping optimization is implemented for data blocks larger than 1MB, and data layout is guided by combining the memory access pattern.

[0043] A GPU-accelerated multi-scalar multiplication (MSM) optimization method, comprising:

[0044] Dynamically model GPU hardware characteristics to generate a four-dimensional parameter vector including memory bandwidth, number of SMs, and cache levels;

[0045] Generate a multi-resolution pre-calculation table based on the window width \(c\) and the layer displacement factor \(\delta\) to achieve fast table lookup in scalar blocks;

[0046] Column-major memory layout is used to store elliptic curve point coordinates, allowing threads in the same thread block to access continuous memory areas;

[0047] The asynchronous stream mechanism is used to separate large and small bucket computing tasks, and prefix and pre-computation are used to eliminate dynamic memory allocation competition.

[0048] Perform multiple rounds of iterative reduction, halving the number of data bits in each round and switching the temporary bucket write path based on the parity bit flag;

[0049] Dynamically adjust thread block dimensions and grid partitioning strategies to optimize SM occupancy and computing throughput based on real-time load.

[0050] The pre-calculation table generation step includes:

[0051] The elliptic curve base point P is 2 k·c Perform exponential displacement to generate a hierarchical displacement point set;

[0052] Dynamically divide the displacement point set of each layer into blocks, and the block coefficient \(\delta\) is adaptively adjusted according to the current GPU memory capacity;

[0053] The pre-computed results are stored continuously according to the window index to form a linear memory structure that supports coalesced access.

[0054] The dynamically adjusting thread block dimension includes:

[0055] According to the number of resident thread warps supported by SM, according to the formula blockDim x =min(1024, SM_capacity·32) adjust thread block width;

[0056] Calculate the grid dimensions using the load balancing model: Where N is the total number of tasks, and load_factor is the dynamic adjustment coefficient based on the number of SMs.

[0057] Also includes cross-architecture optimization steps:

[0058] Enables a warp-level double-buffered pipeline on the Volta architecture to hide register access latency;

[0059] The radix sort bit width is dynamically divided according to the L2 cache capacity to generate a data block structure that adapts to the cache line. BRIEF DESCRIPTION OF THE DRAWINGS

[0060] Figure 1 This is a delay change diagram for the BLS12_381 and BN254 curves under different pre-calculation factors.

[0061] Figure 2 This is a histogram comparing the before and after implementation of the layered displacement pre-calculation factor mechanism. Specific implementation methods

[0062] Combined with the attached figure ( Figure 1-Figure 2 ) is attached to further illustrate the present invention in detail.

[0063] Specific implementation methods

[0064] (1) Data arrangement in the Pippenger algorithm

[0065] In GPU-accelerated MSM operations, data layout has a decisive impact on performance. Its core design goal is to maximize GPU parallel computing efficiency through memory continuity, locality, and load balancing optimization. The following analysis focuses on data structure, storage strategy, and performance optimization:

[0066] Each elliptic curve point in an affine coordinate system, such as a BN254 curve point, consists of two components, x and y. Each component is split into multiple 64-bit unsigned integers, such as a 256-bit coordinate is split into four 64-bit blocks. Column-major order layout: The x and y components of all points are stored consecutively in columns, specifically labeled P0, P1, ..., P n For each point, its coordinate data is recorded in an orderly and repeated structure. Specifically, P0x[0], P1x[0], ..., P n x[0], uses index values ​​to distinguish coordinates in different states. When the index is 0, the horizontal coordinates are recorded in sequence, and the vertical coordinates are recorded as P0y[0], P1y[0], ..., P n y[0]. When the index value becomes 1, there is also a corresponding horizontal coordinate sequence P0x[1], P1x[1], ..., P n x[1] and the ordinate sequence P0y[1], P1y[1], ..., P n y[1]. This data organization method clearly presents the location information of each point under different conditions.

[0067] Each thread is responsible for processing a specific data block for a point. For example, thread 0 processes the x[0] and y[0] blocks for all points. Multiple threads within the same thread block collectively read all blocks for a point, ensuring continuous memory access without strides, meeting GPU coalescing requirements. Column-major layout offers advantages for memory access, eliminating random accesses and avoiding the problem of different threads accessing non-contiguous memory in traditional row-major order. It also reduces memory bandwidth waste and increases resource utilization efficiency.

[0068] (2) Bucket sorting strategy

[0069] Optimization strategies for bucket sorting can be developed from two perspectives: memory layout and computational scheduling. First, we need to clarify the role of the core data structure: the bucket_sizes array stores the capacity of each bucket (i.e., the number of points within the bucket), where bucket_sizes[i] represents the number of points in the i-th bucket. By calling the thrust::sort_by_key function, the buckets are sorted in descending order (using the thrust::greater comparator) using the values ​​in bucket_sizes as keys and single_bucket_indices (an array of bucket indices) as values. After sorting, the bucket indices in single_bucket_indices are sorted from largest to smallest by bucket capacity. This sorting operation is a crucial foundation for subsequent optimizations. Regarding memory layout, the design principle of contiguous storage of buckets within a window is adopted. The bucket index for window w can be represented as w*2c + bucket_id, ensuring that all buckets within the same window are tightly packed in memory. This layout, combined with the characteristics of sorted bucket indices, significantly improves spatial locality. Since large buckets have contiguous indices after sorting, and adjacent buckets within the same window may be frequently accessed during scalar operations, the cache prefetching mechanism can more efficiently load relevant data into the cache, reducing global memory access latency. To optimize atomic operations and memory allocation, a prefix sum method is used to quickly calculate bucket offsets. By precalculating the starting address of each bucket, this method avoids contention issues during dynamic memory allocation and enables intensive access to large buckets to be completed through continuous memory operations. To further improve parallel efficiency, an asynchronous computation mechanism is introduced: by creating independent CUDA streams (such as stream_large_buckets), computation tasks for large buckets and non-large buckets are separated. Specifically, large bucket computations are synchronized with non-large bucket computations using cudaStreamWaitEvent, preventing small buckets from idling while waiting for large buckets to complete. A dynamic resource allocation strategy is employed at the thread scheduling level. NUM_BLOCKS (number of thread blocks) and NUM_THREADS (number of threads per block) are dynamically calculated based on the number of buckets currently being processed. For example, a configuration of 256 threads / block is used for processing smaller buckets. This dynamic adjustment mechanism maximizes streaming multiprocessor (SM) utilization, ensuring a precise match between computing resources and task load, thereby improving overall throughput. Through these multi-level optimizations, systemic performance improvements are achieved, from memory access patterns to computing task scheduling.

[0070] In addition, the reduction phase is optimized, with iterative reduction (Iterative Reduction): the data size is gradually compressed through multiple rounds of reduction, with the number of source bits halved and the number of windows doubled in each round.

[0071] Dual temporary bucket strategy: temp_buckets1 is used for interleaved reduction (bottom window, even phases). temp_buckets2 is used for serial reduction (top window, odd phases). Adaptive parity: Dynamically adjusts the write stride based on the is_odd_c flag to handle boundary conditions.

[0072] Systematic acceleration is achieved through multi-level memory planning, dynamic computation adaptation, and pipeline parallel architecture. The framework includes the following core strategies:

[0073] Hierarchical memory management: A three-level dynamic memory pool is constructed, with differentiated allocation strategies based on object size. Small metadata (<4KB) uses fixed block pre-allocation to reduce fragmentation; medium-sized temporary buffers (4KB-1MB) utilize a buddy algorithm to optimize memory continuity; and large data blocks (>1MB) utilize graphics memory access policies to optimize physical page mapping. The radix sort bit width is dynamically calculated based on the GPU L2 cache capacity, and the data block structure is divided into cache-compatible blocks to ensure data locality during computation.

[0074] Dynamic Scheduling of Computing Resources: Establish a dual-dimensional adaptation mechanism for thread configuration and algorithm selection. Dynamically plan thread block size based on the number of streaming multiprocessors, and maximize SM occupancy through a load balancing model. Design a hybrid sorting strategy selector that dynamically switches between bitonic sorting, radix sorting, and library function sorting based on real-time data characteristics. Set decision thresholds through offline modeling to ensure optimal theoretical time complexity.

[0075] Pipeline parallel architecture: The computation process is broken down into three stages: cardinality partitioning, sub-block sorting, and global merging. Multi-stream parallelism and event synchronization enable computational communication overlap. For ultra-large-scale bucket data, dedicated processing channels are designed to isolate large and small bucket computation tasks, utilizing asynchronous stream mechanisms to eliminate resource contention. The innovative introduction of space-filling curve (Z-Order) rearrangement technology converts linear memory accesses into a two-dimensional spatial locality pattern, significantly improving cache hit efficiency.

[0076] (3) Pre-calculation scheme

[0077] In the Pippenger algorithm, the introduction of the precompute factor (precompute_factor) aims to optimize the computational efficiency of multi-scalar multiplication (MSM) by trading space for time. This solution uses a hierarchical shift architecture to construct the precompute table, and its mathematical essence can be expressed as:

[0078]

[0079] Where k = precomputer factor, Where, P jrepresents the original base point, w is the window width, and c is the dynamically adjusted blocking coefficient. This formula constructs a multi-resolution pre-calculation table through exponential shift operations, allowing the dot product results corresponding to the scalar blocks to be quickly obtained by table lookup during actual calculations.

[0080] Implementation of layered displacement algorithm

[0081] As shown in Table 1, the pre-calculation factor symbol table is given. Algorithm 2 gives the CUDA implementation core logic for generating the pre-calculation table. The key steps are as follows:

[0082] We propose an adaptive parallel optimization methodology based on resource estimation. This methodology, centered on the resource characteristics of current hardware, uses multi-level dynamic task scheduling and parameter tuning to design the appropriate ratio of CUDA core threads to hardware resources, aiming to achieve efficient allocation and utilization of computing resources. This section details the theoretical foundation and technical implementation of this methodology, providing theoretical support and practical reference for optimization design in heterogeneous computing environments. We primarily use the Volta (V100) architecture as an example to maximize task execution efficiency.

[0083] (1) Resource feature modeling

[0084] This methodology first constructs a multi-dimensional feature model to describe key hardware attributes in a parameterized way. This section mainly takes V100 as an example, using the four-dimensional feature vector H arch = <T core ,B mem ,S sm ,C cache > as a characterization function, specifically including the following dimensions:

[0085] Calculate the core type ratio T core The type and ratio of computing cores directly influence task allocation strategies. The Volta architecture relies more on CUDA cores for general-purpose computing. The V100 typically has 5120 CUDA cores.

[0086] Memory bandwidth and hierarchical characteristics B mem Memory bandwidth and layer design determine the efficiency of data transfer. The V100 graphics memory capacity is typically 32GB.

[0087] Streaming multiprocessor execution resource capacity S sm :The number of streaming multiprocessors (SMs) and the warp capacity supported by each SM are key indicators of parallelism. The SM resources of the Volta architecture are generally

[0088] Cache level capacity and replacement strategy C cacheThe size of the L1 and L2 caches and their replacement algorithms have a significant impact on data access efficiency. The Volta architecture's cache design requires additional software intervention to improve data locality.

[0089] Based on the collected data, the system constructs a runtime decision tree. Based on the load formula modeled using the four-dimensional vector, it adjusts the data partitioning strategy to reduce memory pressure. This dynamic decision-making mechanism ensures that the optimization strategy can adapt to runtime hardware changes, thereby improving the robustness and adaptability of the system.

[0090] (2) Building an architectural feature adaptation mechanism

[0091] Based on the modeling of multi-architecture features, this methodology designs a hierarchical parallel control framework to achieve efficient scheduling of computing resources through dynamic adjustment of three levels of granularity (thread block, grid and warp).

[0092] Dynamic Optimization of Thread Block Topology: Thread blocks are the fundamental unit of parallel execution in CUDA, and their topology directly impacts SM resource utilization. This method proposes a dynamic block dimension adjustment strategy to optimize thread block configuration through the following methods: Width Adjustment: The width (x dimension) of a thread block is dynamically adjusted based on the number of resident warps supported by the SM. The Volta architecture's SMs are relatively small, requiring a narrower width to avoid resource contention.

[0093] By dynamically adjusting the thread block topology, the system can optimize resource allocation on different architectures and avoid performance degradation caused by improper thread block configuration. Global optimization of mesh partitioning: To further improve global parallelism, this method designs an adaptive mesh partitioning strategy and calculates the optimal mesh dimension using the following formula:

[0094]

[0095] SM load balancing coefficient C(S sm ) This factor ensures that tasks are evenly distributed across multiple SMs by analyzing the load distribution of the SMs. With a small number of SMs, load balancing is particularly important and requires careful adjustment to avoid local overloads.

[0096] This formula combines architectural characteristics and task scale to automatically generate the optimal grid partitioning scheme, ensuring efficient use of computing resources.

[0097] (3) Specific implementation of the adaptation mechanism

[0098] For example, hardware status perception and real-time monitoring in the algorithm, multi-dimensional monitoring, real-time capture of key indicators, as shown in Table 3, are executed periodically in independent streams (100-500μs intervals) to monitor performance indicators.

[0099] 2. Dynamic parameter space modeling

[0100] A quantitative parameter-performance relationship model is established, and a function is constructed through coefficients to control the thread block size as shown in the formula:

[0101]

[0102] Q(B, T) = α·U sm (B)+β·H cache (T)-γ·C overhead (B, T) (1.4)

[0103] Table 1 Variable table of bucket sorting algorithm for resource estimation

[0104]

[0105] Algorithm 1 Bucket sorting optimization algorithm for resource estimation

[0106]

[0107]

[0108]

[0109] Algorithm 2: Generation of layered displacement pre-calculation table

[0110]

[0111]

[0112] The optimization strategy for the elliptic curve multi-scalar multiplication computing architecture designed in this chapter mainly includes two aspects: one is the bucket sorting optimization method for scale prediction, and the other is the optimization of the task parallel method for resource estimation.

[0113] (1) Comparison of bucket sorting optimization methods

[0114] Since the bucket sort optimization of resource estimation is targeted at data sets with different loads, especially for data sets of different sizes, specific bucket sort optimization is performed. As can be seen in Table 2, there is a more obvious speedup ratio for small-scale data. 5 At a small scale, the speedup can reach 1.2x. For medium-sized data, the speedup is relatively low, around 1.1x. For large-scale radix sorting, the speedup is relatively small compared to the original algorithm. Therefore, the resource-estimated bucket sorting method described in this chapter is primarily optimized for small-scale data.

[0115] Table 2 Comparison of bucket sorting optimization methods

[0116]

[0117] (2) Implementation mechanism of pre-calculation factors for layered displacement

[0118] This chapter also designs an implementation mechanism experiment based on pre-calculation factors, and observes the experiment from the experimental statistics of different pre-calculation factors. Figure 1 A clear observation is that a larger pre-calculation factor is not necessarily better. As the pre-calculation factor changes, the computational latency shows a trend of first decreasing and then increasing, forming a "smiling curve" pattern. This is consistent with the conclusion that an excessively high pre-calculation factor increases overhead, while an excessively low pre-calculation factor fails to meet Pippenger's computational timeframe. Therefore, a suitable pre-calculation factor is crucial. Next, we will conduct an experiment to optimize the layered displacement pre-calculation factor.

[0119] In this study, a hierarchical pre-calculation factor implementation mechanism was designed to address the problems of static pre-calculation strategies and low resource matching in the traditional Pippenger algorithm. By dynamically dividing and adjusting the pre-calculation factors, the flexibility and performance adaptability of the pre-calculation stage were effectively improved. Figure 2 The delay in this study is significantly reduced compared to the one without implementing this algorithm, and the acceleration ratio is about 1.10 times that before the algorithm optimization.

[0120] Task Parallel Method for Resource Estimation

[0121] The parallel resource estimation method is a method designed by combining the two algorithms mentioned above and then combining it with dynamic adjustment of pre-computation. This experiment mainly verifies the parallel resource estimation method.

[0122] As shown in Table 3, the medium-scale scenario (2 16 ~2 20 ): The speedup ratio (1.075-1.020) is limited, indicating that the benefits of parallel optimization are partially offset by the resource allocation and synchronization overhead in terms of data volume. 18 In the large-scale scenario, the time after optimization is only reduced by 0.55ms compared with that before optimization, and the speedup ratio is only 1.005, which shows that the task granularity is relatively fine, and the communication between threads and resource contention become bottlenecks. 22 The speedup ratio significantly increased to 1.130, and the optimization time was reduced by 12.6ms. This shows that as the data scale increases, this method becomes more advantageous: large-scale tasks can more efficiently amortize the fixed overhead, ultimately achieving near-linear speedup.

[0123] The task parallel optimization solutions for resource estimation all achieve certain speedup ratios on different data scales. It is worth noting that relatively excellent results are achieved on large-scale datasets of 22nd power.

[0124] Experimental data show that the task parallel method of resource estimation is effective in large-scale scenarios (2 22 ) shows a more significant optimization effect (speedup ratio 1.13). Its core advantages are: reducing synchronization overhead through dynamic resource allocation and improving parallel efficiency through coarse-grained task division.

[0125] Table 3 Comparison of task parallel methods for resource estimation

[0126]

[0127] In summary, the optimization methods proposed in this paper have achieved certain acceleration effects on data sets of different sizes. The bucket sorting optimization method performs well on small-scale data sets, with an acceleration ratio of up to 1.2 times, but the improvement is limited for large-scale data sets. Dynamic pre-calculation factor adjustment has a significant improvement on the acceleration effect, with an acceleration ratio of about 1.1 times. The task parallel method of resource estimation has limited acceleration effect in medium-scale scenarios (acceleration ratio of about 1.075 to 1.020), but it is effective in large-scale scenarios (2 22 ), achieving a speedup of 1.13, demonstrating the effectiveness of dynamic resource allocation and coarse-grained task partitioning. These optimization strategies collectively provide significant support for improved system performance.

[0128] Elliptic curve multi-scalar multiplication performance comparison experiment

[0129] This experiment compares the typical CPU-based Libsnark task and the GPU-based advanced work GZKP and icicle. The experimental results are shown in Table 4. 16 to 2 24 On a data scale of 1.5, we achieved an average speedup of 1.4 times compared to icicle and 1.7 to 2 times compared to GZKP, improving the computational efficiency of multi-scalar multiplication. This ultimately verified the effectiveness of this method.

[0130] Table 4 MSM comparative experiment table (BN254)

[0131]

Claims

1. A GPU-accelerated multi-scalar multiplication (MSM) optimization system, characterized by: include: A resource estimation module, configured to dynamically analyze GPU memory capacity, number of streaming multiprocessors, and cache characteristics to generate a resource allocation strategy for multi-scalar multiplication tasks; Dynamic task allocation module, configured to dynamically adjust thread block size and grid division based on window width, number of buckets, and hardware resource characteristics to achieve computational load balancing; a layered displacement pre-calculation module configured to construct a multi-resolution pre-calculation table, dynamically adjust the scalar block coefficients by the layered displacement factors, and generate pre-calculated dot product results for continuous storage; The bucket sort optimization module is configured to adopt a memory layout with continuous storage of buckets within a window, combining prefix and offset calculation with an asynchronous stream scheduling mechanism to eliminate atomic operation contention; The iterative reduction module is configured to compress the data size through multiple rounds of reduction and use a dual temporary bucket strategy and parity bit adaptive writing technology to handle boundary conditions; The hardware adaptation module is configured to dynamically select the sorting algorithm and thread block topology based on the GPU architecture characteristics to optimize cache hit rate and computing resource utilization.

2. The system according to claim 1, wherein: The mathematical expression of the layered displacement pre-calculation module is: Among them, P k,i is the hierarchical displacement, k and i are preset values, j is the sequence number, c is the window width, δ is the dynamically adjusted block displacement, P is the elliptic curve base point, and a hierarchical pre-calculation table is generated to support fast point multiplication table lookup operations.

3. The system according to claim 1 or 2, characterized in that The bucket sorting optimization module is further configured as follows: Sort the bucket indexes in descending order of capacity using `thrust::sort_by_key` to generate an ordered sequence of `single_bucket_indices`; Adopting the physical memory continuous layout strategy, the bucket index of window w is mapped to w·2 c + bucket_id, improves spatial locality; Use independent CUDA streams to separate large and small bucket computing tasks, and use `cudaStreamWaitEvent` to achieve asynchronous stream synchronization.

4. The system according to claim 3, characterized in that The iterative reduction module is further configured to: Define temporary buckets `temp_buckets1` and `temp_buckets2` to handle the even and odd bit reduction phases respectively; Dynamically switch the write target bucket according to the reduction round, and control the parity data distribution through the `is_odd_c` flag; A multi-stage pipeline architecture is adopted to decompose the reduction operation into three stages: cardinality partitioning, sub-block sorting and global merging.

5. The system according to claim 1, wherein: The hardware adaptation module is further configured as: Build a four-dimensional feature vector model to dynamically collect CUDA core ratio, memory bandwidth, SM capacity and cache strategy; Dynamically switch between Bitonic sorting, radix sorting, and library function sorting through a runtime decision tree; The Z-Order space filling curve is used to optimize the two-dimensional memory access pattern and improve the L2 cache hit rate.

6. The system according to claim 1, wherein: It also includes a multi-level memory management module, configured as: For data sizes smaller than 4KB, a fixed block pre-allocation strategy is used; Applying the buddy algorithm to optimize memory continuity for 4KB-1MB buffers; Physical page mapping optimization is implemented for data blocks larger than 1MB, and data layout is guided by combining the memory access pattern.

7. A GPU-accelerated multi-scalar multiplication MSM optimization method, characterized in that: include: Dynamically model GPU hardware characteristics to generate a four-dimensional parameter vector including memory bandwidth, number of SMs, and cache levels; Generate a multi-resolution pre-calculation table based on the window width \(c\) and the layer displacement factor \(\delta\) to achieve fast table lookup in scalar blocks; Column-major memory layout is used to store elliptic curve point coordinates, allowing threads in the same thread block to access continuous memory areas; The asynchronous stream mechanism is used to separate large and small bucket computing tasks, and prefix and pre-computation are used to eliminate dynamic memory allocation competition. Perform multiple rounds of iterative reduction, halving the number of data bits in each round and switching the temporary bucket write path based on the parity bit flag; Dynamically adjust thread block dimensions and grid partitioning strategies to optimize SM occupancy and computing throughput based on real-time load.

8. The method according to claim 7, characterized in that The pre-calculation table generation step includes: The elliptic curve base point P is 2 k·c Perform exponential displacement to generate a hierarchical displacement point set; Dynamically divide the displacement point set of each layer into blocks, and the block coefficient \(\delta\) is adaptively adjusted according to the current GPU memory capacity; The pre-computed results are stored continuously according to the window index to form a linear memory structure that supports merged access.

9. The method according to claim 7 or 8, characterized in that The dynamically adjusting thread block dimension includes: According to the number of resident thread warps supported by SM, according to the formula blockDim x =min(1024, SM_capacity·32) adjust thread block width; Calculate the grid dimensions using the load balancing model: Where N is the total number of tasks, and load_factor is the dynamic adjustment coefficient based on the number of SMs.

10. The method according to claim 7, characterized in that Also includes cross-architecture optimization steps: Enables a warp-level double-buffered pipeline on the Volta architecture to hide register access latency; The radix sort bit width is dynamically divided according to the L2 cache capacity to generate a data block structure that adapts to the cache line.

Images