A lookup table method for half-precision floating-point number FP16

Abstract

This invention provides a lookup table method for half-precision floating-point number FP16, including: S1, 1) NNA quantization input; 2) Int16 input: 16 bits output to Lbit, S=0 indicates a positive number, V=0 indicates no overflow; Int16 input, FP16 input: converted to FPi; lookup table value FPi; obtain offset(FPPi), scale_n(FPPi), scale_p(FPPi); add offset, according to the result: positive multiply by scale_p, negative multiply by scale_n; obtain: absolute value converted to integer, take the lower 16 bits and output to Lbit, overflow 16 bits V=1, otherwise V=0, output overflow to V, output sign bit to S; S2, obtain SVLbit 16-bit positive integer; S3, according to LU_Table, SAT_T0, IDX0_T, get: V=1: take the value SAT_T0 (S=1) or SAT_T1 (S=0), V=0: take the high 5 bits of Lbit as idx, S=1: if idx≥IDX0_T1 take the value SAT_T0, S=0: idx+=IDX0_T1, if idx>31 take the value SAT_T1, take the low 11 bits of Lbit as dx, use idx and dx to look up the table and calculate y; S4, lookup result y (FPin); true value output: convert to FP16; quantization output: get ofbit, convert to integer and saturate to ofbit unsigned number; S5, NNA quantization output or XOR RAM port.

Description

Technical Field

[0001] This invention belongs to the field of neural network compression technology, and specifically relates to a lookup table method for half-precision floating-point number FP16. Background Technology

[0002] In recent years, neural networks have become a research hotspot in the field of artificial intelligence. Trained neural network models can be used in technologies such as computer vision and speech recognition.

[0003] FP16 is a half-precision floating-point number format, using 5 bits to represent the exponent and 10 bits to represent the decimal, occupying 2 bytes. Compared to FP32, FP16 consumes only half the memory, making it a more suitable data format for AI computations on mobile devices. While offering lower precision than FP32, it reduces storage space and computational overhead. It is primarily used for computationally intensive tasks such as deep learning and machine learning.

[0004] LUT stands for Explicit Look-Up Table. In the entry generation phase, the output data is pre-calculated and written into the entries. In the lookup phase, the table is looked up for each input signal to find the corresponding output. For complex calculations, the lookup time is much shorter than the computation time, thus improving efficiency.

[0005] However, in the field of model compression, there is currently no effective and systematic FP16 lookup table implementation method that can explain the generation and lookup of table entries, compress the model as much as possible while ensuring model accuracy, improve the model's running efficiency, and help the model be deployed quickly.

[0006] NNA (Neural Network Algorithm) is an algorithm that simulates the information transmission of an animal's neural network through logical reasoning. It is also commonly used in artificial intelligence, where information is stored in a distributed manner, and information processing is achieved through dynamic interactions similar to those between neurons.

[0007] ORAM (Oblivious Random Access Machine) is an encryption scheme that can completely hide data access patterns during I / O operations. ORAM transforms a single file access request into multiple file access requests, thus obfuscating information such as the probability and pattern of file access. By accessing files unrelated to keywords, ORAM disrupts the correlation between probability and keywords, obfuscating access patterns and thereby compromising user privacy. Summary of the Invention

[0008] To address the aforementioned issues, this application aims to propose an implementation method for image and speech recognition technologies, particularly for FP16 lookup tables. The method elaborates on the generation and retrieval of table entries, aiming to compress the model as much as possible while ensuring model accuracy, thereby improving its operational efficiency and facilitating rapid model deployment.

[0009] Specifically, this invention provides a lookup table method for half-precision floating-point numbers (FP16), applicable to image and speech recognition technologies, comprising the following steps:

[0010] S1 performs NNA quantization input, Int16 input, and FP16 input respectively:

[0011] 1) NNA quantization input, i.e., the input of NNA:

[0012] NNA is an on-chip core that improves operator efficiency, including: FRAM, WRAM, and MAC core. FRAM is the on-chip RAM that stores input features inside the NNA, and different chips may have different configurations; WRAM is the on-chip RAM that stores weights inside the NNA; MAC is the multiply-accumulate (MAC) operation; NNA quantization input refers to the input of the NNA.

[0013] 2) Read ORAM address and length, read ORAM port; Int16 input: Int16 means 16-bit integer, which is equivalent to short occupying 2 bytes -32768 to 32767; 16-bit output to Lbit, i.e., the lower 16 bits, S=0, indicating a positive number, V=0, indicating no overflow;

