Method and system for optimizing GPU and accelerating SPHINCS + signature parallel implementation

Abstract

The invention discloses a method and a system for optimizing a GPU (Graphics Processing Unit) to accelerate SPHINCS + signature parallel implementation, and belongs to the technical field of cryptography and GPU parallel computing. In order to solve the problems of computing resource scheduling stiffness, low memory access efficiency and difficulty in optimization caused by an existing single large-kernel structure, a signature process is decomposed into a multi-kernel assembly line, hybrid parallel computing is adopted in an FORS signature stage, and host-device data interaction and GPU occupancy rate optimization are introduced, so that the computing resource scheduling efficiency is improved. And efficient parallel execution of each signature stage is realized. According to the method, the SPHINCS + signature calculation efficiency can be remarkably improved, and complete signature data meeting the standard can be generated.

Description

Technical Field

[0001] This invention belongs to the field of cryptography and GPU parallel computing technology, specifically relating to a method and system for optimizing GPU-accelerated parallel implementation of SPHINCS+ signatures. Background Technology

[0002] SPHINCS+ is a stateless digital signature scheme based on hash functions. Because it effectively resists attacks from both classical and quantum computers, it was selected as one of the algorithms for post-quantum cryptography standardization in 2022. Unlike stateful signature schemes such as XMSS, the stateless nature of SPHINCS+ makes it more advantageous in complex application scenarios involving multiple devices and distributed systems, as it eliminates the need for state synchronization.

[0003] The stateless nature of the signature also leads to higher computational overhead. The SPHINCS+ signature process involves a large number of hash calculations, including building multi-level Merkle trees, such as the multiple trees in FORS (Forest of Random Subsets) signatures and the multi-level trees in XMSS (eXtended Merkle Signature Scheme). To improve its usability, accelerating it using the massively parallel computing capabilities of GPUs has become a key technological approach.

[0004] In the prior art, a parallel implementation scheme for GPU-accelerated SPHINCS+ signature has been disclosed (Chinese patent application CN117040792A). Its core idea is to propose a four-level parallel strategy encompassing "inter-tree parallelism," "inter-node parallelism," "inter-node parallelism," and "hash function parallelism." In its specific code implementation, this scheme tends to adopt a "monopoly large kernel" architecture. That is, the entire SPHINCS+ signature process, from message digest generation to FORS signing, to the construction of multi-level XMSS trees and WOTS+ signing, is implemented in a large and complex device function, and started through a global kernel function. Summary of the Invention

[0005] The purpose of this invention is to propose a method and system for optimizing GPU-accelerated parallel implementation of SPHINCS+ signatures. By decomposing the monolithic large kernel structure into a multi-kernel processing pipeline and employing hybrid parallel computing in the FORS signature stage, combined with GPU resource utilization optimization, this invention addresses the problems of rigid computing resource scheduling, low memory access efficiency, and difficulty in optimizing specific computing stages in existing technologies. The goal is to improve the computational efficiency of SPHINCS+ signatures, reduce GPU execution latency, and achieve efficient parallel processing of each signature stage.

[0006] To achieve the above objectives, the present invention adopts the following technical solution.

[0007] A method for optimizing GPU-accelerated parallel implementation of SPHINCS+ signatures includes the following steps: 1) The original SPHINCS+ signature process, implemented in a single computing core, is decomposed into a processing pipeline consisting of multiple functionally independent computing cores executed on the GPU. The computing cores include a signature preparation core, a FORS underlying construction core for the FORS signature stage, a FORS upper-level construction core and a FORS root hash core, and a signature assembly core. The execution order of each computing core in the processing pipeline is determined based on the data dependencies between each computing stage. 2) During the signature preparation phase, a signature preparation kernel is deployed to generate parameter information including message digest, tree index, and leaf index based on the original message and key, and the generated parameter information is passed to the FORS signature phase; 3) During the FORS signature phase, the FORS underlying construction kernel, the FORS upper-level construction kernel, and the FORS tree root hash kernel are deployed sequentially to execute the node construction and hash calculation of each subtree in parallel on the GPU, generating signature data for the FORS phase. 4) During the XMSS / WOTS+ signing stage, a computation kernel is invoked to perform the XMSS / WOTS+ signing stage, in order to complete the construction of the multi-layer hash tree and the signature calculation, and generate the signature data for the XMSS / WOTS+ stage. 5) During the signature assembly phase, a signature assembly kernel is deployed to combine the original message with the signature data generated in the FORS phase and the signature data generated in the XMSS / WOTS+ phase to generate the final SPHINCS+ signature.

