A Kernel-Secure Dispatch Method and System for Large Illumination Model Matrix Multiplication

Abstract

This invention belongs to the technical field of intelligent lighting operation and maintenance large model inference service, and relates to a secure kernel dispatch method and system for large-scale lighting model matrix multiplication. It aims to solve the problem that existing technologies cannot detect risks related to tensor layout, parallel offset, and environmental drift, leading to kernel dispatch errors or performance degradation. The method includes: responding to matrix multiplication requests, collecting risk-related features, and generating risk fingerprints; if there are associated and verified candidate computation units, execution proceeds in the fast path; otherwise, execution proceeds in the slow path, detecting candidate risks, distinguishing between repairable and unrepairable risks, generating adapter plans and performing layout conversions only for repairable risks, and after verification and performance comparison, marking those that pass as verified and associated with risk fingerprints before execution; if none pass, a fallback method is invoked; the method fails when the risk fingerprint or environment changes. This invention ensures the security and reliability of kernel dispatch in large-scale lighting model inference.

Description

Technical Field

[0001] This invention belongs to the fields of intelligent lighting operation and maintenance large model inference service and GPU computing technology, and particularly relates to a method and system for secure kernel dispatch of large lighting model matrix multiplication. Background Technology

[0002] In large-scale model inference, feedforward networks (FFNs), attention mechanism projection (QKV projection), and routing and expert networks in hybrid expert models (MoEs) all heavily rely on matrix multiplication (GEMM) operations. To improve inference performance, the industry has developed various high-performance matrix multiplication implementations, such as Triton, CUTLASS, TensorRT-LLM, and NVIDIA cuBLASLt libraries. These solutions generally support various high-precision and low-precision (e.g., FP16, BF16, INT8, FP8) GEMM implementations, tiling strategies, register accumulation, and Tensor Core instruction optimizations. Meanwhile, low-precision matrix multiplication formats, represented by block-scaling, significantly improve computational efficiency while maintaining model accuracy by sharing scaling factors at specific block granularities.

[0003] When matrix multiplication (GEMM) is applied to the intelligent operation and maintenance scenario of smart lighting, due to the specific requirements of the intelligent operation and maintenance scenario of smart lighting, the model needs to continuously process alarm texts, equipment status, operation logs, time-series data, etc., which puts forward higher requirements for inference correctness, tail latency and long-term online maintainability.

[0004] There has been much research and publication on the selection and scheduling of matrix multiplication kernels.

[0005] 1. Chinese invention patent application CN113282880A discloses an "Application Programming Interface for Accelerating Matrix Operations," which determines a suitable matrix multiplication algorithm by analyzing matrix multiplication operations. This approach focuses on the static selection of matrix multiplication algorithms and does not address the detection and repair of runtime tensor layout risks or environmental version changes.

[0006] 2. US Patent Application No. US11625605B2 discloses a method for “Selecting computational kernel variants using neural networks,” which trains and utilizes neural networks to generate a list of kernels for computational tasks, thereby identifying the optimal kernel. This method relies on offline training and cannot detect dynamic risks such as tensor parallel offset and discontinuous scaling factor layout in real time.

[0007] 3. PCT patent application WO2020050886A1 discloses a "COMPILER-LEVELGENERAL MATRIX MULTIPLICATION CONFIGURATION OPTIMIZATION", which relates to a method for controlling the execution of GEMM operations on an accelerator containing multiple computing units. It generates execution instructions by determining partition information and computing unit allocation information to optimize block execution. This scheme focuses on block control at the hardware execution level and does not consider kernel dispatch security at the software level.

[0008] 4. US Patent Application No. US9135741B2 discloses an "Interference-driven resource management for GPU-based heterogeneous clusters" that allocates resources by tracking GPU hardware utilization. Although this scheme involves performance degradation handling, it does not include environment fields such as candidate compute unit version, driver signature, and CUDA version in the binding mechanism, and therefore cannot manage cache invalidation caused by environment drift.

[0009] 5. In academic research, QiMeng-GEMM proposes using large language models to understand the architectural characteristics of different hardware platforms and automatically search for optimal GEMM combinations; auto GEMM studies the performance optimization of irregularly shaped matrix multiplication on the kernel; other research introduces compiler-level tuning algorithms based on deep reinforcement learning for GEMM block optimization. The EfficientMeta-kernels framework introduces an adaptive dummy padding scheme for GEMM, using efficient GEMM primitives in conjunction with heterogeneous block sizes. EPS-MoE proposes an expert pipeline scheduler that dynamically selects the best kernel implementation for Group Gemm and Dense Gemm for FFN module computation in the MoE model to optimize inference throughput.

[0010] However, none of the aforementioned existing technologies have systematically solved the matrix multiplication kernel dispatch security problem caused by runtime risk factors such as tensor layout, slice offset, and version environment changes in the specific technical context of large model inference.

[0011] In real-world large-scale model deployments, model weights and scaling factors may originate from different training frameworks, different quantization derivation toolchains, and may undergo different post-processing operations. More importantly, tensor parallelism is widely used to support larger-scale models, meaning each GPU holds only a slice of the weights and scaling factors, accompanied by a global offset. In these complex scenarios, even if the surface features (e.g., dimensions M, N, K) of a matrix multiplication request satisfy the constraints of a high-performance candidate computation unit, the internal properties of its input tensor may still present the following typical risks:

[0012] 1. Layout Risk: Scaling factor tensors may have non-contiguous memory layouts or be transposed views. Directly passing these tensors to high-performance kernels that require contiguous layouts can lead to calculation errors or trigger implicit, time-consuming memory copies. Current technology cannot automatically detect whether the scaling factor is a transposed view or a non-contiguous layout before dispatching, nor can it automatically correct for recoverable risks.

[0013] 2. Slice Offset Risk: In tensor parallel mode, the scaling factor slice on each GPU does not start from logical index 0, but includes a global block offset. If the candidate computation unit does not support this offset, direct calls will access incorrect data. Existing technologies do not include tensor parallel offsets in the risk field of dispatch decisions, making it impossible to determine whether the offset is aligned, whether the candidate computation unit supports it, or to automatically correct it through the adapter.

[0014] 3. Version and Environment Drift Risk: When the versions of the underlying CUDA toolkit, GPU driver, or even the candidate compute units themselves change, previously validated optimal kernel bindings may become unstable, inefficient, or even produce erroneous results. Existing solutions lack a mechanism to bind performance evidence to environmental fields such as candidate compute unit versions, driver signatures, and GPU architectures, and also have no way of determining whether historical caches have become invalid due to environmental changes.

[0015] In summary, existing technologies either focus on optimizing algorithms or configurations, or rely on static, hardware-bound databases, lacking a complete closed loop for perceiving, repairing, verifying, and binding evidence to dynamic runtime risks. When scaling factors or tensor parallel offsets occur in non-contiguous layouts, existing technologies either directly report errors or degenerate into inefficient general implementations, failing to automatically identify repairable risks, perform corrections, and securely reuse high-performance kernels. Therefore, there is an urgent need for a secure kernel dispatch method for large-scale matrix multiplication of lighting models that can systematically solve the above problems for intelligent lighting operation and maintenance large-scale model inference services. Summary of the Invention

[0016] This invention addresses the technical problems in existing technologies, such as the inability to automatically detect risks related to tensor layout, slice offset, and version environment drift in matrix multiplication requests at runtime; the inability to perform automatic adaptation and correctness verification for repairable risks; and the lack of multi-dimensional binding and lifecycle management of performance evidence, risk fingerprints, candidate computation unit versions, and hardware environments, leading to kernel dispatch errors, performance degradation, or continued use of outdated evidence after cache invalidation. It provides a kernel-secure dispatch method and system for large-scale matrix multiplication.

[0017] Therefore, the first objective of this invention is to provide a kernel-secure distribution method for matrix multiplication of large illumination models, which includes the following steps: In response to a matrix multiplication request, risk-related features of the matrix multiplication request are collected. These risk-related features include at least the step size, continuity marker, and transpose view marker of the scaling factor tensor, as well as the tensor parallel offset, to generate a risk fingerprint. If a candidate computation unit with a unique candidate computation unit identifier exists and has been verified as being associated with the risk fingerprint, the fast path is entered, and matrix multiplication is performed directly using the candidate computation unit; otherwise, the slow path is entered. In the slow path, risk detection is performed on candidate computing units to obtain a risk set, which includes repairable and unrepairable risks. An adapter plan is generated for candidate computing units containing only repairable risks. The scaling factor tensor is converted into a target layout according to the adapter plan. The correctness verification of the candidate computing units is performed using the converted scaling factor tensor. Candidate computing units that pass the correctness verification are selected by performance degradation comparison. The selected candidate computing units are marked as verified and bound to the risk fingerprint to form a valid dispatch binding record. The matrix multiplication is performed using the candidate computing units that pass the performance degradation comparison. If no candidate computation unit passes both the correctness check and the performance degradation comparison, the fallback computation unit is called to perform matrix multiplication. After the matrix multiplication request is completed, if the risk-related characteristics corresponding to the valid distribution binding record change, causing the newly generated risk fingerprint to be inconsistent with the risk fingerprint of the valid distribution binding record, or if the operating environment parameters corresponding to the valid distribution binding record change, the valid distribution binding record will be marked as invalid.

[0018] Preferably, the response matrix multiplication request collects risk-related features of the matrix multiplication request and generates a risk fingerprint, including: The risk-related features are normalized into strings in a fixed order, and the hash value of the strings is calculated to obtain the risk fingerprint; The risk-related features include: metadata of the input tensor, metadata of the weight tensor, metadata of the scaling factor tensor, metadata of the output tensor, and runtime environment signature.

[0019] Preferably, the adapter generation plan includes: For a candidate computing unit that contains only the repairable risks, an adapter plan is generated based on the risk set, and the adapter plan corresponds to a unique adapter identifier; The adapter plan allocates the work area and executes the adapter kernel to convert the scaling factor tensor into the target layout, obtaining the converted scaling factor tensor for the candidate computing unit to perform correctness verification. The adapter plan is used to determine: a scaling factor tensor and the workspace size required to perform the adaptation operation, wherein the scaling factor tensor includes the target layout, target step size, and global block offset.

[0020] Preferably, the correctness check includes: Obtain the candidate outputs obtained after the candidate computation unit performs matrix multiplication, and the reference outputs obtained by CPU computation; By comparing the error between the candidate output and the reference output, the maximum absolute error and the relative error are obtained; If the maximum absolute error is less than or equal to the absolute tolerance threshold, or the relative error is less than or equal to the relative tolerance threshold, then the candidate computing unit is determined to have passed the correctness check.

[0021] Preferably, the performance degradation comparison includes: The performance of the candidate computing unit is compared with at least one performance baseline, which includes: a predefined baseline of the security candidate computing unit, the historical median performance of the same risk fingerprint, and the previous version record of the same candidate computing unit; If the performance degradation exceeds a preset threshold, the test will fail.

[0022] As a preferred option, it also includes: Performance evidence is collected from candidate computing units that pass the correctness check, and the collected performance evidence is parsed into a unified format performance evidence record. The performance evidence includes at least runtime latency, kernel execution summary and hardware counter metrics. Associating the performance evidence record with the risk fingerprint and the current operating environment includes: generating an evidence key based on the risk fingerprint, the candidate computing unit identifier, the adapter identifier, and the operating environment signature; binding the performance evidence record with the evidence key and storing it in the evidence database.

[0023] Preferably, marking the valid distribution binding record as invalid when the risk-related characteristics corresponding to the valid distribution binding record change, causing the newly generated risk fingerprint to be inconsistent with the risk fingerprint of the valid distribution binding record, or when the operating environment parameters corresponding to the valid distribution binding record change, includes: A state machine is used to manage dispatch binding records and performance evidence records. The states include at least valid states and invalid states. When the risk-related characteristics or the operating environment are detected to change, the corresponding dispatch binding record is changed from a valid state to an invalid state. When all candidate computation units fail the correctness check and performance degradation comparison, record the risk fingerprint, the reason for failure, and the rollback time.