[0014] 3) Read ORAM address and length, read ORAM port; Int16 input, FP16 input, where FP16 half-precision floating-point number uses 5 bits to represent the exponent and 10 bits to represent the fraction, occupying 2 bytes; further includes:

[0015] 3.1) Convert to FPIN;

[0016] 3.2) Lookup table value (FPin);

[0017] 3.3) Obtain offset(FPin), scale_n(FPin), and scale_p(FPin);

[0018] To activate the lookup table process using SiLU, the SiLU (sigmoid Linear Unit) activation function, also known as the Swish activation function, is a novel activation function proposed in recent years. It is defined as the input multiplied by its sigmoid function value:

[0019] SiLU(x)=x*δ(x),

[0020] in: The lookup table includes:

[0021] start;

[0022] Find the maximum and minimum values ​​of the input vector, x_min and x_max.

[0023] For x_min and x_max, divide them into table_num parts, generating table_num+1 division points; perform SiLU operation and quantization operation on table_num+1 input x to generate table_num+1 output y;

[0024] Output table entries y0 and dy based on table_num+1 outputs; generate offset, scale_n, and scale_p; end.

[0025] 3.4) Add offset, and based on the result: multiply by scale_p for positive values ​​and by scale_n for negative values; where offset is used to adjust the input lookup table value so that when looking up T0, the value is less than 0, S=1, and when looking up T1, the value is greater than or equal to 0, S=0.

[0026] scale_p, scale_n: are used to perform appropriate fixed-point quantization on the lookup values ​​T0(scale_n) and T1(scale_p) of the two regions respectively, so that the lookup value of interest falls on the appropriate entry; the appropriate fixed-point quantization is related to the bit width of quantization and is an empirical value, which is 10 when it is 8 bits;

[0027] 3.5) Based on 3.2) and 3.3), we get:

[0028] A) Convert absolute value to integer;

[0029] B) Take the lower 16 bits and output them to Lbit. If x_fixed_abs_int is the converted integer...

[0030] The code is: int mask_lbit = int(pow(2,16)-1);

[0031] int32_t lbit=int32_t(x_fixed_abs_int&mask_lbit);

[0032] lbit refers to the lower 16 bits;

[0033] C) If overflow occurs, V = 1 for 16 bits; otherwise, V = 0. Output the overflow to V and check if x_fixed_abs_int > 65535.

[0034] D) Output the sign bit to S;

[0035] S2, based on the results of steps 1), 2), and 3) in step S1, obtain SVLbit, a 16-bit positive integer;

[0036] S3, then based on LU_Table, SAT_T0, IDX0_T, we obtain:

[0037] V = 1: The value can be either SAT_T0 (S = 1) or SAT_T1 (S = 0).

[0038] V=0: Take the high 5 bits of Lbit as idx

[0039] S = 1: If idx ≥ IDX0_T1 takes the value SAT_T0,

[0040] S = 0: idx += IDX0_T1, if idx > 31, take the value SAT_T1.

[0041] Take the lower 11 bits of Lbit as dx.

[0042] Using idx and dx to look up tables and calculate y, the linear interpolation formula is:

[0043] y=(1-dx)*y0+dx*y1=dx*(y1-y0)+y0=y0+dx*dy (Formula 1)

[0044] Set up two tables, with each table containing FP16t floating-point numbers;

[0045] in,

[0046] LU_Table: A lookup table containing 32 entries, each containing 2 FP16t numbers. The lookup table is divided into two regions, T0 and T1, and can be configured as two tables or as one table. IDX0_T1: The lookup table is also divided into two regions, T0 and T1. T0's entries are [0, IDX0_T1), and T1's are [IDX0_T1, 32). If IDX0_T1 = 0, there is no T0 region. Negative lookup values ​​use T0, and positive values ​​use T1. Furthermore, IDX0_T is an index for the negative lookup table.

[0047] S4, the lookup result y(FPin) includes:

[0048] Truth output: Convert to FP16;

[0049] Quantization output: Obtain ofbit, convert it into an integer and saturate it to an unsigned ofbit number;

[0050] S5, then write the address according to the ORAM;

[0051] NNA quantization output or write to the ORAM port.

[0052] The following is the code implementation for converting an int32 to an FPIN number:

[0053] __device__uint32_t si32_2_fpin(int32_t si32)

[0054] {

[0055] if (si32 == 0) return 0;

[0056] When si32 is less than 0, sign is 1; otherwise, sign is 0. This is represented as:

[0057] int sign = si32 < 0 ? 1 : 0;

[0058] uint32_t ui32=sign? -si32:si32;

[0059] Define a variable fpin of type i32fph_t:

[0060] i32fph_t fpin;

[0061] fpin.ui=ui32_2_e6m10(ui32);

[0062] fpin.f.sign = sign;

[0063] assert(fpin.f.e+FPIN_E0<=FPIN_E_MAX);

[0064] fpin.f.e+ = FPIN_E0;

[0065] return fpin.ui;

[0066] }