[0008] Further, after step 2) is completed, host-device data interaction is performed, including: Copy the parameter information generated by the kernel for signature preparation from GPU memory to CPU memory; On the CPU side, a FORS address information array is constructed based on the received parameter information; The constructed FORS address information array is copied back to GPU memory to provide input data for the FORS signing stage.

[0009] Furthermore, the parallel execution of the FORS signature phase in step 3) includes: The underlying core of FORS is built using a massively parallel model, where all threads in the GPU computing grid participate in the computation to generate the underlying intermediate nodes of all FORS trees. The FORS upper-layer construction kernel adopts a fine-grained parallel mode of one tree and one thread block. Based on the generated lower-level intermediate nodes, it constructs the remaining upper-layer parts of all FORS trees and generates the root of each FORS tree. The FORS root hash kernel performs concatenation and hash calculation of each FORS root to generate the public key for the FORS stage.

[0010] Furthermore, step 3) optimizes GPU utilization before the underlying kernel of FORS is built and executed, including the following steps: Before starting the kernel, call the CUDA utilization API, and input the kernel function pointer and thread block size; Get the number of thread blocks that can be active simultaneously for each stream processor; The maximum number of thread blocks that can be executed concurrently on the entire GPU is calculated based on the total number of GPU multiprocessors, and this number is used as the upper limit of the kernel grid size.

[0011] Furthermore, the specific processing of the FORS underlying kernel construction in step 3) includes: Each thread uses the Grid-Stride loop to compute the leaf node value corresponding to one or more global leaf indices, based on its globally unique identifier in the GPU computing grid. Perform global synchronization after all leaf nodes have been calculated; Continue using the Grid-Stride loop to perform the first-stage tree reduction calculation on the leaf nodes until the preset switching height H_SWITCH is reached; Global synchronization is performed after each level of reduction calculation is completed to ensure data consistency. Intermediate nodes that reach the height of H_SWITCH are stored in the GPU's global memory to provide input for building the kernel at the upper layer of FORS.

[0012] Furthermore, the switching height H_SWITCH is determined through the following steps: The total height H of the FORS tree and the number of data bytes N per node are determined based on the parameter set of the SPHINCS+ algorithm. The number of bytes of shared memory available for each thread block is determined based on the GPU hardware specifications. A minimum integer Z is determined through iterative calculation, such that... If the available shared memory does not exceed a certain limit, then the smallest integer Z is the switching height H_SWITCH.

[0013] Furthermore, the specific processing of building the kernel at the FORS upper layer in step 3) includes: Set the number of thread blocks in the computation grid to be equal to the number of FORS trees to be constructed; Each thread block determines its corresponding FORS tree index based on its unique identifier. All threads within the thread block work together to build the upper-level nodes of the FORS tree until the root is generated.

[0014] Furthermore, all threads within the thread block collaboratively complete the construction of the upper-level node using a shared memory reduction computation method, including the following steps: Read the data of the intermediate nodes belonging to the tree from global memory and load it into shared memory; The thread performs multiple rounds of reduction calculations in shared memory until the final tree root is generated; After each round of reduction calculation is completed, the synchronization operation within the thread block is performed.

[0015] A system for optimizing GPU-accelerated parallel implementation of SPHINCS+ signatures, used to execute the above method, includes: The computational pipeline scheduling module is used to decompose the SPHINCS+ signature process, which was originally implemented in a single computational core, into a processing pipeline consisting of multiple functionally independent computational cores and executed on the GPU. The computational cores include a signature preparation core, a FORS underlying construction core for the FORS signature stage, a FORS upper-level construction core and a FORS tree root hash core, and a signature assembly core; and to deploy and call each computational core sequentially. The signature preparation module is used to deploy the signature preparation kernel during the signature preparation phase. It generates parameter information, including message digest, tree index and leaf index, based on the original message and key, and passes the generated parameter information to the FORS signature phase. The FORS signature module is used to sequentially deploy the FORS underlying construction kernel, the FORS upper-level construction kernel, and the FORS tree root hash kernel during the FORS signature stage. It is used to execute the node construction and hash calculation of each subtree in parallel on the GPU to generate signature data for the FORS stage. The external signature module is used to construct multi-level hash trees and perform signature calculations during the XMSS / WOTS+ signature stage, generating signature data for the XMSS / WOTS+ stage. The signature assembly module is used to deploy a signature assembly kernel during the signature assembly stage, which combines the original message with the generated signature data from the FORS stage and the signature data from the XMSS / WOTS+ stage to generate the final SPHINCS+ signature.