[0024] The second objective of this invention is to provide a kernel-secure dispatch system for large-scale matrix multiplication of illumination models, comprising: A risk fingerprint generation unit is used to respond to a matrix multiplication request, collect the risk-related features of the matrix multiplication request, and the risk-related features include at least the step size, continuity marker and transpose view marker of the scaling factor tensor, and the tensor parallel offset, to generate a risk fingerprint; The dispatch control unit is configured to select a fast path and directly use the candidate computation unit to perform matrix multiplication if there is a candidate computation unit associated with and verified by the risk fingerprint; otherwise, it selects a slow path. A slow path processing unit is configured to perform risk detection on candidate computation units in the slow path to obtain a risk set, the risk set including repairable risks and unrepairable risks; generate an adapter plan for candidate computation units containing only repairable risks; convert the scaling factor tensor into a target layout according to the adapter plan; perform correctness verification on the candidate computation units using the converted scaling factor tensor; filter candidate computation units that pass the correctness verification through performance degradation comparison; mark the filtered candidate computation units as verified and bind them with the risk fingerprint to form a valid dispatch binding record; and use the candidate computation units that pass the performance degradation comparison to perform the matrix multiplication. The safety fallback unit is configured to call the fallback unit to perform matrix multiplication if no candidate computation unit passes both the correctness check and the performance degradation comparison. The detection management unit is configured to mark the valid distribution binding record as invalid after the matrix multiplication request processing is completed, when the risk-related characteristics corresponding to the valid distribution binding record change, causing the newly generated risk fingerprint to be inconsistent with the risk fingerprint of the valid distribution binding record, or when the operating environment parameters corresponding to the valid distribution binding record change.

[0025] A third objective of this invention is to provide an electronic device comprising: A memory for storing computer programs, the memory including non-volatile memory and random access memory, the non-volatile memory for storing system programs and the random access memory for caching runtime data; A processor is used to execute a program stored in memory to implement any of the steps of the above-mentioned method for kernel-secure dispatching of a large-scale illumination model matrix multiplication.

[0026] The fourth objective of this invention is to provide a computer-readable storage medium storing a computer program, which, when executed by a processor, implements the steps of any of the above-mentioned illumination large model matrix multiplication kernel secure dispatch methods.

[0027] The beneficial effects of this invention are as follows:

[0028] 1. Runtime Risk Awareness and Fingerprint Management: This invention does not advocate for low-precision matrix multiplication calculation formulas, block strategies, or performance profiling tools themselves. Instead, it provides a systematic engineering method for securely and reliably dispatching these high-performance kernels to the illumination large model inference service. Specifically, this invention incorporates multiple dimensions such as tensor shape, stride, dtype, scale layout (continuity, transposed view), storage_offset, TP global offset, candidate compute unit version, GPU architecture, and CUDA / driver signature into the risk fingerprint. This achieves comprehensive awareness and unique identification of dispatch risks, enabling the system to automatically identify hidden risks in matrix multiplication requests at runtime, rather than relying solely on surface features such as M / N / K for kernel selection.

[0029] 2. Automatic adaptation of recoverable risks: For recoverable risks such as scale discontinuity, transposed view, and TP offset mismatch, this invention automatically converts tensors into the layout required by candidate computation cells by generating and executing a CUDA adapter, without degenerating into an inefficient general implementation, thus avoiding manual troubleshooting and manual repair.

[0030] 3. Correctness verification mechanism: This invention performs correctness verification before candidate computation units enter online binding. By comparing the max_abs_error (maximum absolute error), mean_abs_error (mean absolute error), and relative_error (relative error) of the candidate output with the reference output, candidate computation units that produce erroneous outputs are automatically filtered out, fundamentally preventing erroneous kernels from entering the production environment.

[0031] 4. Automatic collection and binding of performance evidence: This invention supports automatic collection of performance evidence in three modes: runtime timing, Nsight Systems and Nsight Compute, and parses it into a unified Perf Evidence Record. The performance evidence is bound to risk fingerprint, candidate computing unit, adapter, candidate computing unit version, GPU architecture and CUDA / driver signature through Evidence Key to form an auditable and reproducible performance evidence chain.

[0032] 5. Performance degradation comparison: This invention compares the performance of candidate computing units with the baseline of secure candidate computing units, the median of the same fingerprint history, and the previous version record of the same candidate computing unit. Only when the performance has not degraded will it be written into the ACTIVE binding, effectively preventing slow candidate computing units or performance degradation versions from entering the online dispatch path.

[0033] 6. Historical Evidence Lifecycle Management: When stride, block shape, TP offset, candidate compute unit version, driver / CUDA signature, or correctness / performance status changes, this invention automatically converts old evidence to STALE, ARCHIVED, or REJECTED status instead of deleting it directly. This prevents obsolete evidence from continuing to affect online distribution while preserving a complete audit trail of performance degradation and version changes.

[0034] 7. Concurrency safety and online availability: This invention avoids data overwriting caused by multiple concurrent inference requests reusing the same temporary video memory through asynchronous execution based on CUDA streams, request-level workspace isolation, CUDA Event lifecycle management, and fallback computing units. At the same time, it ensures that the inference service can still output correct results when the candidate computing unit is abnormal. Attached Figure Description

[0035] The accompanying drawings, which are incorporated in and form part of this specification, illustrate embodiments consistent with the invention and, together with the description, serve to explain the principles of the invention.

[0036] Figure 1 This is a flowchart illustrating a kernel-secure dispatch method for a large-scale illumination model matrix multiplication according to an embodiment of the present invention.

[0037] The accompanying drawings have illustrated specific embodiments of the invention, which will be described in more detail below. These drawings and descriptions are not intended to limit the scope of the invention in any way, but rather to illustrate the concept of the invention to those skilled in the art through reference to particular embodiments. Detailed Implementation

[0038] The present invention will be further described in detail below with reference to the accompanying drawings and specific embodiments. Obviously, the described embodiments are only some, not all, of the embodiments of the present invention. All other embodiments obtained by those skilled in the art based on the described embodiments of the present invention without inventive effort are within the scope of protection of the present invention. Unless otherwise defined, the technical or scientific terms used herein should have the ordinary meaning understood by those skilled in the art to which this invention pertains.

[0039] Terminology Explanation

[0040] Tensor: A fundamental data structure in deep learning, essentially a multidimensional array. Each dimension is called an axis, and the length of each axis is called its shape. In this invention, the input tensor A, weight tensor B, scaling factor tensor scale, and output tensor output are all two-dimensional tensors (matrices).

[0041] CUDA: NVIDIA's parallel computing platform and programming model, allowing developers to write programs that run on NVIDIA GPUs (Graphics Processing Units) using languages ​​such as C / C++. CUDA provides thread hierarchies (grids, thread blocks, threads), memory hierarchies (global memory, shared memory, registers, etc.), and runtime APIs (such as memory management, kernel startup, event synchronization, etc.), forming the foundation for general-purpose GPU computing. The matrix multiplication kernel, scaling factor adapter kernel, memory allocation (CUDA Malloc), event timing (CUDA Event Record), and stream management (CUDA Stream_t) in this invention are all implemented based on CUDA.

[0042] CUDA Stream: An asynchronous execution queue in CUDA. Operations within the same stream are executed sequentially, while operations from different streams can be executed in parallel. The matrix multiplication request in this invention carries a CUDA Stream handle (CUDA Stream_t) to align kernel startup with the caller's asynchronous execution stream.

[0043] GPU Architecture (e.g., sm89): The microarchitecture codename for NVIDIA GPUs. 'sm' stands for Streaming Multiprocessor, and the numbers (e.g., 89, 90) indicate the specific architecture version (e.g., sm89 corresponds to the Ada Lovelace architecture of the RTX 4090, and sm90 corresponds to the Hopper architecture of the H100). The candidate computing unit registry in this invention needs to declare a list of supported architectures to filter kernels adapted to the current GPU.

[0044] Nsight Systems: NVIDIA's official system-level performance analysis tool, used to collect macro-level performance data such as timelines of various operations on the CPU and GPU, kernel execution time, and memory transfers. The performance evidence collection in this invention supports nsys mode, generating an .nsys-rep report via the nsys profile command, and then extracting the kernel summary via nsys stats.

[0045] Nsight Compute: NVIDIA's official kernel-level performance analysis tool, used to collect micro-performance metrics of a single CUDA core, including SM utilization, DRAM throughput, and utilization rate. The performance evidence collection in this invention supports NCU mode, and outputs a CSV file of specified metrics using the `ncu --csv --metrics` command.

[0046] Workspace: A temporarily allocated GPU memory area in CUDA programming used to store intermediate computation results. The AdapterPlan in this invention includes a workspace_bytes field, which specifies the size of temporary GPU memory required for the conversion scaling factor, and is allocated via cudaMalloc.

[0047] Error thresholds (atol / rtol) for correctness verification: In the correctness verification (Guard Profile), two thresholds are used to determine whether the output of a candidate computational unit is correct: atol (absolute tolerance) is the threshold for comparing the maximum absolute error; rtol (relative tolerance) is the threshold for comparing relative errors. A pass is achieved if either max_abs_error ≤ atol or relative_error ≤ rtol. In this embodiment of the invention, atol = 1e-3 and rtol = 1e-3 are assumed to be correct.

[0048] Matrix multiplication (Matmul): In this invention, matrix multiplication specifically refers to the block scaling matrix multiplication operation in large model inference, i.e., calculating C = A × B × scale, where A is the input tensor (activation value), B is the weight tensor, and scale is the scaling factor tensor. The result C (i.e., the output tensor) of this operation is passed to subsequent layers of large model inference (such as activation functions, the next linear layer, etc.) to complete the calculation of feedforward networks, attention projections, or expert networks.

[0049] The 99th percentile latency (p99 latency) refers to the latency value at the 99th percentile after sorting the latency of multiple executions of the same kernel (or the same computational task) from smallest to largest. In other words, 99% of the executions have a latency less than or equal to this value, and only 1% of the executions exceed this value. This metric is often used to measure the tail latency of a system, reflecting the impact of the slowest requests on the user experience. Compared to average latency, p99 latency is more sensitive to occasional performance fluctuations, resource contention, or cache misses, making it more suitable as one of the criteria for judging performance degradation.

[0050] SLA Threshold: In this invention, it refers to a pre-configured performance upper limit (e.g., maximum allowable average latency or 99th percentile latency) by the system administrator or operations personnel. It serves as a performance baseline when historical performance evidence is lacking. When the evidence base does not contain historical performance evidence records for the same risk fingerprint (i.e., the first request encountering that risk characteristic), the system cannot calculate the "historical median of the same risk fingerprint." In this case, the preset SLA threshold is used to determine whether the candidate computing unit meets the minimum performance requirements. For example, if the administrator configures sla_avg_ms=5.0 (average latency must not exceed 5 milliseconds), and the average latency of the candidate computing unit exceeds this value, performance degradation is determined, and the candidate computing unit is rejected. The SLA threshold is typically more stringent than the baseline for secure candidate computing units (e.g., the measured latency of the CPU reference implementation), representing the upper limit acceptable to the business.

[0051] Probability normalization is a mathematical processing method that transforms a set of arbitrary numerical values ​​(which may be positive or negative, and the range of values ​​is not fixed) into a set of non-negative probability values ​​that sum to 1. In simpler terms, it "compresses" this set of values ​​proportionally to the range of 0 to 1, ensuring that they add up to exactly 1. Each value then represents the probability (confidence level) of a certain outcome. For example, if the model outputs three original scores [2.0, 1.0, 0.5], after probability normalization, they become [0.59, 0.24, 0.17], with a sum of 1, corresponding to the confidence levels of three different failures.

[0052] Example 1

[0053] This embodiment provides a secure kernel dispatch method for matrix multiplication in a large-scale lighting model. The invention takes alarm identification, fault attribution, or work order-assisted dispatch inference services within a large-scale intelligent lighting operation and maintenance model as its application background. It uses the gated projection layer (gate_proj) of a feedforward network employing block scaling quantization in this type of large-scale model as a specific matrix multiplication scenario. It constructs a matrix multiplication request scenario where the scaling factor tensor has a transposed view and non-contiguous storage, and also exhibits parallel global tensor offsets. Specifically, in this scenario: 1. Matrix multiplication dimension M=64 (number of tokens): represents the number of street light devices to be processed in the current batch (or the number of monitoring points within the time window); 2. N_local=256 (current GPU column number): indicates the number of intermediate features for the fault features currently being output by the GPU (e.g., the number of intermediate features for 256 fine-grained fault categories such as "lamp head failure", "driver malfunction", "communication interruption", "power supply failure", and "excessive light decay"). 3. N_global=512 (number of global columns): represents the global fault feature dimension (total dimension of tensor parallelism); 4. K=128 (number of channels); represents the status characteristic dimensions of each street light device (such as voltage, current, power factor, cumulative running time, recent alarm code, etc., 128-dimensional features); 5. The block scaling parameters are block_n=128 and block_k=64, where block_n represents a block of 128 elements in N dimensions (fault feature dimension) and block_k represents a block of 64 elements in K dimensions (state feature dimension). The shape of the weight tensor B is [N,K]=[256,128], and the logical shape of the scaling factor tensor is [N / block_n,K / block_k]=[2,2], meaning that each block shares a scaling factor. In this embodiment, the candidate computation unit demo_block_scaled_backend declares the expected block shape as expected_block_n=128 and expected_block_k=64, which is consistent with the requested block scaling parameters. 6. Tensor parallelism is configured as tp_size=2, tp_rank=1, global_n_offset=256, where tp_size represents the total number of tensor parallel cards, tp_rank represents the current GPU's index, and global_n_offset represents the current GPU's starting offset in the global N-dimensional space. Since the total global dimension N_global=512, and each card is allocated 256 columns, when tp_rank=1, the current GPU is responsible for slicing the fault features in the global dimensions 256 to 511. 7. The scaling factor has a logical shape of [Nblk_local=2, Kblk=2] (where Nblk_local represents the number of local N-dimensional blocks, and Kblk represents the number of K-dimensional blocks), but the physical storage uses a transposed view (index formula index=n_global+k_global×Nblk_global, where Nblk_global=4 represents the number of global N-dimensional blocks, and n_global and k_global represent the global N-block index and the global K-block index, respectively), and the stride is non-contiguous (stride=[1,4]). Candidate computation units are required to have a continuous [Nblk_local,Kblk] scaling factor layout and do not support tensor parallel offsets.