[0067] The code implementation of function ui32_2_e6m10 is as follows:

[0068] __device__uint32_t ui32_2_e6m10(uint32_t ui32){

[0069] int lead_zero=LEAD_ZERO_UI32(ui32);

[0070] Assume m is in msb, i.e., min msb: uint32_t mantis = ui32< <lead_zero;

[0071] All 12 bits of msb are 1: const uint32_t msb12one = 0xfff00000;

[0072] If msb12one&mantis is always equal to msb12one, then execute: This is represented as: if((msb12one&mantis)==msb12one){

[0073] Rounding, the condition is true when the first 12 bits of mantis are all 0. (Rounding cause C:)

[0074] lead_zero--;

[0075] mantis >>=1;

[0076] }

[0077] assert(lead_zero>=0&&lead_zero<=31);

[0078] m in the lower 12 bits: mantis >> = 20;

[0079] Take the 12th bit of mantis as rnd, the rounding bit:

[0080] uint32_t rnd = mantis&1;

[0081] m in the lower 11 bits: mantis >> = 1;

[0082] Add rnd, doing round:

[0083] mantis+ = rnd;

[0084] const uint32_t bit11=0x00000400;

[0085] In m, msb must be 1: assert((mantis&bit11)==bit11);

[0086] mantis& = ~bit11;

[0087] assert(mantis<=0x000003ff);

[0088] int exp = 31 - lead_zero;

[0089] i32fph_t fp;

[0090] fp.fe = exp;

[0091] fp.fm = mantis;

[0092] return fp.ui;

[0093] }

[0094] Step S1, 3.3), further includes:

[0095] `offset`, `scale_n`, and `scale_p` are parameters required when using a dual-table architecture: `offset` is the dividing point between negative and positive entries, and `table_neg_num` is the number of negative entries, obtained through optimization methods. The specific steps are as follows:

[0096] scale_n=std::pow(2,16) / (offset-x_min)*

[0097] table_neg_num / 32(Formula 1)

[0098] scale_p=std::pow(2,16) / (x_min-table_zeropoint_value)*

[0099] (32-table_neg_num) / 32(Formula 2)

[0100] start;

[0101] The input vector X is generated at any time, and its maximum and minimum values ​​are x_min and x_max, respectively.

[0102] Calculate the output vector Y_ori = SiLU(X);

[0103] TABLE_ITEMS_STEP=32 Initialization Loss=std::pow(2,31)K=0offset=0table_pos_num=TABLE_ITEMS_STEP;

[0104] Determine K <TABLE_ITEMS_STEP,

[0105] If NO, set table_pos_num = TABLE_ITEMS_STEP; End.

[0106] If Yes, then table_pos_num = TABLE_ITEMS_STEP - k and table_neg_num = k, generating the current table entry; look up the table to generate output Y_lut, calculate the difference between Y_lut and Y_ori (loss_current); and determine the loss_current. <loss;

[0107] If Yes, then loss = loss_current, offset = k, table_pos_num = TABLE_ITEMS_STEP - k; K = K + 1, return the judgment K. <TABLE_ITEMS_STEP;

[0108] If NO, K = K + 1, return the value of K. <TABLE_ITEMS_STEP。

[0109] In step S3, the lookup table is divided into two regions, T0 and T1, and preferably two tables are configured.

[0110] In step S3

[0111] table_dy stores delta_y = y<i+1> -y That is, each subsequent term minus the preceding term.

[0112] table_y0 stores y That is, each preceding term,

[0113] Currently, there are a total of 32 table entries, which can accommodate 2 tables;

[0114] Take the high i bits of the L-bit integer, add idx0 to get idx, where idx0 is the starting index of this table; get: dy = table_dy[idx]; y0 = table_y0[idx]

[0115] Take the lower x bits of Lbit = Lbit - i bits and shift them right by x bits to form a fixed-point number dx, that is, dx is a fraction in [0,1) with x bits as the number of significant bits;

[0116] Calculate: y = y0 + dx * dy.

[0117] The L-bit integer is currently 16 bits, and the high i-bit is currently 5 bits;

[0118] The lower xbit = Lbit - ibit bits of Lbit are currently 11 bits;

[0119] The dx is a decimal in the range [0,1), with a minimum value of 0 and a maximum value of 1-2-xbit.