[0016] Furthermore, the FORS signature module includes: The underlying building unit is used to execute the FORS underlying building kernel and perform underlying node calculations for each FORS subtree; The upper-level building unit is used to execute the FORS upper-level building kernel, perform upper-level node calculations for each FORS subtree until the root of the tree is generated; The root hash unit is used to execute the FORS root hash kernel, which concatenates and hashes the FORS roots generated by the upper-level building units to obtain the final FORS public key data.

[0017] The present invention has achieved the following beneficial effects.

[0018] 1. This invention uses a multi-core layered pipeline architecture to split the original single-core SPHINCS+ signature process into functionally independent computing cores, thereby achieving sequential parallelism and task separation of each signature stage and improving GPU resource scheduling efficiency.

[0019] 2. This invention employs a hybrid parallel strategy in the FORS signature phase, with the lower-level nodes using large-scale Grid-for-all parallel computation and the upper-level nodes using Block-per-Tree combined with shared memory reduction, thereby improving the throughput and parallel efficiency of tree structure construction.

[0020] 3. This invention uses host-device data interaction to construct a FORS address information array, providing efficient input data for the GPU signature stage and enabling CPU and GPU collaborative computing.

[0021] 4. This invention sets a dynamic switching height H_SWITCH, optimizes the switching of parallel computing from the lower layer to the upper layer based on the GPU shared memory capacity, balances global memory access and shared memory utilization, and improves computing performance.

[0022] 5. This invention introduces a GPU utilization optimization module, which calculates the optimal thread block and computing grid parameters to avoid thread idleness and storage bottlenecks, thereby further improving the overall signature efficiency.

[0023] 6. This invention combines the outputs of each stage into a complete SPHINCS+ signature at the end of the pipeline by using a signature assembly kernel, thereby achieving closed-loop output and standardized data encapsulation.

[0024] 7. The present invention can achieve efficient execution of SPHINCS+ signature on GPU, support multi-device and multi-task parallel computing, and solve the performance bottleneck of traditional monolithic large kernel in large-scale applications. Attached Figure Description

[0025] Figure 1 This is an overall flowchart of the method for optimizing GPU-accelerated parallel implementation of SPHINCS+ signatures according to the present invention; Figure 2 A schematic diagram illustrating the process of implementing a hybrid parallel approach for FORS signatures. Detailed Implementation

[0026] The technical solutions of the embodiments of the present invention will be clearly and completely described below with reference to the accompanying drawings. Obviously, the described embodiments are only some, not all, of the embodiments of the present invention. Based on the embodiments of the present invention, all other embodiments obtained by those skilled in the art without creative effort are within the scope of protection of the present invention.

[0027] In the description of this invention, it should be understood that the terms "comprising" and "including" indicate the presence of the described features, integrals, steps, operations, elements, and / or components, but do not exclude the presence or addition of one or more other features, integrals, steps, operations, elements, components, and / or collections thereof. It should also be understood that the terminology used in this specification is for the purpose of describing particular embodiments only and is not intended to limit the invention. As used in this specification and the appended claims, the singular forms "a," "an," and "the" are intended to include the plural forms unless the context clearly indicates otherwise. It should also be further understood that the term "and / or" as used in this specification and the appended claims refers to any combination and all possible combinations of one or more of the associated listed items, and includes such combinations. It should be understood that although terms such as first, second, third, etc., may be used in embodiments of this invention to describe preset ranges, these preset ranges should not be limited to these terms. These terms are only used to distinguish preset ranges from one another. Depending on the context, the word "if," as used herein, can be interpreted as "when," "when," "in response to determination," or "in response to detection."

[0028] The accompanying drawings illustrate various structural schematic diagrams according to embodiments disclosed in this invention. These drawings are not to scale, and some details have been enlarged for clarity, and some details may have been omitted. The shapes of the various regions and layers shown in the drawings, as well as their relative sizes and positional relationships, are merely exemplary and may deviate from reality due to manufacturing tolerances or technical limitations. Furthermore, those skilled in the art can design regions / layers with different shapes, sizes, and relative positions as needed.

[0029] This invention provides a method and system for optimizing GPU-accelerated parallel implementation of SPHINCS+ signatures. The core of this method is to decompose the original monolithic kernel architecture into a multi-kernel processing pipeline structure and adopt a hybrid parallel implementation for the critical FORS signature stage, thereby effectively solving the performance bottleneck problem existing in the background technology.

[0030] Figure 1The overall process of optimizing the GPU-accelerated parallel implementation of SPHINCS+ signatures is shown below.

[0031] 1. Construct a multi-core computing pipeline.