[0054] The aforementioned layout risks and tensor parallel offset risks are precisely the objects that this method needs to automatically detect and repair.

[0055] It should be noted that the scale tensor in this embodiment simultaneously has a logical view pointer and an underlying storage base address. `scale.data_ptr` represents the starting address of the local logical view passed in by the inference framework, `scale.storage_base_ptr` represents the base address of the underlying storage region for the scaling factor, and `scale.storage_offset` represents the element offset of the logical view relative to the underlying storage base address. During adapter execution, `storage_base_ptr`, `storage_offset`, `global_n_block_offset`, and `global_k_block_offset` are used together to calculate the source index, avoiding the direct overlay of global offsets onto the local logical view pointer, which could cause out-of-bounds errors. This embodiment will elaborate on each step of the secure dispatch method based on the above scenario.

[0056] An embodiment of the present invention provides a kernel-secure dispatch method for large-scale illumination model matrix multiplication, such as... Figure 1 As shown, the steps include the following.

[0057] S1, responding to the matrix multiplication request, collect the risk-related features of the matrix multiplication request, the risk-related features include at least the step size of the scaling factor tensor, the continuity marker and the transpose view marker, and the tensor parallel offset, and generate a risk fingerprint.

[0058] First, risk-related features are collected. These features include metadata for the input tensor, weight tensor, scaling factor tensor, and output tensor. Each tensor's metadata includes shape, stride, storage offset, data type, continuity marker, and transpose view marker; as well as a runtime environment signature, which consists of the CUDA runtime version, CUDA driver API version, and NVIDIA driver version, in the format "cuda12.8_driver570". Additionally, matrix dimensions M, N, and K, block shapes block_n and block_k, tensor parallel offsets (global_n_offset and global_k_offset), and GPU architecture are also collected. These fields together constitute the complete risk-related features. Then, the risk-related features are normalized into strings in a fixed order, and the hash value of the strings is calculated to obtain the risk fingerprint.

[0059] Specifically, in this embodiment, a local binding table is created during system initialization. This table stores dispatch binding records; among them, dispatch binding records with an active status are valid records. The local binding table is initially empty; dispatch binding records refer to verified and usable candidate computing units and their associated information.