[0120] Therefore, the advantage of this application is that: the present invention proposes a lookup table implementation method for FP16, and elaborates on the generation and lookup of table entries. While ensuring the accuracy of the model, it compresses the model as much as possible, improves the running efficiency of the model, and helps the model to be deployed quickly. Attached Figure Description

[0121] The accompanying drawings, which are provided to further illustrate the invention and form part of this application, are not intended to limit the scope of the invention.

[0122] Figure 1 This is a schematic diagram of the NNA structure.

[0123] Figure 2 This is a flowchart illustrating the process of obtaining offset(FPin), scale_n(FPin), and scale_p(FPin).

[0124] Figure 3 This is a schematic diagram illustrating the specific steps of the optimization method in step S1, 3), 3.3).

[0125] Figure 4 This is a schematic diagram of the overall process of this application.

[0126] Figure 5 and Figure 6 This is a schematic diagram of the code implementation of step S1, 3), 3.1). Detailed Implementation

[0127] To better understand the technical content and advantages of the present invention, the present invention will now be described in further detail with reference to the accompanying drawings.

[0128] This invention relates to a lookup table method for half-precision floating-point numbers (FP16). It proposes an implementation method for FP16 lookup tables, and elaborates on the generation and lookup of table entries. While ensuring model accuracy, it compresses the model as much as possible, improves the model's running efficiency, and helps the model to be deployed quickly.

[0129] like Figure 4 The diagram shows the overall steps for implementing the FP16 lookup proposed in this invention, including: S1, performing NNA quantization input, Int16 input, and FP16 input respectively:

[0130] 1) NNA quantization input, i.e., the input of NNA:

[0131] NNA is the core on the Ingenic chip that improves operator efficiency, such as... Figure 1 As shown, it includes: FRAM, WRAM, and MACcore. FRAM is the on-chip RAM inside the NNA that stores input features, and different chips may have different configurations; WRAM is the on-chip RAM inside the NNA that stores weights; MAC is Multiply Accumulate (MAC); NNA quantization input refers to the input of the NNA.

[0132] 2) Read ORAM address and length, read ORAM port; Int16 input: Int16 means 16-bit integer, which is equivalent to short occupying 2 bytes -32768 to 32767; 16-bit output to Lbit, i.e., the lower 16 bits, S=0, indicating a positive number, V=0, indicating no overflow;

[0133] 3) Read ORAM address and length, read ORAM port; Int16 input, FP16 input, where FP16 half-precision floating-point number uses 5 bits to represent the exponent and 10 bits to represent the fraction, occupying 2 bytes; further includes:

[0134] 3.1) Convert to FPi; such as Figure 5 and Figure 6 As shown:

[0135] The following is the code implementation for converting an int32 to an FPIN number:

[0136] __device__uint32_t si32_2_fpin(int32_t si32)

[0137] {

[0138] if (si32 == 0) return 0;

[0139] This is C code, which means that when si32 is less than 0, sign is 1, otherwise sign is 0, expressed as:

[0140] int sign = si32 < 0? 1 : 0;

[0141] uint32_t ui32 = sign? -si32 : si32;

[0142] Define a variable fpin of type i32fph_t:

[0143] i32fph_t fpin;

[0144] fpin.ui = ui32_2_e6m10(ui32);

[0145] fpin.f.sign = sign;

[0146] assert(fpin.f.e + FPIN_E0 <= FPIN_E_MAX);

[0147] fpin.f.e += FPIN_E0;

[0148] return fpin.ui;

[0149] }

[0150] The code implementation of the function ui32_2_e6m10 is as follows:

[0151] __device__ uint32_t ui32_2_e6m10(uint32_t ui32) {

[0152] int lead_zero = LEAD_ZERO_UI32(ui32);

[0153] Assume m is at the msb, i.e., min msb: uint32_t mantis = ui32 << lead_zero; The 12 bits of the msb are all 1: const uint32_t msb12one = 0xfff00000;

[0154] If msb12one & mantis is always equal to msb12one, then execute: expressed as: if ((msb12one & mantis) == msb12one) {

[0155] Rounding, the condition is true when the first 12 bits of mantis are all 0. (Rounding cause C:)

[0156] lead_zero--;

[0157] mantis >>=1;

[0158] }

[0159] assert(lead_zero>=0&&lead_zero<=31);

[0160] m in the lower 12 bits: mantis >> = 20;

[0161] Take the 12th bit of mantis as rnd, the rounding bit:

[0162] uint32_t rnd = mantis&1;

[0163] m in the lower 11 bits: mantis >> = 1;

[0164] Add rnd, doing round:

[0165] mantis+ = rnd;

[0166] const uint32_t bit11=0x00000400;

[0167] In m, msb must be 1: assert((mantis&bit11)==bit11);

[0168] mantis& = ~bit11;

[0169] assert(mantis<=0x000003ff);

[0170] int exp = 31 - lead_zero;

[0171] i32fph_t fp;

[0172] fp.fe = exp;

[0173] fp.fm = mantis;

[0174] return fp.ui;

[0175] }