[0032] The original overall signature process, implemented in a single computing core, is decomposed into a processing pipeline consisting of multiple functionally independent computing cores. These cores include a signature preparation core, a FORS low-level construction core, a FORS high-level construction core, a FORS root hash core, and a signature assembly core. The FORS low-level construction core, the FORS high-level construction core, and the FORS root hash core together constitute the FORS (Forest of Random Subsets) signature module. Based on the data dependencies at each stage, the various computing cores in the processing pipeline are sequentially started and executed on the GPU.

[0033] 2. Phase 1: Signature preparation for the kernel.

[0034] At the very beginning of the processing pipeline, such as Figure 1 Phase one, as shown, involves deploying and executing the signature preparation kernel. This kernel is responsible for generating a message digest (mhash), a tree index, and leaf indices based on the original message (m) and the private key (sk), which serve as input data for subsequent kernel computations. The signature preparation kernel is configured for lightweight execution, such as using single-threaded blocks and single-threaded execution, to reduce resource consumption and improve response speed.

[0035] After the kernel signature preparation is complete, please refer to [link / reference]. Figure 1 The interaction between Phase 1 and Phase 2 executes the following host-device data interaction steps.

[0036] (1) Copy the calculated core parameters (including tree index and leaf index) from the device-side GPU memory to the host-side CPU memory.

[0037] (2) The host CPU constructs the address information array required for the FORS signature stage in the host memory based on the received parameters.

[0038] (3) Copy the completed FORS address information array from the host memory back to the device memory to provide input data support for the subsequent FORS signature stage.

[0039] 3. Phase Two: FORS Signature Module.

[0040] like Figure 1Phase Two is shown. In the processing pipeline, the FORS signature module employs a hybrid parallel implementation for the FORS signature stage. This involves deploying a parallel execution system on the GPU, consisting of multiple functionally independent computational kernels, including the FORS low-level construction kernel, the FORS high-level construction kernel, and the FORS root hash kernel. This achieves coordinated execution of inter-task parallelism and intra-data parallelism, significantly improving computational efficiency while ensuring signature correctness. The specific implementation details are discussed below. Figure 2 Please provide a detailed explanation.

[0041] 4. Phase Three: XMSS / WOTS+ Kernel.

[0042] like Figure 1 Phase three is shown. In the processing pipeline, after the FORS signature phase, a computational kernel from an external computation module is integrated and invoked to execute the XMSS / WOTS+ signature phase. This kernel utilizes the output of the preceding phases to complete the construction of subsequent multi-layered hash trees and signature computation, thereby realizing the complete SPHINCS+ signature structure generation process.

[0043] 5. Phase Four: Signing and Assembling the Kernel.

[0044] like Figure 1 Phase four is shown. After each signature calculation stage in the processing pipeline has been completed, the signature assembly kernel is deployed and executed. This kernel is used to combine and encapsulate the signature data generated in each stage and the original message m according to the SPHINCS+ standard specification, and finally generate a complete signature sm to achieve closed-loop output of the entire SPHINCS+ signature process.

[0045] Figure 2 The process of implementing FORS signatures in a hybrid parallel manner is demonstrated, and the details are as follows.

[0046] (1) Deploy and execute the underlying kernel of FORS.

[0047] In this step, a FORS underlying construction kernel is deployed and executed. This kernel employs a "Grid-for-all" massively parallel mode, where all threads in the GPU computing grid participate in the underlying construction task of the entire FORS tree to achieve high-throughput parallel computation. This mode enables the simultaneous generation of all underlying intermediate nodes of the FORS tree, significantly improving the execution efficiency of the tree structure initialization phase. The specific execution process of the FORS underlying construction kernel includes the following sub-steps.

[0048] (1-1) Each thread in the kernel calculates the leaf node value corresponding to one or more global leaf indices using a Grid-Stride loop based on its globally unique identifier (tid, thread identifier) ​​in the GPU computing grid. The calculation results are stored in the GPU global memory for subsequent tree reduction calculations.

[0049] (1-2) After all leaf nodes have been computed, all threads in the kernel perform a global synchronization operation (grid.sync()) to ensure that the computation results of the leaf nodes are visible to all threads in the global memory, thereby ensuring the consistency of the input data for subsequent tree reduction computation.

[0050] (1-3) Each thread continues to use the Grid-Stride loop to perform the first stage of tree reduction calculation on the aforementioned leaf nodes, calculating the parent nodes layer by layer until the preset switching height H_SWITCH is reached. After each layer of reduction calculation is completed, all threads in the kernel perform a global synchronization operation to ensure that the calculation results of the previous layer can be correctly accessed when calculating the next layer.

[0051] (1-4) After completing the reduction calculation and reaching the switching height H_SWITCH, the kernel stores the intermediate node data of all FORS trees at that height in the GPU's global memory as input data for subsequent kernels (such as the upper-level FORS reduction kernel or signature aggregation kernel).