[0060] This step first converts the raw matrix multiplication parameters submitted by the inference framework into a unified structure, MatmulDispatchRequest, which encapsulates all the input information for a single matrix multiplication request. The following steps will use this structure for risk detection and dispatch. Specifically, the meaning of each data item in the MatmulDispatchRequest structure is as follows: 1. Input tensor A (activation, shape [M,K]): represents the activation value to be calculated, such as the operating status characteristics of street light equipment in the batch (voltage, current, power factor, embedding vector of the most recent alarm code). 2. Weight tensor B (weight, shape [N,K]): represents the weight matrix of the fully connected layer, which has been overquantized and split into tensors in parallel; 3. Scaling factor tensor (scale(block scaling factor, shape [N / block_n, K / block_k], where N represents the output feature dimension, K represents the inner product dimension, and block_n and block_k represent the block scaling size in N and K dimensions, respectively): used to scale the result of each block in low-precision matrix multiplication; 4. Output tensor (shape [M,N]): Used to store the calculation result of matrix multiplication. After subsequent probability normalization processing, the result can be used to obtain the confidence level of various faults (such as "lamp head fault" confidence level, "driver abnormality" confidence level, "communication interruption" confidence level, etc.), which is used to determine whether to dispatch a maintenance work order or trigger an alarm.

[0061] The data type, layout, and storage attributes of the above four tensors together determine whether the candidate computation unit can perform matrix multiplication correctly and efficiently. Therefore, the metadata of the above four tensors is used as the core field of risk fingerprint.

[0062] 5. Matrix dimensions M, N, K: where M is the number of rows in the input tensor A, i.e., the sequence length; N is the number of rows in the weight tensor B, i.e., the output feature dimension; and K is the number of columns in A and B, i.e., the inner product dimension. 6. Block shapes block_n and block_k: represent the block scaling size in N and K dimensions, respectively; 7. Model name (model_name), layer ID (layer_id), and layer type (layer_type): These are used to identify the model and specific layer to which the current matrix multiplication belongs, facilitating customized dispatch strategies by layer; 8. TensorParallelInfo, including: tp_size (total number of tensor parallel GPUs), tp_rank (current GPU rank), global_n_offset (current GPU's starting offset in the global N-dimensional plane), global_k_offset (current GPU's starting offset in the global K-dimensional plane), shard_axis (sharding axis, 0 indicates sharding along the M-axis, 1 indicates sharding along the N-axis, 2 indicates sharding along the K-axis, -1 indicates no sharding; in this embodiment, sharding along the N-axis is used as an example, and other sharding methods can be deduced similarly); and CUDA stream handle (used to align with the caller's asynchronous execution flow, ensuring kernel startup is synchronized with the stream). global_n_offset and global_k_offset are collectively referred to as tensor parallel offsets.

[0063] Collect metadata for each tensor, including: 1. Shape: The length of each dimension; 2. Stride: The interval between adjacent logical elements in the underlying storage, in units of the number of elements, used to determine continuity and layout; if the underlying framework returns a byte stride, the system first divides it by the number of bytes of the corresponding data type to convert it into an element stride before participating in risk fingerprint generation and adapter index calculation. 3. Storage offset: The offset of the tensor data pointer relative to the base pointer, used for handling view tiles; 4. Data type (dtype): For example, FP32, FP16, BF16, INT8; 5. Continuity marker (is_contiguous): Whether memory is stored contiguously, that is, whether adjacent logical elements are physically adjacent; 6. Is_transposed_view flag: Whether the tensor was generated by a transpose operation, resulting in a reversed step order.

[0064] These tensor metadata elements constitute the main body of the risk fields. Among them, the step size, continuity marker, and transpose view marker directly reflect the layout details of the scaling factor and are the core fields that distinguish the risk fingerprint of this invention from the prior art.

[0065] In addition, this step also collects risk fields related to the runtime environment, including the version of the candidate computing unit, GPU architecture, CUDA version, and driver signature. All risk fields (tensor metadata and environment information) are normalized into strings in a fixed order, and then hash values ​​are calculated to obtain the risk fingerprint. The driver signature (driver_cuda_signature) is generated by combining the current system's CUDA runtime version (obtained through cudaRuntimeGetVersion) and GPU driver version (obtained through cudaDriverGetVersion), in the format "cuda<major runtime version>.<minor runtime version>_driver<major driver version>.<minor driver version>", for example, "cuda12.8_driver570". This signature is used to detect cache invalidation caused by environmental changes.

[0066] Unlike existing technologies that only use M / N / K or data type selection kernels, the risk fingerprint of this invention incorporates layout details of scaling factors (step size, continuity marker, transpose view marker) and tensor parallel offset as dynamic risk fields, enabling requests that appear identical but have different internal risks to be distinguished.

[0067] S101, Receive the inference framework request and extract the request information.

[0068] The following information is read from the inference framework: pointers to input tensor A, weight tensor B, scaling factor tensor (scale), and output tensor; values ​​of matrix dimensions M, N, and K; values ​​of block shapes (block_n and block_k); model name (model_name), layer identifier (layer_id), and layer type (layer_type); tensor parallelism information: total number of parallel GPUs (tp_size), current GPU rank (tp_rank), current GPU's initial offset in the global N-dimensional plane (global_n_offset), current GPU's initial offset in the global K-dimensional plane (global_k_offset), and shard axis; and the CUDA stream handle (stream). This information is then populated into the predefined data structure `MatmulDispatchRequest`.

[0069] In this embodiment, the tensor A input to the inference framework has a shape of [64, 128], the tensor B has a shape of [256, 128], the scale logical shape is [2, 2] but the physical storage is a transposed view, and the output shape is [64, 256]. Specific parameter values ​​are: matrix row count M = 64, output column count N = 256, inner product dimension K = 128, N-dimensional block shape block_n = 128, K-dimensional block shape block_k = 64, model name model_name="lighting_fault_model", layer identifier layer_id = 12, layer type layer_type="gate_proj", total number of parallel tensor cards tp_size = 2, current GPU rank tp_rank = 1, current GPU's starting offset in the global N-dimensional dimension global_n_offset = 256, current GPU's starting offset in the global K-dimensional dimension global_k_offset = 0, and shard axis shard_axis = 1 (sharding along the N-axis). The CUDA stream is the default stream (nullptr). This completes the MatmulDispatchRequest structure.

[0070] S102, collect tensor metadata from risk-related features.

[0071] Based on the tensor pointers and shapes obtained in MatmulDispatchRequest in step S101, for each tensor, determine whether it is stored continuously by its shape and element stride: if for all dimensions i, stride[i] == stride[i+1] × shape[i+1], and the stride of the last dimension is 1, then it is stored continuously.

[0072] For the scaling factor tensor, the system does not simply rely on shape to determine whether it is a transposed view. Instead, it considers the tensor view source, stride order, storage_offset, and whether the last dimension stride is 1 to determine whether it belongs to a non-standard layout. If the tensor is generated by view operations such as transpose (exchanging the tensor dimension order), permute (rearranging dimensions in a specified order), or as_strided (modifying the tensor's stride and shape), or if logically adjacent elements are not contiguous in the underlying storage, or if a low-dimensional logical index corresponds to a smaller stride while a high-dimensional logical index corresponds to a cross-block stride, then it is marked as is_non_standard_scale_layout (non-standard layout flag, indicating that the scaling factor tensor layout does not meet the requirements of the candidate computation unit) or is_transposed_or_strided (transposed or strided view flag, indicating that the tensor was generated by transpose, permutation, or strided view operations).

[0073] The storage offset `storage_offset` is obtained by the element offset between the tensor's underlying storage base address `storage_base_ptr` and the logical view start pointer `data_ptr`. This is used to populate TensorMeta structures for tensors A, B, scale, and output, respectively. These structures describe the tensor's memory layout attributes and include the following fields: underlying storage base address `storage_base_ptr`, logical view start pointer `data_ptr`, shape, element stride, storage offset `storage_offset`, data type `dtype`, contiguous flag `is_contiguous`, non-standard layout flag `is_non_standard_scale_layout`, or transposed / stretched view flag `is_transposed_or_strided`. `storage_base_ptr` is used by the adapter for global offset index calculation, and `data_ptr` is used for direct access under no-offset or localized layout conditions. These fields are key inputs for generating risk fingerprints.

[0074] Determining if a view is transposed by step size order: For a 2D tensor, if stride[0] == 1 and stride[1] == shape[0], then it is a transposed view. The storage offset (storage_offset) is obtained by the difference between the tensor's base pointer and data pointer.

[0075] This results in the TensorMeta structure being populated for tensors A, B, scale, and output. This structure describes the memory layout attributes of the tensors and includes the following fields: data pointer (data_ptr), shape, stride, storage offset (storage_offset), data type (dtype), contiguousness flag (is_contiguous), and whether it's a non-standard layout flag (is_non_standard_scale_layout) or a transposed / straddled view flag (is_transposed_or_strided). The data pointer (data_ptr) points to the starting address of the tensor in GPU memory, and the meanings of the remaining fields are the same as defined in step S1. These fields are key inputs for generating the risk fingerprint.

[0076] In this embodiment: Tensors A and B are both stored contiguously: is_contiguous=true, is_transposed_view=false; The scaling factor tensor scale: shape=[2,2], but the physical storage stride=[1,4] (because there are 4 N blocks globally), storage offset=0, is_contiguous=false (not satisfying the continuity condition: the last dimension stride[1]=4≠1), is_transposed_view=true (because stride[0]=1 is less than stride[1]=4); The output tensor is stored contiguously.

[0077] S103, the ability to read candidate computational units.

[0078] The capability descriptions of candidate computing units are read from a candidate computing unit registry (which can be registered via configuration file or code). A candidate computing unit refers to a low-level computing unit or software module capable of performing matrix multiplication operations, such as a matrix multiplication kernel wrapper implemented based on libraries like Triton, CUTLASS, TensorRT-LLM, and cuBLASLt. Each candidate computing unit declares its supported computing capabilities and constraints (such as GPU architecture, data type, block shape, etc.). The security dispatch layer then selects candidate computing units that meet the current request conditions for correctness and performance verification.

[0079] Based on the current GPU architecture (obtained via cudaGetDeviceProperties), the requested data type, and the block shape, the system filters matching candidate compute unit capability descriptions from the registry, resulting in a BackendCapability structure. The BackendCapability structure declares the compute capabilities and constraints supported by a candidate compute unit, containing the following fields: 1. backend_id: A unique identifier for the candidate computation unit; 2. backend_version: The version number of the candidate computation unit; 3. supported_arch: A list of supported GPU architectures (e.g., 89 corresponds to RTX4090, 90 corresponds to H100). 4. supported_dtype: List of supported data types (e.g., FP32, FP16, BF16, INT8); 5. expected_block_n and expected_block_k: The expected N-dimensional and K-dimensional block shapes of the candidate computational units; 6. require_contiguous_scale: Whether to require the scaling factor tensor to be stored contiguously in memory; 7. support_tp_offset: Whether tensor parallel offsets (i.e., non-zero global_n_offset or global_k_offset) are supported.

[0080] Subsequent risk detection will determine whether the request is safe and executable based on the above fields.

[0081] In this embodiment, the capabilities of demo_block_scaled_backend are read: backend_version="demo_backend_v2", supported_arch contains 89 (current GPU is RTX4090), supported_dtype contains FP32, require_contiguous_scale=true, support_tp_offset=false, expected_block_n=128, expected_block_k=64.

[0082] S104, Generate and hash the risk fingerprint field.

[0083] The request risk fields collected in steps S101 to S103 are concatenated into a normalized string in a fixed order, which serves as the original string for the risk fingerprint. These request risk fields include: matrix dimensions M, N, and K; block shapes block_n and block_k; shape, stride, and data type dtype of input tensor A; shape, stride, and data type dtype of weight tensor B; shape, stride, storage offset, contiguousness flag is_contiguous, non-standard layout flag is_non_standard_scale_layout, or transposed / straddled view flag is_transposed_or_strided; shape and stride of output tensor; tensor parallel offsets global_n_offset and global_k_offset; shard axis; GPU architecture gpu_arch; and runtime environment signature consisting of CUDA runtime version, CUDA Driver API version, and NVIDIA driver version.

[0084] The risk fingerprint does not include backend_id and backend_version. backend_id, adapter_id, and backend_version are used to generate the EvidenceKey later. Therefore, the risk fingerprint identifies the risk characteristics of a single matrix multiplication request itself, while the EvidenceKey identifies performance evidence for a candidate computational unit under a specific adapter and environment.

[0085] The following is an example format for concatenating strings: "64|256|128|128|64|fp32|A_shape=[64,128]|A_stride=[128,1]|B_shape=[256,128]|B_stride=[128,1]|scale_s hape=[2,2]|scale_stride=[1,4]|scale_storage_offset=0|scale_is_contiguous=0|scale_is_non_standard_lay out=1|output_shape=[64,256]|output_stride=[256,1]|global_n_offset=256|global_k_offset=0|shard_axis=1 |gpu_arch=89|runtime_environment_signature=cuda_runtime_12.8|cuda_driver_api_12.8|nvidia_driver_570".

[0086] The normalized string is then hashed using SHA-256, BLAKE3, or other deterministic hashing algorithms, and a fixed-length hexadecimal string is taken as the risk fingerprint. This hash result remains stable across different processes, compilers, and machines. The system can prefix the hash value with "fp_" to obtain the final risk fingerprint string, `fingerprint_hash`.

[0087] In this embodiment, the calculated risk fingerprint is the example value "fp_a1b2c3d4e5f6". The risk fingerprint uniquely identifies the risk characteristics of the current request and is passed to step S2 as the final result of step S1.

[0088] S2. If there is a candidate computation unit with a unique candidate computation unit identifier that is associated with the risk fingerprint and has been verified, enter the fast path and directly use the candidate computation unit to perform matrix multiplication; otherwise, enter the slow path.

[0089] Specifically, the local binding table stores verified dispatch binding records, where records with an active status are valid dispatch binding records. The binding table is queried using the risk fingerprint as the key. If the retrieved record is active and the candidate compute unit version, CUDA, and driver environment signature in the record match the current runtime environment signature, then a match is found and the system enters the fast path; otherwise, it enters the slow path.

[0090] This step uses the risk fingerprint (fingerprint_hash) generated in step S1 to query the binding table in the local cache. The binding table is a key-value pair container, where the key is the risk fingerprint and the value is the dispatch binding record (DispatchBinding). The DispatchBinding structure stores the identifiers of verified candidate computation units and their associated information, including the following fields: 1. backend_id: A unique identifier for the candidate computation unit; 2. backend_version: The version number of the candidate computation unit; 3. adapter_plan: The adapter plan describes how to translate the scaling factor into the layout required by the candidate compute units; 4. guard_profile: Correctness verification results, recording error indicators and whether the verification passed; 5. perf_profile: Performance profile, recording performance metrics such as latency and throughput; 6. evidence_key: The evidence key is used to bind performance evidence to risk fingerprints, candidate compute units, adapters, versions, and hardware environments; 7. state: Bound state, such as ACTIVE or STALE.

[0091] If a binding record exists with an ACTIVE status, and the candidate compute unit version, adapter version, GPU architecture, CUDA runtime version, CUDA Driver API version, and NVIDIA driver version in the record are all consistent with the current environment, then the binding is directly reused, skipping subsequent detection, adaptation, and verification processes to achieve rapid deployment. The binding stores not only the candidate compute unit identifier but also the adapter plan, workspace requirements, evidence key, and environment version signature. Version consistency is verified upon a match to prevent incorrect reuse after environment changes.

[0092] In the specific implementation, the fingerprint_hash is used as the key to query the bound table (e.g., std::unordered_map).<std::string,DispatchBinding> If a record is found whose state equals BindingState::ACTIVE, whose backend_version equals the version number of the current candidate compute unit, and whose driver_cuda_signature equals the CUDA / driver environment signature of the current environment, then a hit is determined, and the subsequent slow path is skipped (steps S301 to S103). The system increments the hit count hit_count by 1, updates the last used time last_used_at_ms, and returns the binding. Otherwise, the system proceeds to the slow path in step S301.

[0093] In this embodiment, since it is the first execution, there is no corresponding risk fingerprint in the binding table, so it is determined to be a miss, and the process switches to the slow path (step S3). If a miss is found, matrix multiplication is performed directly using the candidate calculation unit in the returned binding, and the subsequent slow path is skipped (step S3).

[0094] S3, in the slow path, the candidate computation units are checked for correctness, and the candidate computation units that pass the correctness check are selected by performance degradation comparison. The selected candidate computation units are marked as verified and associated with the risk fingerprint, and the matrix multiplication is performed using the candidate computation units that pass the performance degradation comparison.

[0095] When step S2 fails (i.e. there is no valid ACTIVE binding in the binding table), this step executes the complete slow path, screening candidate computing units, detecting risks, planning and executing adapters, verifying correctness, collecting and parsing performance evidence, binding evidence, and comparing degradation based on the risk fingerprint. Finally, the verified candidate computing units are written into the ACTIVE binding for subsequent reuse.

[0096] In the slow path, based on the risk fingerprint, the following steps are performed:

[0097] S301, Filter candidate computation units.

[0098] In this embodiment, multiple candidate computing units are pre-registered into the adaptation framework through a registration interface, forming a list of registered candidate computing units.

[0099] If step S2 fails, a complete detection, adaptation, and verification process needs to be executed. This step, based on the current hardware and the basic constraints of the request, filters out candidate computing units that may support the request from the list of registered candidate computing units. The filtering is based not only on data type and architecture but also on block shape and scaling factor layout requirements, eliminating obviously mismatched candidate computing units in advance.

[0100] During implementation, all registered candidate computation units are traversed. Registered candidate computation units refer to instances added to the security dispatch layer via the `register_backend` interface. `register_backend` is a registration method provided by the security dispatch layer, used to add candidate computation units (objects implementing the `IMatmulBackend` interface) to an internal list. `IMatmulBackend` is an abstract interface class that declares four methods that candidate computation units must implement: `capability()` (returns a description of the candidate computation unit's capabilities), `plan_adapter()` (plans an adapter), `can_direct_run()` (determines whether it can be run directly), and `run()` (performs matrix multiplication).

[0101] Call the `capability()` method of the candidate computation unit to obtain its capability description. This method returns a `BackendCapability` structure (see step S103). If the candidate computation unit's supported architecture list includes the current GPU architecture, and its supported A, B, scale, and output data types respectively include the requested corresponding data types, and the candidate computation unit's expected block shape matches the requested `block_n` and `block_k`, then add the candidate computation unit to the candidate list. This yields a list of candidate computation units of type `std::vector`. <std::shared_ptr <imatmulbackend>>

[0102] In this embodiment, demo_block_scaled_backend (a candidate computational unit that implements a block scaling matrix multiplication demonstration kernel) has been pre-registered to the secure dispatch layer. Only demo_block_scaled_backend in the secure dispatch layer registry simultaneously meets the requirements of the current GPU architecture (sm89 in this embodiment), data type FP32, and desired block shape of 128×64; therefore, the candidate list only includes this candidate computational unit.

[0103] S302, perform risk detection on the candidate computing units to obtain a risk set, which includes repairable risks and unrepairable risks.

[0104] Examine the compatibility of candidate computing units with computing resources and data formats, and distinguish between irreparable and repairable risks. Repairable risks are those that can be compensated for or converted through the adapter plan, while irreparable risks cannot be resolved through the adapter plan.

[0105] The computing resources include GPU architecture.

[0106] The data format includes data type, tensor layout, and tensor parallel offset. Specifically, the tensor layout includes: the continuity of the weight tensor, the continuity of the scaling factor tensor, and the view form of the scaling factor tensor (e.g., whether it is a transposed view or a straddle view).

[0107] During detection, the fields in the risk fingerprint are compared item by item with the capability constraints of the candidate computation unit to identify one or more risk flags (RiskFlags). Risk flags describe multiple risk factors that may coexist in the same matrix multiplication request, such as scale factor transpose / step view, non-continuous scale factor, tensor parallel offset, block shape mismatch, and unsupported data type. This embodiment does not limit the detection result to a single enumeration value, but instead generates a risk set (RiskSet), which can contain zero or one or more risk flags simultaneously.

[0108] RiskFlag is an enumeration type, and its complete values ​​and meanings are as follows: SAFE_DIRECT_RUN: Risk-free, can be run directly, no adapter required; SCALE_TRANSPOSE_STRIDE_MISMATCH: The scaling factor is a transposed view, and the step order does not match the continuous layout required by the candidate computational cells, which is a recoverable risk. NON_CONTIGUOUS_SCALE: The scaling factor is stored non-contiguously in memory, while candidate computation units require contiguous storage, which is a recoverable risk. NON_CONTIGUOUS_WEIGHT: The weight tensor is discontinuous while the candidate computational unit is required to be continuous, which is an irreparable risk. TP_OFFSET_REQUIRED: There is a non-zero tensor parallel offset that the candidate computation unit does not support, which is a recoverable risk; TP_OFFSET_NOT_ALIGNED: Tensor parallel offsets are not aligned to block size, which is an unrecoverable risk; UNSUPPORTED_DTYPE: The requested data type is not in the list of candidate compute units supported; UNSUPPORTED_BLOCK_SHAPE: The requested block shape does not match the expected shape of the candidate computational unit; UNSUPPORTED_ARCH: The current GPU architecture is not in the list of architectures supported by the candidate compute units.

[0109] The specific inspection items and their corresponding risk markers are as follows: If the GPU architecture is not in the list of architectures supported by the candidate computing unit, an unrepairable risk flag "architecture not supported" is generated. That is, if the GPU architecture is not in supported_arch, UNSUPPORTED_ARCH (unrepairable) is added to RiskSet. If any data type is not in the list of data types supported by the candidate computation unit, an unrepairable risk flag "data type not supported" is generated. That is, if any data type is not supported, UNSUPPORTED_DTYPE (unrepairable) is added to RiskSet. If the block shape does not match the expected block shape of the candidate computation unit, an unrepairable risk flag "block shape mismatch" is generated. That is, if the block shape does not match, UNSUPPORTED_BLOCK_SHAPE (unrepairable) is added to the RiskSet. If the weight tensor is discontinuous but the candidate computation unit requires continuity, an unrepairable risk marker "weight discontinuous" is generated. That is, if the weight is discontinuous but the candidate computation unit requires continuity, NON_CONTIGUOUS_WEIGHT (unrepairable) is added to RiskSet. If the tensor parallel offset is not aligned to the block size, an unrepairable risk flag "tensor parallel offset not aligned" is generated. That is, if the TP offset is not aligned to the block size (global_n_offset%block_n!=0 or global_k_offset%block_k!=0), TP_OFFSET_NOT_ALIGNED (unrepairable) is added to the RiskSet. If the tensor parallel offset is non-zero and the candidate computation unit does not support tensor parallel offset, a repairable risk flag "requires tensor parallel offset" is generated. That is, if the TP offset is non-zero and the candidate computation unit does not support TP offset, TP_OFFSET_REQUIRED (repairable) is added to the RiskSet. If the scaling factor tensor is a transposed view, straddle view, or other non-standard layout, a repairable risk flag "scaling factor layout is non-standard" is generated. That is, if the scaling factor is a transposed, straddle, or other non-standard layout, SCALE_TRANSPOSE_STRIDE_MISMATCH (repairable) is added to the RiskSet. If the scaling factor tensor is discontinuous but the candidate computation unit requires continuity, a repairable risk marker "scaling factor discontinuous" is generated. That is, if the scaling factor is discontinuous but the candidate computation unit requires continuity, NON_CONTIGUOUS_SCALE (repairable) is added to the RiskSet.

[0110] After the inspection is complete, a risk set, RiskSet, is obtained. The risk set is in the form: {flags:[RiskFlag1,RiskFlag2,...],has_unrepairable:bool,need_adapter:bool}. It is formed as follows: all detected RiskFlags are sequentially stored in the flags list; if any unrepairable risk flag is added, has_unrepairable is set to true, otherwise false; if flags is not empty and has_unrepairable is false, need_adapter is set to true, otherwise false.

[0111] The risk set contains zero or one or more of the aforementioned risk markers and records whether there are any irreparable risks.

[0112] Decision-making logic: If RiskSet is empty (i.e., flags are empty), it is determined to be SAFE_DIRECT_RUN, and the candidate computation unit is run directly without an adapter.

[0113] If RiskSet is not empty and has_unrepairable is false (i.e., RiskSet only contains repairable risks), then for the candidate computational unit that only contains repairable risks, an adapter plan is planned based on its RiskSet.

[0114] If has_unrepairable is true, then the candidate computation unit is rejected, and other candidate computation units are evaluated.

[0115] In this embodiment, the current candidate computation unit requires a continuous scaling factor and does not support TP offset. The requested scaling factor is a transposed / stepped view and is not continuous, with global_n_offset=256 (non-zero). The detection result RiskSet simultaneously contains TP_OFFSET_REQUIRED, SCALE_TRANSPOSE_STRIDE_MISMATCH, and NON_CONTIGUOUS_SCALE, and does not contain unrepairable risks. Therefore, it is determined to be a set of repairable risks and enters the adapter planning stage.

[0116] S303, for candidate computing units that contain only the repairable risks, generate an adapter plan based on the risk set, wherein the adapter plan corresponds to a unique adapter identifier.

[0117] The adapter plan determines the transformation rules for converting the current input data into the layout required by the candidate computing units, as well as the workspace size required to perform the transformation.

[0118] The adapter plan corresponds to a unique adapter identifier, which is used to determine: the scaling factor tensor and the work area size required to perform the adaptation operation, wherein the scaling factor tensor includes the target layout, the target step size, and the global block offset.

[0119] Specifically, the adapter plan starts with the underlying storage base address of the scaling factor tensor, combines the element offset of the logical view relative to the base address, and the N block offset and K block offset in the global block, multiplies them by the step size of the source data in the N-dimensional and K-dimensional dimensions respectively, and then sums them to calculate the source index. Meanwhile, the target index is calculated by multiplying the local N-block index and the local K-block index by the step sizes of the N-dimensional and K-dimensional dimensions in the target storage, respectively, and then summing the results. Finally, the element at the source index position in the underlying storage base address is assigned to the target index position in the target storage; Wherein, the local N-block index and the local K-block index are the N-dimensional and K-dimensional coordinates of the current processing block within its respective block, respectively.

[0120] During implementation, the `plan_adapter()` method of the candidate computation unit is called. This method calculates the following parameters based on the risk fingerprint: The number of local N blocks is nblocks_local = N / block_n (assuming N is divisible by block_n); The number of local K blocks is kblocks_local = K / block_k (assuming K is divisible by block_k); The shape of the target continuous scaling factor is [nblocks_local, kblocks_local], and the step size is [kblocks_local, 1]. The source stride is read from scale_stride: src_stride_n=scale_stride[0], src_stride_k=scale_stride[1]; Target stride: dst_stride_n=kblocks_local, dst_stride_k=1; Global block offset: global_n_block_offset = global_n_offset / block_n, global_k_block_offset = global_k_offset / block_k; Workspace size: workspace_bytes = sizeof(dtype) × nblocks_local × kblocks_local. Here, sizeof(dtype) represents the number of bytes occupied by each element in the scaling factor tensor. The system obtains this value based on the scaling factor's data type dtype (e.g., 4 bytes for FP32, 2 bytes for FP16) through a predefined type-size mapping table or CUDA runtime functions (such as cudaDataTypeSize). If no adapter is needed, then need_adapter=false.

[0121] After the above calculations are completed, the Adapter Plan structure is obtained. The Adapter Plan describes how to convert the source scaling factor into the layout required by the candidate computational units, and includes the following fields: need_adapter (whether an adapter is needed), adapter_id (the adapter identifier), dst_scale_shape and dst_scale_stride (the shape and stride of the target scaling factor), src_stride_n and src_stride_k (the stride of the source scaling factor in N and K dimensions), dst_stride_n and dst_stride_k (the stride of the target scaling factor in N and K dimensions), global_n_block_offset and global_k_block_offset (the offsets of the global N and K blocks), and workspace_bytes (the number of bytes of GPU workspace to be allocated).

[0122] It's important to note that the Adapter Plan does not directly and permanently store a single workspace pointer, `dst_scale_ptr`, as the shared output address for all requests. The Adapter Plan only records the adaptation rules and the required workspace size. During online execution, the security dispatch layer retrieves the workspace pointer `dst_scale_ptr` exclusively for the current request from the WorkspacePool based on the current GPU device, CUDA stream, and concurrent request sequence number. This workspace is first written by the Adapter kernel within the same CUDA stream, and then read by the candidate compute unit. After a request is completed, the system records the last usage point of the workspace through a CUDA Event. Only after the Event is completed can the workspace be reused by subsequent requests. This avoids scaling factor overwriting caused by multiple concurrent requests reusing the same workspace pointer `dst_scale_ptr`.

[0123] In this embodiment, nblocks_local=256 / 128=2, kblocks_local=128 / 64=2. src_stride_n=1, src_stride_k=4 (because there are 4 N blocks globally). dst_stride_n=2, dst_stride_k=1. global_n_block_offset=256 / 128=2, global_k_block_offset=0. workspace_bytes=4×2×2=16 bytes (assuming FP32 type, each element is 4 bytes). In the generated AdapterPlan, adapter_id="scale_relayout_tp_adapter", need_adapter=true.

[0124] S304, allocate workspaces according to the adapter plan, and execute the adapter kernel to convert the scaling factor tensor into the target layout to obtain the converted scaling factor tensor for the candidate computing unit to perform correctness verification.

[0125] Based on the adapter plan, a workspace is allocated, and the CUDA adapter kernel converts the scaling factor tensor into a continuous layout, resulting in a continuous layout scaling factor tensor. This continuous layout scaling factor tensor serves as the input data that meets the requirements of the candidate computation unit, allowing the candidate computation unit to perform correctness checks.

[0126] In this embodiment, the system maintains a workspace pool to manage the allocation and reclamation of GPU memory workspaces.

[0127] This step, based on the Adapter Plan generated in step S303, retrieves the GPU memory workspace currently requested exclusively from the WorkspacePool and starts the CUDA kernel to convert the source scaling factor into a contiguous layout required by the candidate compute units. The adapter kernel simultaneously handles transpose / step views, non-contiguous scaling factors, and tensor parallel offsets, avoiding multiple memory copies.

[0128] In implementation, the `acquire_adapter_workspace(device_id, stream, fingerprint_hash, workspace_bytes)` function is first called. This function retrieves a pointer `dst_scale_ptr` for a currently requested exclusive workspace from the WorkspacePool based on the current GPU device, CUDA stream, risk fingerprint, and required workspace size. If no reusable video memory exists in the workspace pool, the CUDA runtime functions `cudaMalloc` or `cudaMallocAsync` are called to allocate video memory of size `plan.workspace_bytes`; if a workspace already confirmed by a CUDA event to be no longer used by a previous request exists, that workspace is reused.

[0129] Then the CUDA adapter kernel is invoked. This kernel uses a two-dimensional grid and two-dimensional thread block configuration, for example, each thread block contains 16×16 threads, and each thread is responsible for calculating one element of the target scaling factor. Internally, the kernel calculates the source index and performs the assignment according to the following formula: src_index=storage_offset+(n_local+global_n_block_offset)×src_stride_n+(k_local+global_k_block_offset)×src_stride_k dst_index=n_local×dst_stride_n+k_local×dst_stride_k dst[dst_index]=src_storage_base[src_index] Wherein, src_index represents the one-dimensional index position of the target element in the source scaling factor tensor in the underlying storage, dst_index represents the one-dimensional index position of the corresponding element in the target scaling factor tensor, src_storage_base is the underlying storage base address of the scaling factor, storage_base_ptr, not the already offset local logical view pointer data_ptr; storage_offset is the element offset of the logical view relative to the underlying storage base address; n_local and k_local are the local N-block index and local K-block index processed by the thread, respectively; dst_stride_n and dst_stride_k are the element stride of the target scaling factor; src_stride_n and src_stride_k are the element stride of the source scaling factor; global_n_block_offset and global_k_block_offset are collectively referred to as the global block offset.

[0130] The Adapter kernel and candidate matrix multiplication kernel are submitted to the same CUDA stream carried in the request. The sequential semantics of the same stream ensure that Adapter writes precede matrix multiplication reads. The online dispatch path does not call cudaDeviceSynchronize; it only waits for the operation related to the request to complete through cudaStreamSynchronize(stream) or CUDA Event in correctness verification, performance verification, or anomaly diagnosis modes.

[0131] At this point, the continuous scaling factor workspace exclusively for the current request is obtained, with its workspace pointer being dst_scale_ptr. This pointer is only valid within the lifetime of this request or the protection of this CUDA Event and is not written to the ACTIVE binding as a long-term field shared by all requests.

[0132] In this embodiment, a 16-byte current request exclusive workspace is obtained from the WorkspacePool. The scale_relayout_tp_adapter_kernel (a CUDA adapter kernel used to convert source scaling factor tensors into target contiguous layouts required by candidate compute units) is started, with a grid size of (2,2) and each thread block size of (16,16). The source scaling factor is physically stored as a transposed / stride view, with src_stride_n=1, src_stride_k=4, global_n_block_offset=2, and global_k_block_offset=0. This kernel calculates the source index based on storage_base_ptr and storage_offset, and writes the four values ​​corresponding to the global block positions (2,0), (2,1), (3,0), and (3,1) in the source to the local block positions (0,0), (0,1), (1,0), and (1,1) in the contiguous target. After execution, the current requested exclusive work area pointer dst_scale_ptr points to the continuous scaling factor layout [[s20,s21],[s30,s31]], where s20 represents the scaling factor value of the 2nd N block and the 0th K block globally, and so on.

[0133] S305, use the input data to perform matrix multiplication of the candidate calculation unit to obtain candidate output.

[0134] This step uses the continuously scaling factor obtained after adaptation (i.e., the continuously scaling factor tensor in the GPU memory pointed to by dst_scale_ptr in the adapter plan), combined with the original input tensor A and weight tensor B, to call the matrix multiplication kernel of the candidate computation unit to perform calculations and obtain candidate outputs.

[0135] During implementation, the `run()` method of the candidate computation unit is called. `run()` is a method that the candidate computation unit must provide when implementing the `IMatmulBackend` interface, used to actually perform matrix multiplication operations. When called, if the `need_adapter` field in the adapter plan is true, the `dst_scale_ptr` obtained and populated for the current request in step S304 is passed as the scaling factor parameter; otherwise, the original scaling factor `scale` is passed. `dst_scale_ptr` is managed by the WorkspacePool according to the request or CUDA Event lifecycle and is not a fixed pointer shared by all requests. Internally, the candidate computation unit starts its GEMM (Generic Matrix Multiplication) kernel. After execution, the output tensor `output` is filled with the matrix multiplication result as the candidate output.

[0136] In this embodiment, the `demo_block_scaled_backend::run()` method is called. This method internally calls the `launch_block_scaled_matmul` kernel (a CUDA kernel that implements block scaling matrix multiplication) and performs calculations using continuous scaling factors. This kernel uses a naive implementation (only for demonstrating the dispatch process of this invention, and its performance is not claimed), and the calculation method is as follows:

[0137] Where A[m][k] represents the element in the m-th row and k-th column of the input tensor A, B[n][k] represents the element in the n-th row and k-th column of the weight tensor B, and C[m][n] represents the element in the m-th row and n-th column of the output tensor; the summation variable k ranges from 0 to K-1. This indicates rounding down; `scale` is the scaling factor tensor. For block indexes in N dimensions, Here is the block index in the K-dimensional space. The calculation result C[m][n] is the element value in the m-th row and n-th column of the output tensor. The entire output tensor is filled with matrix multiplication results as candidate outputs. The unnormalized initial scores of various faults in the matrix multiplication results can be converted into fault type confidence scores after subsequent probability normalization processing, which are used for alarm identification and work order dispatching in the lighting operation and maintenance system.

[0138] S306, Perform a correctness check on the candidate output to obtain the correctness check result.

[0139] This step verifies the correctness of the candidate output to prevent erroneous kernels from entering the production environment. The system automatically compares the error between the candidate output and the reference output. If the error is within a preset tolerance range, the correctness verification result is considered passed, requiring no manual inspection.

[0140] Specifically, this step includes the following sub-steps:

[0141] S3061, obtain the candidate output obtained after the candidate computation unit performs matrix multiplication, and the reference output obtained by the CPU.

[0142] The reference output can be obtained through CPU three-layer loop calculation, or through representative mini-batch verification, sampled element comparison, block verification, or secure GPU reference candidate computation unit. For small-scale matrices in this embodiment, the reference output can be accurately calculated by CPU three-layer loop; for large-size matrices in real large models, correctness verification can be performed using representative mini-batch verification, sampled element comparison, block verification, or comparison with secure GPU reference candidate computation unit, avoiding CPU full matrix multiplication blocking online services.

[0143] In implementation, a temporary output buffer (GPU) is first allocated, and candidate computation units are called (same as step S305) to obtain candidate outputs. Then, the candidate outputs are copied to host memory. The reference matrix multiplication is implemented on the CPU using a three-loop system: the outer loop iterates through row indices m from 0 to M-1 (M is the number of rows in the output tensor), the middle loop iterates through column indices n from 0 to N-1 (N is the number of columns in the output tensor), and the inner loop iterates through inner product indices k from 0 to K-1 (K is the inner product dimension). For each (m, n, k), the correct scaling factor value is calculated based on the tensor parallel offsets global_n_offset and global_k_offset, and these are accumulated to obtain the reference output.

[0144] S3062, compare the error between the candidate output and the reference output to obtain the maximum absolute error and relative error.

[0145] Calculate the absolute error between the candidate output and the reference output element by element, and update the maximum error and the sum of accumulated errors. Calculate the mean absolute error (sum of accumulated errors divided by the total number of elements) and the relative error (sum of accumulated errors divided by the sum of the absolute values ​​of the reference values).

[0146] S3063, if the maximum absolute error is less than or equal to the absolute tolerance threshold, or the relative error is less than or equal to the relative tolerance threshold, then the candidate calculation unit is determined to have passed the correctness check.

[0147] If the maximum absolute error max_abs_error ≤ atol (e.g., atol = 1e-3) or the relative error relative_error ≤ rtol (e.g., rtol = 1e-3), then set accepted to true; otherwise, set it to false. Finally, release the temporary buffer.

[0148] The final result is a GuardProfile structure. GuardProfile records the comparison results of the correctness check, containing the following fields: max_abs_error (maximum absolute error), mean_abs_error (mean absolute error), relative_error (relative error), and accepted (whether the check passed). The passing condition is: max_abs_error ≤ atol or relative_error ≤ rtol.

[0149] If accepted is true, the candidate computational unit passes the correctness verification and proceeds to step S307 (performance evidence collection).

[0150] If accepted is false, the candidate computation unit is rejected, and the process returns to step S301 to continue evaluating the next candidate computation unit; if there are no more candidate computation units, an exception is thrown or the implementation is downgraded to the default implementation.

[0151] In this embodiment, the candidate output is compared with the CPU reference output. The maximum absolute error is 8.2e-5, and the relative error is 1.1e-4, both of which are less than the threshold. Therefore, accepted=true. If an erroneous candidate computation unit appears (e.g., tensor parallel offset is incorrectly handled), the maximum error may reach tens or even hundreds, and the correctness check will reject the candidate computation unit.

[0152] S307, collect performance evidence for candidate computing units that have passed the correctness check, the performance evidence including at least runtime latency, kernel execution summary and hardware counter.

[0153] For candidate computing units that pass the correctness check (accepted = true in step S306), this step automatically collects performance evidence of their execution under the current hardware and request configuration, which is used in subsequent steps (such as step S308) to select the optimal candidate computing unit. The performance evidence includes runtime latency, kernel execution summary, and hardware counters, supports multiple collection modes, and automatically generates original files for subsequent parsing without manual intervention.

[0154] The performance evidence collection supports multiple collection modes, including at least: runtime direct measurement mode (runtime mode), system-level tracing mode (nsys mode), and hardware indicator collection mode (ncu mode). The runtime direct measurement mode records kernel execution time through CUDA events; The system-level tracking mode generates a timeline report; The hardware metrics collection mode collects instruction throughput, cache hit rate, and memory access efficiency.

[0155] Among them, runtime mode can be used for low-overhead online slow path sampling; nsys mode and ncu mode are mainly used for offline verification, gray-scale traffic windows, or anomaly diagnosis tasks, and are not forced to execute in the synchronous path of the main production request. Specifically, the goals of the above three modes are as follows: 1. Runtime mode: Performs multiple rounds of timing based on CUDA events to obtain kernel latency distribution; 2. nsys mode: Invokes NVIDIA Nsight Systems tools to obtain system-level timelines and kernel execution summaries; 3. NCU mode: Calls the NVIDIA Nsight Compute tool to obtain kernel-level hardware counter metrics (such as SM utilization, video memory throughput, etc.).

[0156] The system specifies one or more modes through external configuration parameters (such as the command-line parameter --mode or the environment variable EVIDENCE_MODE). The implementation methods of the three modes are described below.

[0157] 1. Evidence collection using direct measurement mode (runtime mode).

[0158] Multi-round timing is performed using CUDA events. First, a warm-up is executed several times (e.g., 5 times) to bring the GPU to a stable state. Then, the process is repeated a specified number of times (e.g., 30 times), recording the number of milliseconds from cudaEventRecord(start) to cudaEventRecord(stop) each time. Finally, the average, 50th percentile (p50), 95th percentile (p95), and 99th percentile (p99) are calculated, outputting a performance evidence line in the format "PERFevidencefingerprint=...avg_ms=...p50_ms=...p95_ms=...p99_ms=...kernel_calls=...", recording information such as risk fingerprint, average latency, latency at each percentile, and kernel call count. This results in a log file containing the performance evidence line generated in runtime mode, which serves as one of the performance evidence files.

[0159] In this embodiment, the process was executed 30 times, and the average latency was measured to be 0.234ms, with a p99 latency of 0.287ms.

[0160] 2. System-level tracing mode (nsys mode) evidence collection.

[0161] The same matrix multiplication task was started using the Nsight Systems tool (command-line nsys profile) to generate an .nsys-rep report file. Then, nsys stats --reportcuda_gpu_kern_sum was run to extract the CUDA kernel summary and save it as a text file. This yielded both the report file and the kernel summary statistics file, which together form part of the performance evidence file generated in nsys mode.

[0162] In this embodiment, the generated report shows the number of calls and average time taken for block_scaled_matmul_kernel and scale_relayout_tp_adapter_kernel.

[0163] 3. Hardware indicator acquisition mode (NCU mode) verification.

[0164] Use the Nsight Compute tool (command line: ncu --csv --metrics) to collect metrics for specified hardware counters. The collected metrics include: 1. sm__throughput.avg.pct_of_peak_sustained_elapsed: Percentage of peak throughput for streaming multiprocessor; 2. dram__throughput.avg.pct_of_peak_sustained_elapsed: Percentage of video memory throughput relative to peak value; 3. smsp__warps_active.avg.pct_of_peak_sustained_active: Percentage of active warps.

[0165] The command outputs a metric data file containing the aforementioned hardware counter values ​​as one of the performance evidence files.

[0166] In this embodiment, the measured SM utilization rate was 78.3%, the DRAM throughput was 412GB / s (peak utilization 45%), and the utilization rate was 82.1%.

[0167] Step S307 ultimately yields performance evidence files generated in three modes: log files generated in runtime mode, report and kernel summary files generated in nsys mode, and metric data files generated in ncu mode. These files will serve as input for performance evidence step S308.

[0168] S308, the collected performance evidence is parsed into a unified format performance evidence record.

[0169] This step parses the performance evidence files of different formats collected in step S307 into a unified PerfEvidenceRecord data structure as performance evidence records. Unified field definitions are used, missing metrics are not falsified, and original rows are retained for auditing purposes.

[0170] In implementation, the parsing script is invoked to process the performance evidence files generated in the three modes respectively:

[0171] 1. For log files generated in runtime mode, extract the performance evidence lines that contain the "PERFevidence" identifier, and parse out values ​​such as average latency avg_ms, latency at each percentile (p50_ms, p95_ms, p99_ms), and kernel call count kernel_calls.

[0172] 2. For the reports and kernel summary files generated in nsys mode, use regular expressions to match the lines containing the specified kernel name (e.g., block_scaled_matmul), extract the "Avg(ns)" column (average time in nanoseconds) and convert it to milliseconds, and simultaneously accumulate the number of kernel calls.

[0173] 3. For the indicator data file generated in NCU mode, search for the indicator name (e.g., sm__throughput.avg.pct_of_peak_sustained_elapsed) line by line and extract the corresponding value.

[0174] If a field cannot be parsed (e.g., nsys did not find a matching kernel), the value of that field is set to -1, and the reason for the missing data is recorded in the `metric_missing` field. This ultimately results in an instance of the `PerfEvidenceRecord` structure (e.g., a JSON object). `PerfEvidenceRecord` is a unified performance evidence record that integrates metrics from three collection modes, including the following fields: `avg_ms` (average latency, milliseconds), `p50_ms` (50th percentile latency), `p95_ms` (95th percentile latency), `p99_ms` (99th percentile latency), `kernel_calls` (kernel call count), `sm_util_pct` (SM utilization percentage), `dram_throughput_gbps` (memory throughput, GB / s), and `occupancy_pct` (thread bundle occupancy percentage). Additionally, it includes a `metric_missing` flag (recording the reason for missing metrics) and `raw_lines` (retained raw data lines for auditing). Metrics that cannot be obtained are not falsified.

[0175] In this embodiment, the parsed values ​​are: avg_ms=0.234, p99_ms=0.287, kernel_calls=2, sm_util_pct=78.3, dram_throughput_gbps=412, occupancy_pct=82.1. All metrics were successfully parsed, with no missing metrics.

[0176] S309, associate the performance evidence record with the risk fingerprint and the current operating environment.

[0177] An evidence key is generated based on the risk fingerprint, the identifier of the candidate computing unit, the adapter identifier, and the runtime environment signature. The performance evidence record is then bound to the evidence key and stored in the evidence database.

[0178] Specifically, in this embodiment, the EvidenceKey further incorporates the candidate computing unit identifier, adapter identifier, candidate computing unit version signature, GPU architecture, and CUDA / driver environment signature into the risk fingerprint. This is used to distinguish performance evidence generated by different candidate computing units, different adapters, different candidate computing unit versions, and different operating environments under the same request risk characteristics. Then, the following six fields are concatenated in a fixed order, separated by a vertical bar "|": fingerprint_hash (risk fingerprint hash value) backend_id (candidate computation unit identifier) adapter_id (Adapter Identifier) backend_version_signature (candidate computation unit version signature) gpu_arch (GPU architecture, such as sm89) driver_cuda_signature (CUDA / driver environment signature) The concatenated string has the following format: The string `fingerprint_hash|backend_id|adapter_id|backend_version_signature|gpu_arch|driver_cuda_signature` is the Evidence Key. The SHA256 hash of this string is calculated, and the first 16 bits are used as the `evidence_id`. Then, the `PerfEvidenceRecord` (containing the parsed performance metrics, evidence key, status, etc.) is written to the evidence repository (e.g., the JSON file `evidence_store.json` or the SQLite table `perf_evidence_record`), and the evidence key and evidence ID are recorded. This key ensures that evidence from different environments or different candidate computation units under the same risk fingerprint can be distinguished. The evidence repository, parallel to the binding table, is used to store performance evidence records, and a state machine is used to manage the lifecycle of the records.

[0179] In this embodiment, the values ​​of each field are as follows: fingerprint_hash="fp_a1b2c3d4e5f6" backend_id="demo_block_scaled_backend" adapter_id="scale_relayout_tp_adapter" backend_version_signature="demo_backend_v2" gpu_arch="sm89" driver_cuda_signature="cuda12.8_driver570" The generated evidence key is: The evidence key "fp_a1b2c3d4e5f6|demo_block_scaled_backend|scale_relayout_tp_adapter|demo_backend_v2|sm89|cuda12.8_driver570" has the evidence ID "ev_8f3a2b1c". This evidence key serves as the unique identifier for the current performance evidence record.

[0180] S310, Based on the associated performance evidence, a performance degradation comparison is performed on the candidate computing unit to obtain the performance degradation comparison result.

[0181] The performance of the candidate computing unit is compared with at least one performance baseline, which includes: a predefined baseline of the security candidate computing unit, the historical median performance of the same risk fingerprint, and the previous version record of the same candidate computing unit.

[0182] 1. Predefined baseline for safe candidate computing units: a predefined acceptable performance threshold or the measured latency of the CPU reference implementation.

[0183] 2. Historical Median of the Same Risk Fingerprint: Calculated from performance evidence records in the evidence base that have the same risk fingerprint and are in a valid or bound state. Specifically, all performance evidence records with the same risk fingerprint and a state of ACTIVE or BOUND are retrieved from the evidence base, and the average latency avg_ms (milliseconds) of each record is extracted as a performance metric. These avg_ms values ​​are sorted from smallest to largest, and the median is calculated: if the number of records is odd, the middle value is used; if the number of records is even, the arithmetic mean of the two middle values ​​is used. This median serves as the performance baseline, used for comparison with the average latency of the current candidate computation unit. If the number of valid records with the same risk fingerprint in the evidence base is less than a preset threshold (e.g., less than 3), the historical median is not used as the baseline; instead, the baseline of the security candidate computation unit or the configured SLA threshold is used.

[0184] 3. Previous version record of the same candidate computing unit: Using the candidate computing unit identifier in the evidence key, query the evidence database for the most recent valid record with the same computing unit identifier but a different version.

[0185] In implementation, the regression percentage of the candidate computing unit's performance relative to each baseline is calculated using the formula (candidate-baseline) / baseline, where candidate represents the performance index value of the current candidate computing unit and baseline represents the performance baseline value currently used.

[0186] If the performance degradation exceeds a preset threshold, the candidate computation unit is rejected. Specifically, if any regression percentage exceeds a preset threshold (in this embodiment, the threshold for average latency (avg) is 10%, and the threshold for the 99th percentile latency (p99) is 15%), performance degradation is determined, and the candidate computation unit is rejected; otherwise, it is accepted. The preset thresholds include the average latency threshold and the 99th percentile latency threshold. If the candidate computation unit is rejected, the process returns to step S301 to continue evaluating the next candidate computation unit.

[0187] When there is no historical record for the first time, the system uses the measured latency of the safe candidate computation unit, the configured SLA threshold, or historical statistical values ​​of the same shape as the initial baseline. If no reliable baseline exists, the candidate computation unit only enters the BOUND state and is allowed to be upgraded to ACTIVE after accumulating a preset number of samples. The historical median of the same risk fingerprint is calculated from the evidence base by querying performance evidence records with the same fingerprint_hash and a state of ACTIVE or BOUND. The previous version record of the same candidate computation unit is obtained from the evidence base by querying the most recent valid record with the same backend_id but a different version. After comparison, a Boolean value "accepted" (indicating whether it passes the degradation judgment) and the percentage of each regression are obtained.

[0188] In this embodiment, since it is the first execution and there is no historical record, the system presets a baseline for the safe candidate computation unit: safe_avg = 10.0 milliseconds (which is much greater than the actual average latency of the candidate computation unit, 0.234 milliseconds), so it is determined to be without degradation. If in a subsequent execution, the average latency of the candidate computation unit for the same risk fingerprint rises to 0.28 milliseconds (the regression percentage relative to the historical median of 0.234 milliseconds is (0.28-0.234) / 0.234 ≈ 19.6%, exceeding the 10% threshold), then the candidate computation unit is rejected and will not be written into the ACTIVE binding.

[0189] S311, the candidate computation unit that passes the correctness check and the performance degradation comparison result and its associated information are stored in the binding table as a valid dispatch binding record, and the candidate computation unit is used to perform the current matrix multiplication.

[0190] For candidate computation units that pass the correctness check (step S306) and performance degradation comparison (step S310), this step uses the candidate computation unit to execute the current matrix multiplication request. The calculation result is written to the output tensor, and the current request is completed. The system then returns to the inference framework to continue subsequent calculations. Simultaneously, the binding information of the candidate computation unit is written to the local binding table, and its status is set to ACTIVE. Subsequent requests with the same risk fingerprint will directly reuse this binding, eliminating the need to repeatedly execute the slow path.

[0191] It's important to note that the ACTIVE binding stores the candidate compute unit identifier, adapter plan, workspace requirements, evidence key, and environment version signature, rather than a fixed dst_scale_ptr shared by all concurrent requests. During actual execution, the system still obtains an exclusive workspace for the current request from the WorkspacePool to ensure concurrency safety.

[0192] In implementation, a `DispatchBinding` object (dispatch binding record) is created, and the following fields are set: 1. `state` is set to `BindingState::ACTIVE` (indicating currently valid); 2. `created_at_ms` and `last_used_at_ms` are set to the current timestamp (milliseconds); 3. `hit_count` is initialized to 0. This object is then stored in the `binding_table_`, with the key being the risk fingerprint `fingerprint_hash` and the value being the binding record. The binding table maintains the mapping relationship from risk fingerprints to candidate computation units and corresponding adapter plans.

[0193] In this embodiment, the write binding is as follows: the risk fingerprint "fp_a1b2c3d4e5f6" is mapped to the candidate computing unit demo_block_scaled_backend and its corresponding adapter plan (adapter_id="scale_relayout_tp_adapter").

[0194] When a subsequent matrix multiplication request with the same risk characteristics is received, step S2 will directly retrieve the ACTIVE binding from the local binding table and perform fast dispatch. This means it will directly use the candidate computation unit and adapter plan of the binding to perform the matrix multiplication operation (the same candidate computation unit used in this request), and write the calculation result to the output tensor. This process skips steps S301 to S311 (all steps of the slow path). After fast dispatch is completed, the dispatch layer finishes processing this request and returns to the inference framework to continue subsequent calculations.

[0195] S4. If no candidate computation unit passes both the correctness check and the performance degradation comparison, call the fallback computation unit to perform matrix multiplication.

[0196] When all candidate computation units fail the correctness check and performance degradation comparison, the risk fingerprint, failure reason, and rollback time are recorded to ensure that output correctness takes precedence over performance.

[0197] The fallback computation unit is typically a CPU reference implementation or a GPU implementation that guarantees correctness but has lower performance. Its characteristic is that it does not rely on any optimizations, only guaranteeing the correctness of the calculation result. The execution of the fallback computation unit bypasses the fast or slow path of this method, directly performing matrix multiplication.

[0198] In implementation, if none of the candidate computation units pass verification, the registered `run()` method of `fallback_backend` is called. This method receives the original matrix multiplication request `MatmulDispatchRequest`, processes it, and directly obtains the correct output tensor `output`.

[0199] In this embodiment, if no candidate computation unit passes verification, a fallback computation unit is invoked to compute the output. The fallback computation unit can be a CPU reference implementation or a known correct but low-performance GPU general-purpose matrix multiplication implementation. In small-scale verification or offline diagnostics, the CPU reference implementation can be invoked to compute the output and the result copied back to the GPU output tensor. In online inference services, a safe GPU fallback candidate computation unit is invoked first to avoid blocking production requests due to CPU full matrix multiplication and host-to-device copying. The system also records warning logs, risk fingerprints, failed candidate computation units, failure reasons, and fallback time for maintenance personnel to investigate.

[0200] S5, after the matrix multiplication request is completed, if the risk-related characteristics corresponding to the valid distribution binding record change, causing the newly generated risk fingerprint to be inconsistent with the risk fingerprint of the valid distribution binding record, or if the operating environment parameters corresponding to the valid distribution binding record change, the valid distribution binding record is marked as invalid.

[0201] A state machine is used to manage dispatch binding records and performance evidence records. The states include at least valid and invalid states. When a risk-related characteristic or a change in the operating environment is detected, the corresponding dispatch binding record is changed from a valid state to an invalid state. Historical data is not deleted; the state machine is used to transition the state to preserve the complete audit trail. The states of the state machine are as follows: CANDIDATE: Candidate record, not yet resolved; meaning: candidate record, not yet resolved. PARSED: Parsed, performance evidence has been standardized to the PerfEvidenceRecord format; meaning: the performance evidence has been parsed and standardized to a formatted performance evidence record. BOUND: Bound, performance evidence has been generated and bound to risk fingerprint, candidate compute unit, adapter, version and hardware environment; EvidenceKey means: It has been bound, and the performance evidence has been bound to the risk fingerprint, candidate compute unit identifier, adapter identifier, candidate compute unit version signature, GPU architecture, CUDA and driver signature to generate an evidence key; ACTIVE: Currently valid and can be used for rapid distribution; meaning: active, currently valid, and can be used for rapid distribution. STALE: Expired but retained; cannot be used for rapid deployment, used for historical comparison; meaning: invalid, became invalid due to changes in risk fields or runtime environment versions; ARCHIVED: Historical archives have been replaced by new ACTIVE records and are used only for auditing purposes; meaning: archived, historical data is retained for auditing purposes. REJECTED: Rejected because the correctness check failed or the performance degradation comparison failed; meaning: rejected because the correctness check or performance degradation comparison failed.

[0202] When a change is detected in any field of the metadata or the version of the runtime environment signature, the binding record of the corresponding risk fingerprint is changed from ACTIVE to STALE, and the corresponding record in the evidence base is changed to STALE, ARCHIVED, or REJECTED, while the original data is retained for auditing.

[0203] Change events include: scaling factor stride changes, scaling factor storage_offset changes, scaling factor non-standard layout markup changes, block shape changes, tensor parallel offset (TPoffset) changes, candidate compute unit version upgrades, adapter version changes, GPU architecture changes, CUDA runtime version changes, CUDA Driver API version changes, NVIDIA driver version changes, correctness check failures, and performance degradation comparison rejections.

[0204] In implementation, the `invalidate(fingerprint_hash, reason)` method is called externally (where `reason` represents the reason for invalidation, such as changes in the environment signature, upgrades to the candidate computation unit version, or failure of correctness verification). This iterates through the binding table and evidence database, updates the status, and records the reason for invalidation and the time. The specific rules are as follows: Change the binding status of the corresponding risk fingerprint from ACTIVE to STALE; Change the record status of the corresponding evidence_key in the evidence database from BOUND or ACTIVE to STALE or REJECTED (depending on the reason for failure). Once a new binding is established, the corresponding old STALE record will be converted to ARCHIVED.

[0205] In this embodiment, assuming that the subsequent candidate computing unit version is upgraded to demo_backend_v3, the system detects a change in backend_version_signature and calls invalidate to change the binding status of the original risk fingerprint "fp_a1b2c3d4e5f6" from ACTIVE to STALE, and the original evidence record to ARCHIVED. The next time a matrix multiplication request with the same risk feature is received, because the binding has expired (the status is not ACTIVE), it will re-enter the slow path (step S301), re-acquire evidence, and establish a new binding.

[0206] The above is a complete description of Embodiment 1 of the present invention. This embodiment details how to complete a closed loop from receiving a matrix multiplication request, through risk fingerprint generation, binding query, slow path filtering, risk detection, adapter planning and execution, correctness verification, performance verification, evidence binding, degradation comparison, binding activation, cache invalidation management, to security rollback. Those skilled in the art, based on the description in this embodiment and in conjunction with the publicly available engineering code package, can implement the present invention without any inventive effort. Although this embodiment uses a smart lighting gating projection layer as an example, the method can be directly extended to other inference scenarios that rely on block scaling matrix multiplication (such as alarm timing prediction, energy consumption optimization strategy recommendation, etc.).

[0207] Example 2

[0208] This embodiment provides a kernel-secure dispatch system for large-scale illumination matrix multiplication, used to implement the kernel-secure dispatch method for large-scale illumination matrix multiplication described in Embodiment 1. The system includes:

[0209] 1. A risk fingerprint generation unit, used to respond to a matrix multiplication request, collect the risk-related features of the matrix multiplication request, the risk-related features including at least the step size of the scaling factor tensor, the continuity marker and the transpose view marker, and the tensor parallel offset, to generate a risk fingerprint.

[0210] The risk fingerprint generation unit responds to the matrix multiplication request, collects multi-dimensional features related to distribution security (including tensor layout attributes, storage attributes, tensor parallel offset, and runtime environment information) from the matrix multiplication request, and generates the corresponding risk fingerprint after standardizing the above features into a unified format. This risk fingerprint is used as a unique key for subsequent binding queries.

[0211] 2. The dispatch control unit is configured to select a fast path and directly use the candidate computation unit to perform matrix multiplication if there is a candidate computation unit associated with and verified by the risk fingerprint; otherwise, a slow path is selected.

[0212] The distribution control unit queries the binding storage unit using the risk fingerprint as the key. If a distribution binding record associated with the risk fingerprint and in a valid state exists, the candidate calculation unit in the binding record is directly reused to perform matrix multiplication, skipping the verification process; otherwise, the process switches to the slow path, and the slow path processing unit takes over the subsequent operations.

[0213] 3. A slow path processing unit, configured to perform risk detection on candidate computing units in the slow path to obtain a risk set, the risk set including repairable risks and unrepairable risks; generate an adapter plan for candidate computing units containing only repairable risks; convert the scaling factor tensor into a target layout according to the adapter plan; perform correctness verification on the candidate computing units using the converted scaling factor tensor; filter candidate computing units that pass the correctness verification through performance degradation comparison; mark the filtered candidate computing units as verified and bind them to the risk fingerprint to form a valid dispatch binding record; and use the candidate computing units that pass the performance degradation comparison to perform the matrix multiplication.

[0214] The slow path processing unit includes: 1) Candidate filtering module, used to filter candidate computing units that meet the current request conditions from the registered candidate computing units based on basic constraints such as the current GPU architecture, data type and block shape; 2) Risk detection module, used to perform risk detection on candidate computing units to obtain a risk set including repairable and unrepairable risks; 3) Adapter module, used to generate an adapter plan based on the risk set for candidate computational units containing only repairable risks, allocate work areas according to the adapter plan and execute adapter kernels to transform tensor layout; 4) Candidate execution module, used to call candidate calculation unit to perform matrix multiplication and obtain candidate output; 5) Correctness verification module, used to obtain candidate output and reference output, compare the error between candidate output and reference output, and determine whether candidate calculation unit passes the correctness verification based on whether the error is within the preset tolerance range; 6) Performance evidence collection and binding module, used to collect performance evidence from candidate computing units that have passed the correctness verification and parse it into a unified format performance evidence record. The performance evidence record is then associated with the risk fingerprint and the current operating environment to generate an evidence key and stored in the binding storage unit. 7) Degradation comparison module, used to compare the performance of candidate computing units with a preset performance baseline based on the associated performance evidence, and to determine whether performance degradation has occurred; 8) Binding activation module, used to mark the candidate computing unit that has passed the correctness check and performance degradation comparison as verified and associated with the risk fingerprint and store it in the binding storage unit, and at the same time use the candidate computing unit to perform the current matrix multiplication.