[0176] 3.2) Look up the value FPIN from the table;

[0177] 3.3) Obtain offset(FPin), scale_n(FPin), and scale_p(FPin);

[0178] like Figure 2 The diagram illustrates the lookup table process for activation using SiLU. The SiLU (sigmoid Linear Unit) activation function, also known as the Swish activation function, is a novel activation function proposed in recent years. It is defined as the input multiplied by its sigmoid function value.

[0179] SiLU(x) = x*δ(x);

[0180] in: The table lookup process includes:

[0181] start;

[0182] Find the maximum and minimum values ​​of the input vector, x_min and x_max.

[0183] For x_min and x_max, divide them into table_num parts, generating table_num+1 split points;

[0184] Perform SiLU and quantization operations on table_num+1 inputs x to generate table_num+1 outputs y;

[0185] Output table entries y0 and dy based on table_num+1 outputs; generate offset, scale_n, and scale_p; end.

[0186] 3.3) further includes:

[0187] `offset`, `scale_n`, and `scale_p` are parameters required when using a dual-table architecture: `offset` is the dividing point between negative and positive entries, and `table_neg_num` is the number of negative entries, obtained through optimization methods. The specific steps are as follows... Figure 3 As shown:

[0188] scale_n=std::pow(2,16) / (offset-x_min)*

[0189] table_neg_num / 32(Formula 1)

[0190] scale_p=std::pow(2,16) / (x_min-table_zeropoint_value)*(32-table_neg_num) / 32(Formula 2)

[0191] start;

[0192] The input vector X is generated at any time, and its maximum and minimum values ​​are x_min and x_max, respectively.

[0193] Calculate the output vector Y_ori = SiLU(X);

[0194] TABLE_ITEMS_STEP=32 Initialization Loss=std::pow(2,31)K=0offset=0table_pos_num=TABLE_ITEMS_STEP;

[0195] Determine K <TABLE_ITEMS_STEP,

[0196] If NO, set table_pos_num = TABLE_ITEMS_STEP; End.

[0197] If Yes, then table_pos_num = TABLE_ITEMS_STEP - k and table_neg_num = k, generating the current table entry; look up the table to generate output Y_lut, calculate the difference between Y_lut and Y_ori (loss_current); and determine the loss_current. <loss;

[0198] If Yes, then loss = loss_current, offset = k, table_pos_num = TABLE_ITEMS_STEP - k; K = K + 1, return the judgment K. <TABLE_ITEMS_STEP;

[0199] If NO, K = K + 1, return the value of K. <TABLE_ITEMS_STEP。

[0200] 3.4) Add offset, and based on the result: multiply by scale_p for positive values ​​and by scale_n for negative values; where offset is used to adjust the input lookup table value so that when looking up T0, the value is less than 0, S=1, and when looking up T1, the value is greater than or equal to 0, S=0.

[0201] scale_p, scale_n: are used to perform appropriate fixed-point quantization on the lookup values ​​T0(scale_n) and T1(scale_p) of the two regions respectively, so that the lookup value of interest falls on the appropriate entry; the appropriate fixed-point quantization is related to the bit width of quantization and is an empirical value, which is 10 when it is 8 bits;

[0202] 3.5) Based on 3.2) and 3.3), we get:

[0203] A) Convert absolute value to integer;

[0204] B) Take the lower 16 bits and output them to Lbit. If x_fixed_abs_int is the converted integer...

[0205] The code is: int mask_lbit = int(pow(2,16)-1);

[0206] int32_t lbit=int32_t(x_fixed_abs_int&mask_lbit);

[0207] lbit refers to the lower 16 bits;

[0208] C) If overflow occurs, V = 1 for 16 bits; otherwise, V = 0. Output the overflow to V and check if x_fixed_abs_int > 65535.

[0209] D) Output the sign bit to S;

[0210] S2, based on the results of steps 1), 2), and 3) in step S1, obtain SVLbit, a 16-bit positive integer;

[0211] S3, then based on LU_Table, SAT_T0, IDX0_T, we obtain:

[0212] V = 1: The value can be either SAT_T0 (S = 1) or SAT_T1 (S = 0).