[0052] (2) Deploy and execute the upper layer of the FORS kernel.

[0053] After deploying and executing the FORS lower-level build kernel, a FORS upper-level build kernel is further deployed and executed. This kernel adopts a fine-grained parallel mode of "Block-per-Tree" (one thread block per tree), and completes the upper-level construction of each FORS tree based on the generated lower-level intermediate node data, until the final root node of each tree is generated. The deployment and execution of the FORS upper-level build kernel specifically includes the following sub-steps.

[0054] (2-1) Set the computation grid dimension when the kernel starts up so that the number of its thread blocks is equal to the number of FORS trees to be constructed. Each thread block is responsible for building an independent FORS tree, implementing a "one tree, one thread block" scheduling strategy, thereby achieving fine-grained parallel control and data locality optimization.

[0055] (2-2) During kernel execution, each thread block determines the index of the specific FORS tree it is constructing based on its unique identifier (blockIdx.x) in the computation grid. This index is used to distinguish the tree construction task boundaries between different thread blocks, ensuring that thread blocks do not interfere with each other and that data access is independent.

[0056] (2-3) After determining the target tree index, all threads within each thread block collaboratively and specifically complete the construction of the upper part of that particular FORS tree. This process employs a reduction computation method based on shared memory, specifically including the following refinement operations.

[0057] (a) The thread within the thread block reads the intermediate node data generated above, which belongs only to the specific tree, from the GPU global memory.

[0058] (b) Load the read intermediate node data into the shared memory of the thread block to reduce global memory access latency and improve data processing efficiency.

[0059] (c) The threads within the thread block perform multiple rounds of tree node reduction calculations based entirely on the data in shared memory, generating parent nodes at each level sequentially from bottom to top until the final root of the FORS tree is calculated.

[0060] (d) After each round of reduction calculation is completed, perform a synchronization operation (__syncthreads()) within the thread block to ensure that all threads within the same thread block complete the calculation of the current level before entering the next round of reduction calculation, thus ensuring the correctness of the tree structure construction and data consistency.

[0061] By employing the aforementioned shared memory-based reduction mechanism and "Block-per-Tree" parallel mode, the number of global memory accesses can be significantly reduced, improving the parallel computing efficiency and synchronization reliability during the upper-layer node construction phase. Ultimately, the output of the FORS upper-layer construction kernel is the root node data of all FORS trees, providing the basic input for hash aggregation and authentication path generation in the subsequent signature phase.

[0062] (3) Deploy and execute the FORS tree root hash kernel.

[0063] After deploying and executing the FORS upper-layer construction kernel, a FORS tree root hash kernel is further deployed and executed. This kernel's function is to concatenate and hash the root node data of all generated FORS trees to generate the final FORS public key. This public key serves as input to the external computing module (i.e., the XMSS / WOTS+ kernel), providing core input data for the subsequent generation of the signature chain.

[0064] In this kernel, all root data is read from the GPU's global memory in a predetermined order, processed by a hash function, and then output as a fixed-length public key value. This hash calculation can be based on hash algorithms specified in the SPHINCS+ standard parameter set (such as SHA256, SHAKE256, etc.), and the computational throughput can be further improved in the GPU through parallelized message compression operations.

[0065] Figure 2 The invention also demonstrates the concept and determination method of the switching height (H_SWITCH). H_SWITCH is the boundary level for the FORS tree to switch from massively parallel (Grid-for-all) mode to fine-grained parallel (Block-per-Tree) mode, and its selection directly affects the balance between global memory access overhead and shared memory utilization. The method for determining the switching height includes the following calculation steps.

[0066] (3-1) Based on the parameter set of the SPHINCS+ algorithm, determine the height of the FORS tree (SPHINCS_FORS_HEIGHT) and the number of bytes of data in a single node (SPX_N). Among them, SPHINCS_FORS_HEIGHT represents the total number of levels in each FORS tree, and SPX_N represents the storage length of the hash value of each tree node in bytes.

[0067] (3-2) Based on the GPU hardware specifications, determine the total number of bytes of shared memory available for each thread block in order to evaluate the upper limit of data that a single thread block can accommodate in parallel reduction computation.

[0068] (3-3) Through iterative calculation, determine a minimum integer H_SWITCH such that the following condition is met: The shared memory capacity is not exceeded, meaning all node data above the switching height can be fully loaded into shared memory. This ensures that the upper-layer kernel of FORS completes all reduction calculations in shared memory, maximizing parallel efficiency. This dynamic switching mechanism can automatically adjust according to the shared memory capacity of different GPU models to achieve an optimal balance between computational performance and resource consumption.