[0215] 4. The safety fallback unit is configured to call the fallback unit to perform matrix multiplication if no candidate computation unit passes both the correctness check and the performance degradation comparison.

[0216] After the slow path processing unit traverses all candidate computation units, if no candidate passes both the correctness check and performance degradation comparison simultaneously, the safety rollback unit calls the fallback computation unit to perform matrix multiplication, ensuring that the inference service can always output the correct result and recording failure information for maintenance personnel to troubleshoot.

[0217] 5. The detection management unit is configured to mark the valid distribution binding record as invalid after the matrix multiplication request processing is completed, when the risk-related characteristics corresponding to the valid distribution binding record change, causing the newly generated risk fingerprint to be inconsistent with the risk fingerprint of the valid distribution binding record, or when the operating environment parameters corresponding to the valid distribution binding record change.

[0218] The detection management unit continuously monitors risk-related characteristics and changes in the operating environment version. When changes are detected in key fields such as tensor layout attributes, tensor parallel offset, candidate computation unit version, CUDA version, or driver version, the detection management unit changes the corresponding dispatch binding record in the binding storage unit from a valid state to an invalid state, and at the same time changes the corresponding performance evidence record to an archived or invalid state to prevent outdated data from continuing to affect online dispatch, while retaining a complete audit trail.