[0213] V=0: Take the high 5 bits of Lbit as idx

[0214] S = 1: If idx ≥ IDX0_T1 takes the value SAT_T0,

[0215] S = 0: idx += IDX0_T1, if idx > 31, take the value SAT_T1.

[0216] Take the lower 11 bits of Lbit as dx.

[0217] Using idx and dx to look up tables and calculate y, the linear interpolation formula is:

[0218] y=(1-dx)*y0+dx*y1=dx*(y1-y0)+y0=y0+dx*dy (Formula 1)

[0219] Set up two tables, with each table containing FP16t floating-point numbers;

[0220] in,

[0221] LU_Table: Lookup table containing 32 entries, each containing 2 FP16t numbers; the lookup table is divided into 2 regions, T0 and T1, and can be configured as 2 tables or as 1 table, with 2 tables being the preferred configuration;

[0222] IDX0_T1: The lookup table is divided into two regions, T0 and T1. The entry for T0 is [0, IDX0_T1), and for T1 it is [IDX0_T1, 32). If IDX0_T1 = 0, then there is no T0 region. Negative lookup values ​​use T0, and positive values ​​use T1. In addition, IDX0_T is an index for the negative table.

[0223] In step S3

[0224] table_dy stores delta_y = y<i+1> -y That is, each subsequent term minus the preceding term.

[0225] table_y0 stores y That is, each preceding term,

[0226] Currently, there are a total of 32 table entries, which can accommodate 2 tables;

[0227] Take the high i bits of the L-bit integer, add idx0 to get idx, where idx0 is the starting index of this table; get: dy = table_dy[idx]; y0 = table_y0[idx]

[0228] Take the lower x bits of Lbit = Lbit - i bits and shift them right by x bits to form a fixed-point number dx, that is, dx is a fraction in [0,1) with x bits as the number of significant bits;

[0229] Calculate: y = y0 + dx * dy.

[0230] The L-bit integer is currently 16 bits, and the high i-bit is currently 5 bits;

[0231] The lower xbit = Lbit - ibit bits of Lbit are currently 11 bits;

[0232] The dx is a decimal in the range [0,1), with a minimum value of 0 and a maximum value of 1-2-xbit.

[0233] S4, the lookup result y(FPin) includes:

[0234] Truth output: Convert to FP16;

[0235] Quantization output: Obtain ofbit, convert it into an integer and saturate it to an unsigned ofbit number;

[0236] S5, then write the address according to the ORAM;

[0237] NNA quantization output or write to the ORAM port.

[0238] The variables are explained below:

[0239] LU_Table: A lookup table containing 32 entries, each containing two FP16t numbers. The lookup table is divided into two regions, T0 and T1, and can be configured as two tables or as a single table. IDX0_T1: The lookup table is divided into two regions, T0 and T1. T0's entries are [0, IDX0_T1), and T1's entries are [IDX0_T1, 32). If IDX0_T1 = 0, there is no T0 region. Negative lookup values ​​use T0, and positive values ​​use T1.

[0240] offset: Used to adjust the input lookup values ​​so that when looking up T0, the values ​​are less than 0 (S=1), and when looking up T1, the values ​​are greater than or equal to 0 (S=0).

[0241] `scale_p`, `scale_n`: Used to perform appropriate fixed-point quantization on the lookup values ​​T0(scale_n) and T1(scale_p) of the two regions, ensuring that the lookup values ​​of interest fall on the appropriate entries. `SAT_T0`, `SAT_T1`: Saturation values ​​for T0 and T1. The lookup output takes the saturation value when the fixed-point quantization overflows (V=1) or `idx` goes out of bounds. The saturation value is in the form of: quantized output takes INT16 (containing 12, 8, 4), and true value output takes FP16.

[0242] Although embodiments of this application have been described, those skilled in the art, upon learning the basic inventive concept, can make other changes and modifications to these embodiments. Therefore, the appended claims are intended to be interpreted as including the preferred embodiments as well as all changes and modifications falling within the scope of the embodiments of this application.

[0243] Finally, it should be noted that in this document, relational terms such as "first" and "second" are used only to distinguish one entity or operation from another, and do not necessarily require or imply any such actual relationship or order between these entities or operations. Furthermore, the terms "comprising," "including," or any other variations thereof are intended to cover non-exclusive inclusion, such that a process, method, article, or terminal device that comprises a list of elements includes not only those elements but also other elements not expressly listed, or elements inherent to such a process, method, article, or terminal device. Without further limitations, an element defined by the phrase "comprising one..." does not exclude the presence of other identical elements in the process, method, article, or terminal device that includes said element.