[0069] In addition, when executing each computing kernel in the pipeline, especially before starting the FORS underlying build kernel, a GPU utilization optimization step is also included. This step is used to improve GPU resource utilization by pre-calculating kernel startup parameters, avoiding thread idleness and storage bottlenecks. The specific process is as follows.

[0070] (1) Before starting the kernel, the function pointer and block size of the kernel to be started are input by calling the CUDA utilization analysis application interface (API), and the kernel startup configuration parameters that can maximize the utilization of GPU hardware are pre-calculated, including block size, register utilization, shared memory allocation, etc.

[0071] (2) Obtain the number of thread blocks that can be active simultaneously on each stream processor (Multiprocessor) through the API.

[0072] (3) Based on the total number of multiprocessors of the GPU device and the number of active thread blocks obtained, calculate the maximum number of thread blocks that can be executed concurrently on the entire GPU, and use this number as the upper limit of the computing grid size when the kernel starts, so as to make full use of GPU computing resources.

[0073] Through the above steps, the GPU's computing resources can be optimally allocated, enabling each computing core to run efficiently in the pipeline and improving the overall signature generation performance.

[0074] Taking the SPHINCS+-SHA256-192s-robust parameter using the SHA256 hash function as an example, when using the common monolithic kernel implementation method described in the background technology, the signature generation time of the GPU (model: NVIDIA RTX 5090D) is 2.231 ms. After adopting the implementation method based on multi-core pipeline and hybrid parallel structure described in this invention, under the same hardware and parameter conditions, the GPU signature generation time is shortened to 1.514 ms. Therefore, this invention significantly improves the execution efficiency of SPHINCS+ signatures on the GPU through multi-core layered parallel design and shared memory allocation optimization, solving the performance bottleneck problem under the traditional monolithic kernel structure.

[0075] In another embodiment of the present invention, a system for optimizing GPU-accelerated parallel implementation of SPHINCS+ signatures is provided. This system can be used to implement the method for optimizing GPU-accelerated parallel implementation of SPHINCS+ signatures described in the present invention, and includes the following modules.

[0076] The computational pipeline scheduling module is used to decompose the SPHINCS+ signature process, which was originally implemented in a single computational core, into a processing pipeline consisting of multiple functionally independent computational cores and executed on the GPU. The computational cores include a signature preparation core, a FORS underlying construction core for the FORS signature stage, a FORS upper-level construction core and a FORS tree root hash core, and a signature assembly core; and to deploy and call each computational core sequentially.

[0077] The signature preparation module is used to deploy the signature preparation kernel during the signature preparation phase. It generates parameter information, including message digest, tree index, and leaf index, based on the original message and key, and passes the generated parameter information to the FORS signature phase.

[0078] The FORS signature module is used to sequentially deploy the FORS underlying construction kernel, the FORS upper-level construction kernel, and the FORS tree root hash kernel during the FORS signature stage. It is used to execute the node construction and hash calculation of each subtree in parallel on the GPU to generate the signature data for the FORS stage.

[0079] The external signature module is used to construct multi-layer hash trees and perform signature calculations during the XMSS / WOTS+ signature stage, generating signature data for the XMSS / WOTS+ stage.

[0080] The signature assembly module is used to deploy a signature assembly kernel during the signature assembly stage, which combines the original message with the generated signature data from the FORS stage and the signature data from the XMSS / WOTS+ stage to generate the final SPHINCS+ signature.

[0081] The FORS signature module includes the following units.

[0082] The underlying building unit is used to execute the FORS underlying building kernel and perform underlying node calculations for each FORS subtree.

[0083] The upper-level building unit is used to execute the FORS upper-level building kernel, perform upper-level node calculations for each FORS subtree until the root tree is generated.

[0084] The root hash unit is used to execute the FORS root hash kernel, which concatenates and hashes the FORS roots generated by the upper-level building units to obtain the final FORS public key data.

[0085] In another embodiment of the present invention, a terminal device is provided, comprising a processor and a memory. The memory stores a computer program, which includes program instructions. The processor executes the program instructions stored in the computer storage medium. The processor may be a combination of a central processing unit and a graphics processing unit, or other general-purpose processors, digital signal processors, application-specific integrated circuits, etc. It is the computing and control core of the terminal, suitable for implementing one or more instructions, specifically suitable for loading and executing one or more instructions to implement a corresponding method flow or corresponding function. The processor described in this embodiment of the present invention can be used to execute the operation of the optimized GPU-accelerated SPHINCS+ signature parallel implementation method described in this invention, including: decomposing the overall signature process of a single computing core into a processing pipeline composed of multiple functionally independent computing cores; deploying a hybrid parallel implementation method composed of multiple independent cores for the FORS signature stage in the processing pipeline; integrating and calling a computing core for executing the XMSS / WOTS+ signature stage in the processing pipeline; and deploying and executing a signature assembly core after all signature calculation stages of the pipeline have been completed.