[0219] The aforementioned units work together to achieve secure distribution of the matrix multiplication kernel of the large lighting model.

[0220] For details on the specific implementation of the operations performed by the above units, please refer to the relevant step descriptions in Example 1, which will not be repeated here.

[0221] Example 3

[0222] This embodiment provides an electronic device, the electronic device comprising: A memory for storing computer programs, the memory including non-volatile memory and random access memory, the non-volatile memory for storing system programs and the random access memory for caching runtime data; A processor is configured to execute a computer program stored in the memory to implement the steps of a kernel-secure dispatch method for a large illumination model matrix multiplication as described in Embodiment 1.

[0223] For details on the specific implementation of each step and related explanations, please refer to Example 1, which will not be repeated here.

[0224] The processor can be a general-purpose processor, including a central processing unit, a network processor, etc.; it can also be a digital signal processor, an application-specific integrated circuit, a field-programmable gate array, or other programmable logic devices.

[0225] Example 4

[0226] This embodiment provides a computer-readable storage medium storing a computer program that, when executed by a processor, implements the steps of a kernel-secure dispatch method for a large-scale illumination model matrix multiplication as described in Embodiment 1.

[0227] For details on the specific implementation of each step of this method and related explanations, please refer to Example 1, which will not be elaborated here.

[0228] The computer-readable storage medium can be any tangible medium that contains or stores program instructions, such as: USB flash drive, portable hard drive, read-only memory (ROM), random access memory (RAM), magnetic disk or optical disk, etc.