[0244] The above description is merely a preferred embodiment of the present invention and is not intended to limit the present invention. For those skilled in the art, various modifications and variations can be made to the embodiments of the present invention. Any modifications, equivalent substitutions, improvements, etc., made within the spirit and principles of the present invention should be included within the protection scope of the present invention.

Claims

1. A lookup table method for half-precision floating-point numbers (FP16), characterized in that, The method is applicable to image and speech recognition technology and includes the following steps: S1 performs NNA quantization input, Int16 input, and FP16 input respectively: 1) NNA quantization input, i.e., the input of NNA: The NNA is an on-chip core for improving operator efficiency, including: FRAM, WRAM, and MAC core. FRAM is the on-chip RAM that stores input features inside the NNA, and different chips may have different configurations; WRAM is the on-chip RAM that stores weights inside the NNA; MAC is the multiply-accumulate-add operation; NNA quantization input refers to the input of the NNA. 2) Read ORAM address and length, read ORAM port; Int16 input: Int16 is a 16-bit integer, equivalent to a short occupying 2 bytes -32768 to 32767; The 16-bit output is to Lbit (the lower 16 bits), S=0 indicates a positive number, and V=0 indicates no overflow; 3) Read ORAM address and length, read ORAM port; Int16 input, FP16 input, where FP16 half-precision floating-point number uses 5 bits to represent the exponent and 10 bits to represent the fraction, occupying 2 bytes; further includes: 3.1) Convert to FPIN; 3.2) Look up the value FPIN from the table; 3.3) Obtain offset(FPin), scale_n(FPin), and scale_p(FPin); For the lookup table process to be activated by SiLU, the SiLU activation function, also known as the Swish activation function, is defined as the input multiplied by its sigmoid function value: SiLU(x)=x*δ(x), in: The table lookup process includes: start; Find the maximum and minimum values ​​of the input vector, x_min and x_max. For x_min and x_max, divide them into table_num parts, generating table_num+1 split points; Perform SiLU and quantization operations on table_num+1 inputs x to generate table_num+1 outputs y; Output table entries y0 and dy based on table_num+1 outputs; generate offset, scale_n, and scale_p; end. 3.4) Add offset, and based on the result: multiply by scale_p for positive results, and by scale_n for negative results; where, The offset is used to adjust the input lookup value so that when looking up T0, the value is less than 0, S=1, and when looking up T1, the value is greater than or equal to 0, S=0. scale_p, scale_n: are used to perform appropriate fixed-point quantization on the lookup values ​​T0(scale_n) and T1(scale_p) of the two regions respectively, so that the lookup value of interest falls on the appropriate entry; the appropriate fixed-point quantization is related to the bit width of quantization and is an empirical value, which is 10 when it is 8 bits; 3.5) Based on 3.2) and 3.3), we get: A) Convert absolute value to integer; B) Take the lower 16 bits and output them to Lbit. If x_fixed_abs_int is the converted integer... The code is: int mask_lbit = int(pow(2,16)-1); int32_t lbit=int32_t(x_fixed_abs_int&mask_lbit); lbit refers to the lower 16 bits; C) If overflow occurs, V = 1 for 16 bits; otherwise, V = 0. Output the overflow to V and check if x_fixed_abs_int > 65535. D) Output the sign bit to S; S2, based on the results of steps 1), 2), and 3) in step S1, obtain SVLbit, a 16-bit positive integer; S3, then based on LU_Table, SAT_T0, IDX0_T, we obtain: V = 1: The value can be either SAT_T0 (S = 1) or SAT_T1 (S = 0). V=0: Take the high 5 bits of Lbit as idx S = 1: If idx ≥ IDX0_T1 takes the value SAT_T0, S = 0: idx += IDX0_T1, if idx > 31, take the value SAT_T1. Take the lower 11 bits of Lbit as dx. Using idx and dx to look up tables and calculate y, the linear interpolation formula is: y=(1-dx)*y0+dx*y1=dx*(y1-y0)+y0=y0+dx*dy (Formula 1) Set up two tables, with each table containing FP16t floating-point numbers; in, LU_Table: A lookup table containing 32 entries, each containing 2 FP16t numbers. The lookup table is divided into two regions, T0 and T1, and can be configured as two tables or as one table. IDX0_T1: The lookup table is also divided into two regions, T0 and T1. T0's entries are [0, IDX0_T1), and T1's are [IDX0_T1, 32). If IDX0_T1 = 0, there is no T0 region. Negative lookup values ​​use T0, and positive values ​​use T1. Furthermore, IDX0_T is an index for the negative lookup table. S4, the lookup result y(FPin) includes: Truth output: Convert to FP16; Quantization output: Obtain ofbit, convert it into an integer and saturate it to an unsigned ofbit number; S5, then write the address according to the ORAM; NNA quantization output or write to the ORAM port.