[0086] In another embodiment of the present invention, a storage medium is provided, specifically a computer-readable storage medium, which is a memory device in a terminal device for storing programs and data. It is understood that the computer-readable storage medium here can include both the built-in storage medium in the terminal device and extended storage media supported by the terminal device. The computer-readable storage medium provides storage space that stores the terminal's operating system. Furthermore, the storage space also stores one or more instructions suitable for loading and execution by a processor, which can be one or more computer programs (including program code). It should be noted that the computer-readable storage medium here can be high-speed RAM or non-volatile memory, such as at least one disk storage device.

[0087] One or more instructions stored in a computer-readable storage medium can be loaded and executed by the processor to implement the corresponding steps of the optimized GPU-accelerated SPHINCS+ signature parallel implementation method in the above embodiments; one or more instructions in the computer-readable storage medium are loaded by the processor and executed as follows: executing the pipeline architecture deployment step, decomposing the single computing core process into a multi-core pipeline; executing the FORS signature stage implementation step, deploying a hybrid parallel method; executing the external module integration step, calling the XMSS / WOTS+ computing kernel; executing the signature data assembly step, generating the final signature.

[0088] Those skilled in the art will clearly understand that, for the sake of convenience and brevity, the above-described division of functional units and modules is used as an example. In practical applications, the above functions can be assigned to different functional units and modules as needed, that is, the internal structure of the device can be divided into different functional units or modules to complete all or part of the functions described above.

[0089] In the above embodiments, the descriptions of each embodiment have different focuses. For parts that are not described in detail or recorded in a certain embodiment, please refer to the relevant descriptions of other embodiments.

[0090] Those skilled in the art will recognize that the units and method steps of the various examples described in conjunction with the embodiments disclosed in this invention can be implemented in electronic hardware, or a combination of computer software and electronic hardware. Whether these functions are implemented in hardware or software depends on the specific application and design constraints of the technical solution. Those skilled in the art can use different methods to implement the described functions for each specific application, but such implementations should not be considered beyond the scope of this invention.

[0091] The above content is only for illustrating the technical concept of the present invention and should not be construed as limiting the scope of protection of the present invention. Any modifications made to the technical solution based on the technical concept proposed in this invention shall fall within the scope of protection of the claims of this invention.

Claims

1. A method for optimizing GPU-accelerated parallel implementation of SPHINCS+ signatures, characterized in that, Includes the following steps: 1) The overall SPHINCS+ signature process, which was originally implemented in a single computing kernel, is decomposed into a processing pipeline consisting of multiple functionally independent computing kernels and executed on the GPU. The computing kernels include a signature preparation kernel, a FORS underlying construction kernel for the FORS signature stage, a FORS upper-level construction kernel and a FORS tree root hash kernel, and a signature assembly kernel. And based on the data dependencies between each computation stage, the execution order of each computation kernel in the processing pipeline is determined; 2) During the signature preparation phase, a signature preparation kernel is deployed to generate parameter information including message digest, tree index, and leaf index based on the original message and key, and the generated parameter information is passed to the FORS signature phase; 3) During the FORS signature phase, the FORS underlying construction kernel, the FORS upper-level construction kernel, and the FORS tree root hash kernel are deployed sequentially to execute the node construction and hash calculation of each subtree in parallel on the GPU, generating signature data for the FORS phase. 4) During the XMSS / WOTS+ signing stage, a computation kernel is invoked to perform the XMSS / WOTS+ signing stage, in order to complete the construction of the multi-layer hash tree and the signature calculation, and generate the signature data for the XMSS / WOTS+ stage. 5) During the signature assembly phase, a signature assembly kernel is deployed to combine the original message with the signature data generated in the FORS phase and the signature data generated in the XMSS / WOTS+ phase to generate the final SPHINCS+ signature.

2. The method as described in claim 1, characterized in that, Step 2) After execution, perform host-device data interaction, including: Copy the parameter information generated by the kernel for signature preparation from GPU memory to CPU memory; On the CPU side, a FORS address information array is constructed based on the received parameter information; The constructed FORS address information array is copied back to GPU memory to provide input data for the FORS signing stage.