[0229] The above description is merely an embodiment of the present invention and is not intended to limit the invention. Various modifications and variations can be made to the present invention by those skilled in the art. Any modifications, equivalent substitutions, improvements, etc., made within the spirit and principle of the present invention should be included within the scope of the claims of the present invention.< / imatmulbackend>

Claims

1. A kernel-secure distribution method for matrix multiplication of large illumination models, characterized in that, include: In response to a matrix multiplication request, risk-related features of the matrix multiplication request are collected. These risk-related features include at least the step size, continuity marker, and transpose view marker of the scaling factor tensor, as well as the tensor parallel offset, to generate a risk fingerprint. If a candidate computation unit with a unique candidate computation unit identifier exists and has been verified as associated with the risk fingerprint, the fast path is entered, and matrix multiplication is performed directly using the candidate computation unit; otherwise, the slow path is entered. In the slow path, risk detection is performed on candidate computing units to obtain a risk set, which includes repairable and unrepairable risks. An adapter plan is generated for candidate computing units containing only repairable risks. The scaling factor tensor is converted into a target layout according to the adapter plan. The correctness verification of the candidate computing units is performed using the converted scaling factor tensor. Candidate computing units that pass the correctness verification are selected by performance degradation comparison. The selected candidate computing units are marked as verified and bound to the risk fingerprint to form a valid dispatch binding record. The matrix multiplication is performed using the candidate computing units that pass the performance degradation comparison. If no candidate computation unit passes both the correctness check and the performance degradation comparison, the fallback computation unit is called to perform matrix multiplication. After the matrix multiplication request is completed, if the risk-related characteristics corresponding to the valid distribution binding record change, causing the newly generated risk fingerprint to be inconsistent with the risk fingerprint of the valid distribution binding record, or if the operating environment parameters corresponding to the valid distribution binding record change, the valid distribution binding record will be marked as invalid.