2. The lookup table method for half-precision floating-point number FP16 according to claim 1, characterized in that, Step S1, 3.1), further includes: The following is the code implementation for converting an int32 to an FPIN number: This is C code, indicating that when si32 is less than 0, sign is 1, and otherwise sign is 0. It is represented as: int sign = si32 < 0 ? 1 : 0; uint32_t ui32=sign? -si32:si32; Define a variable fpin of type i32fph_t: i32fph_t fpin; fpin.ui=ui32_2_e6m10(ui32); fpin.f.sign = sign; assert(fpin.f.e+FPIN_E0<=FPIN_E_MAX); fpin.f.e+ = FPIN_E0; return fpin.ui; } The code implementation of function ui32_2_e6m10 is as follows: __device__uint32_t ui32_2_e6m10(uint32_t ui32){ int lead_zero=LEAD_ZERO_UI32(ui32); Assume m is in msb, i.e., min msb: uint32_t mantis = ui32< <lead_zero; All 12 bits of msb are 1: const uint32_t msb12one = 0xfff00000; If msb 12one&mantis is always equal to msb 12one, then the following is executed: This is represented as: if((msb12one&mantis)==msb12one){ Rounding, the condition is true when the first 12 bits of mantis are all 0. (Rounding cause C:) lead_zero--; mantis >>=1; } assert(lead_zero>=0&&lead_zero<=31); m in the lower 12 bits: mantis >> = 20; Take the 12th bit of mantis as rnd, the rounding bit: uint32_t rnd = mantis&1; m in the lower 11 bits: mantis >> = 1; 3. The lookup table method for half-precision floating-point number FP16 according to claim 1, characterized in that, Step S1, 3.3), further includes: offset, scale_n, and scale_p are the parameters required when using a dual-table architecture. The offset is the dividing point between negative and positive entries, and table_neg_num is the number of negative entries, obtained through an optimization method. The specific steps are as follows: scale_n=std::pow(2,16) / (offset-x_min)* table_neg_num / 32(Formula 1) scale_p=std::pow(2,16) / (x_min-table_zeropoint_value)* (32-table_neg_num) / 2(Formula 2) start; The input vector X is generated at any time, and its maximum and minimum values ​​are x_min and x_max, respectively. Calculate the output vector Y_ori = SiLU(X); TABLE_ITEMS_STEP=32 Initialization Loss=std::pow(2,31)K=0offset=0table_pos_num=TABLE_ITEMS_STEP; Determine K <TABLE_ITEMS_STEP, If NO, set table_pos_num = TABLE_ITEMS_STEP; End. If Yes, then table_pos_num = TABLE_ITEMS_STEP - k and table_neg_num = k, generating the current table entry; look up the table to generate output Y_lut, calculate the difference between Y_lut and Y_ori (loss_current); and determine the loss_current. <loss; If Yes, loss=loss_current offset=k table_pos_num = TABLE_ITEMS_STEP - k; K = K + 1, return the result of the judgment of K. <TABLE_ITEMS_STEP; If NO, K = K + 1, return the value of K. <TABLE_ITEMS_STEP。 4. The lookup table method for half-precision floating-point number FP16 according to claim 1, characterized in that, In step S3, the lookup table is divided into two regions, T0 and T1, and preferably two tables are configured.

5. The lookup table method for half-precision floating-point number FP16 according to claim 1, characterized in that, In step S3 table_dy delta_y=y<i+1> -and That is, each subsequent term minus the preceding term. table_y0 stores y That is, each preceding term, Currently, there are a total of 32 table entries, which can accommodate 2 tables; Take the high i bits of the L-bit integer, add idx0 to get idx, where idx0 is the starting index of this table; get: dy = table_dy[idx]; y0 = table_y0[idx] Take the lower x bits of Lbit = Lbit - i bits and shift them right by x bits to form a fixed-point number dx, that is, dx is a fraction in [0,1) with x bits as the number of significant bits; Calculate: y = y0 + dx * dy.

6. The lookup table method for half-precision floating-point number FP16 according to claim 5, characterized in that, The L-bit integer is currently 16 bits, and the high i-bit is currently 5 bits; The lower xbit = Lbit - ibit bits of Lbit are currently 11 bits; The dx is a decimal in the range [0,1), with a minimum value of 0 and a maximum value of 1-2-xbit.

Images