3. The method as described in claim 1, characterized in that, The parallel execution of the FORS signature phase in step 3) includes: The underlying core of FORS is built using a massively parallel model, where all threads in the GPU computing grid participate in the computation to generate the underlying intermediate nodes of all FORS trees. The FORS upper-layer construction kernel adopts a fine-grained parallel mode of one tree and one thread block. Based on the generated lower-level intermediate nodes, it constructs the remaining upper-layer parts of all FORS trees and generates the root of each FORS tree. The FORS root hash kernel performs concatenation and hash calculation of each FORS root to generate the public key for the FORS stage.

4. The method as described in claim 3, characterized in that, Step 3) optimizes GPU utilization before the underlying kernel of FORS is built and executed, including the following steps: Before starting the kernel, call the CUDA utilization API, and input the kernel function pointer and thread block size; Get the number of thread blocks that can be active simultaneously for each stream processor; The maximum number of thread blocks that can be executed concurrently on the entire GPU is calculated based on the total number of GPU multiprocessors, and this number is used as the upper limit of the kernel grid size.

5. The method as described in claim 3, characterized in that, Step 3) involves the following specific processing steps in building the FORS underlying kernel: Each thread uses the Grid-Stride loop to compute the leaf node value corresponding to one or more global leaf indices, based on its globally unique identifier in the GPU computing grid. Perform global synchronization after all leaf nodes have been calculated; Continue using the Grid-Stride loop to perform the first-stage tree reduction calculation on the leaf nodes until the preset switching height H_SWITCH is reached; Global synchronization is performed after each level of reduction calculation is completed to ensure data consistency. Intermediate nodes that reach the height of H_SWITCH are stored in the GPU's global memory to provide input for building the kernel at the upper layer of FORS.

6. The method as described in claim 5, characterized in that, The switching height H_SWITCH is determined through the following steps: The total height H of the FORS tree and the number of data bytes N per node are determined based on the parameter set of the SPHINCS+ algorithm. The number of bytes of shared memory available for each thread block is determined based on the GPU hardware specifications. A minimum integer Z is determined through iterative calculation, such that... If the available shared memory does not exceed a certain limit, then the smallest integer Z is the switching height H_SWITCH.

7. The method as described in claim 3, characterized in that, Step 3) involves the following specific processes for building the kernel at the FORS upper layer: Set the number of thread blocks in the computation grid to be equal to the number of FORS trees to be constructed; Each thread block determines its corresponding FORS tree index based on its unique identifier. All threads within the thread block work together to build the upper-level nodes of the FORS tree until the root is generated.

8. The method as described in claim 7, characterized in that, All threads within a thread block collaborate to construct the upper-level node using a shared memory reduction computation method, including the following steps: Read the data of the intermediate nodes belonging to the tree from global memory and load it into shared memory; The thread performs multiple rounds of reduction calculations in shared memory until the final tree root is generated; After each round of reduction calculation is completed, the synchronization operation within the thread block is performed.

9. A system for optimizing GPU-accelerated parallel implementation of SPHINCS+ signatures, characterized in that, include: The computational pipeline scheduling module is used to decompose the SPHINCS+ signature process, which was originally implemented in a single computational core, into a processing pipeline consisting of multiple functionally independent computational cores and executed on the GPU. The computational cores include a signature preparation kernel, a FORS underlying construction kernel for the FORS signature stage, a FORS upper-layer construction kernel and a FORS tree root hash kernel, and a signature assembly kernel. And deploy and invoke each computing kernel sequentially; The signature preparation module is used to deploy the signature preparation kernel during the signature preparation phase. It generates parameter information, including message digest, tree index and leaf index, based on the original message and key, and passes the generated parameter information to the FORS signature phase. The FORS signature module is used to sequentially deploy the FORS underlying construction kernel, the FORS upper-level construction kernel, and the FORS tree root hash kernel during the FORS signature stage. It is used to execute the node construction and hash calculation of each subtree in parallel on the GPU to generate signature data for the FORS stage. The external signature module is used to construct multi-level hash trees and perform signature calculations during the XMSS / WOTS+ signature stage, generating signature data for the XMSS / WOTS+ stage. The signature assembly module is used to deploy a signature assembly kernel during the signature assembly stage, which combines the original message with the generated signature data from the FORS stage and the signature data from the XMSS / WOTS+ stage to generate the final SPHINCS+ signature.

10. The system as described in claim 9, characterized in that, The FORS signature module includes: The underlying building unit is used to execute the FORS underlying building kernel and perform underlying node calculations for each FORS subtree; The upper-level building unit is used to execute the FORS upper-level building kernel, perform upper-level node calculations for each FORS subtree until the root of the tree is generated; The root hash unit is used to execute the FORS root hash kernel, which concatenates and hashes the FORS roots generated by the upper-level building units to obtain the final FORS public key data.

Images