2. The method according to claim 1, characterized in that, The response matrix multiplication request collects risk-related features of the matrix multiplication request and generates a risk fingerprint, including: The risk-related features are normalized into strings in a fixed order, and the hash value of the strings is calculated to obtain the risk fingerprint; The risk-related features include: metadata of the input tensor, metadata of the weight tensor, metadata of the scaling factor tensor, metadata of the output tensor, and runtime environment signature.

3. The method according to claim 1, characterized in that, The adapter generation plan includes: For a candidate computing unit that contains only the repairable risks, an adapter plan is generated based on the risk set, and the adapter plan corresponds to a unique adapter identifier; The adapter plan allocates the work area and executes the adapter kernel to convert the scaling factor tensor into the target layout, obtaining the converted scaling factor tensor for the candidate computing unit to perform correctness verification. The adapter plan is used to determine: a scaling factor tensor and the workspace size required to perform the adaptation operation, wherein the scaling factor tensor includes the target layout, target step size, and global block offset.

4. The method according to claim 1, characterized in that, The correctness check includes: Obtain the candidate outputs obtained after the candidate computation unit performs matrix multiplication, and the reference outputs obtained by CPU computation; By comparing the error between the candidate output and the reference output, the maximum absolute error and the relative error are obtained; If the maximum absolute error is less than or equal to the absolute tolerance threshold, or the relative error is less than or equal to the relative tolerance threshold, then the candidate computing unit is determined to have passed the correctness check.

5. The method according to claim 1, characterized in that, The performance degradation comparison includes: The performance of the candidate computing unit is compared with at least one performance baseline, which includes: a predefined baseline of the security candidate computing unit, the historical median performance of the same risk fingerprint, and the previous version record of the same candidate computing unit; If the performance degradation exceeds a preset threshold, the test will fail.

6. The method according to claim 2 or 3, characterized in that, Also includes: Performance evidence is collected from candidate computing units that pass the correctness check, and the collected performance evidence is parsed into a unified format performance evidence record. The performance evidence includes at least runtime latency, kernel execution summary and hardware counter metrics. Associating the performance evidence record with the risk fingerprint and the current operating environment includes: generating an evidence key based on the risk fingerprint, the candidate computing unit identifier, the adapter identifier, and the operating environment signature; binding the performance evidence record with the evidence key and storing it in the evidence database.

7. The method according to claim 6, characterized in that, When the risk-related characteristics corresponding to the valid distribution binding record change, causing the newly generated risk fingerprint to be inconsistent with the risk fingerprint of the valid distribution binding record, or when the operating environment parameters corresponding to the valid distribution binding record change, marking the valid distribution binding record as invalid includes: A state machine is used to manage dispatch binding records and performance evidence records. The states include at least valid states and invalid states. When the risk-related characteristics or the operating environment are detected to change, the corresponding dispatch binding record is changed from a valid state to an invalid state. When all candidate computation units fail the correctness check and performance degradation comparison, record the risk fingerprint, the reason for failure, and the rollback time.

8. A kernel-secure dispatch system for large-scale illumination matrix multiplication, characterized in that, include: A risk fingerprint generation unit is used to respond to a matrix multiplication request, collect the risk-related features of the matrix multiplication request, and the risk-related features include at least the step size, continuity marker and transpose view marker of the scaling factor tensor, and the tensor parallel offset, to generate a risk fingerprint; The dispatch control unit is configured to select a fast path and directly use the candidate computation unit to perform matrix multiplication if there is a candidate computation unit associated with and verified by the risk fingerprint; otherwise, it selects a slow path. A slow path processing unit is configured to perform risk detection on candidate computing units in the slow path to obtain a risk set, the risk set including repairable risks and unrepairable risks; An adapter plan is generated for candidate compute units containing only repairable risks. The scaling factor tensor is converted into a target layout according to the adapter plan. The correctness verification of the candidate compute units is performed using the converted scaling factor tensor. Candidate compute units that pass the correctness verification are selected by performance degradation comparison. The selected candidate compute units are marked as verified and bound to the risk fingerprint to form a valid dispatch binding record. The matrix multiplication is performed using the candidate compute units that pass the performance degradation comparison. The safety fallback unit is configured to call the fallback unit to perform matrix multiplication if no candidate computation unit passes both the correctness check and the performance degradation comparison. The detection management unit is configured to mark the valid distribution binding record as invalid after the matrix multiplication request processing is completed, when the risk-related characteristics corresponding to the valid distribution binding record change, causing the newly generated risk fingerprint to be inconsistent with the risk fingerprint of the valid distribution binding record, or when the operating environment parameters corresponding to the valid distribution binding record change.

9. An electronic device, characterized in that, include: Memory, used to store computer programs; A processor for executing a computer program stored in the memory to implement the steps of the method according to any one of claims 1 to 7.

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

Images