[h2]Function Description
Converts precision based on the data types of the source and destination operand tensors.
Before understanding the precision conversion rules, you need to know the basics about the representation modes of floating point numbers and the rounding rules of binary numbers.
Representation modes of floating-point numbers:
The half type has 16 bits, including 1 sign bit (S), 5 exponent bits (E), and 10 mantissa bits (M).
When the exponent bits (E) are not all 0s or all 1s, the value is:
(-1)S * 2E - 15 * (1 + M)
When the exponent bits (E) are all 0s, the value is:
(-1)S * 2-14 * M
When the exponent bits (E) are all 1s and the mantissa bits (M) are all 0s, the value is ±inf (depending on the sign bit). When the exponent bits (E) are all 1s and the mantissa bits (M) are not all 0s, the value is Not-a-Number (NaN).
In the preceding figure, S = 0, E = 15, and M = 2-1 + 2-2, indicating that the value is 1.75.
The float type has 32 bits, including 1 sign bit (S), 8 exponent bits (E), and 23 mantissa bits (M).
When the exponent bits (E) are not all 0s or all 1s, the value is:
(-1)S * 2E - 127 * (1 + M)
When the exponent bits (E) are all 0s, the value is:
(-1)S * 2-126 * M
When the exponent bits (E) are all 1s and the mantissa bits (M) are all 0s, the value is ±inf (depending on the sign bit). When the exponent bits (E) are all 1s and the mantissa bits (M) are not all 0s, the value is Not-a-Number (NaN).
In the preceding figure, S = 0, E = 127, and M = 2-1 + 2-2, indicating that the value is 1.75.
The bfloat16_t type has 16 bits, including 1 sign bit (S), 8 exponent bits (E), and 7 mantissa bits (M).
When the exponent bits (E) are not all 0s or all 1s, the value is:
(-1)S * 2E - 127 * (1 + M)
When the exponent bits (E) are all 0s, the value is:
(-1)S * 2-126 * M
When the exponent bits (E) are all 1s and the mantissa bits (M) are all 0s, the value is ±inf (depending on the sign bit). When the exponent bits (E) are all 1s and the mantissa bits (M) are not all 0s, the value is Not-a-Number (NaN).
In the preceding figure, S = 0, E = 127, and M = 2-1 + 2-2, indicating that the value is 1.75.
Rounding in binary mode is similar to that in decimal mode.
In CAST_RINT mode, if the first bit to be rounded is 0, no carry is performed. If the first bit to be rounded is 1 and the subsequent bits are not all 0s, carry is performed. If the first bit is 1 and all subsequent bits are 0, no carry is performed when the last M bit is 0, and carry is performed when the last M bit is 1.
In CAST_FLOOR mode, if bit S is 0, no carry is performed. If bit S is 1 and the bits to be rounded are all 0s, carry is not performed; in other cases, carry is performed.
In CAST_CEIL mode, if bit S is 1, no carry is performed. If bit S is 0 and the bits to be rounded are all 0s, carry is not performed; in other cases, carry is performed.
In CAST_ROUND mode, if the first bit to be rounded is 0, no carry is performed; in other cases, carry is performed.
In CAST_TRUNC mode, no carry is performed.
In CAST_ODD mode, if the bits to be rounded are all 0s, no carry is performed. If the bits to be rounded are not all 0s, no carry is performed when the last M bit is 1, and carry is performed when the last M bit is 0.
The following table describes the precision conversion rules.
Table 1 Precision conversion rules
| src | dst | Precision Conversion Rule |
|---|---|---|
| float | float | Rounds src based on round_mode (precision conversion mode. For details, see the** round_mode** parameter in Parameters) and stores the result in dst in float format. For example, in the case of input 0.5: The output is 0.0 in CAST_RINT mode, 0.0 in CAST_FLOOR mode, 1.0 in CAST_CEIL mode, 1.0 in CAST_ROUND mode, and 0.0 in CAST_TRUNC mode. |
| float | half | Rounds src according to round_mode and writes the result in half type to dst (the overflow part is saturated). Example, for input 0.5 + 2-12, it is represented as 2-1 x (1 + 2-11) in float format. Therefore, E = –1 + 127 = 126, and M = 2-11.
The exponent bits of the half type can represent 2-1, meaning E = –1 + 15 = 14. However, the half type has only 10 mantissa bits. Therefore, the gray part needs to be rounded. In CAST_RINT mode, the result mantissa is 0000000000, E = 14, and M = 0, indicating that the value is 0.5. In CAST_FLOOR mode, the result mantissa is 0000000000, E = 14, and M = 0, indicating that the value is 0.5. In CAST_CEIL mode, the result mantissa is 0000000001, E = 14, and M = 2-10, indicating that the value is 0.5 + 2-11. In CAST_ROUND mode, the result mantissa is 0000000001, E = 14, and M = 2-10, indicating that the value is 0.5 + 2-11. In CAST_TRUNC mode, the result mantissa is 0000000000, E = 14, and M = 0, indicating that the value is 0.5. In CAST_ODD mode, the result mantissa is 0000000001, E = 14, and M = 2-10, indicating that the value is 0.5 + 2-11. |
| float | int64_t | Rounds src according to round_mode and writes the result in int64_t format to dst (the overflow part is saturated). Example, for input 222 + 0.5: The output is 222 in CAST_RINT mode, 222 in CAST_FLOOR mode, 222 + 1 in CAST_CEIL mode, 222 + 1 in CAST_ROUND mode, and 222 in CAST_TRUNC mode. |
| float | int32_t | Rounds src according to round_mode and writes the result in int32_t format to dst (the overflow part is saturated). Example, for input 222 + 0.5: The output is 222 in CAST_RINT mode, 222 in CAST_FLOOR mode, 222 + 1 in CAST_CEIL mode, 222 + 1 in CAST_ROUND mode, and 222 in CAST_TRUNC mode. |
| float | int16_t | Rounds src according to round_mode and writes the result in int16_t format to dst (the overflow part is saturated). Example, for input 222 + 0.5: The output is 215 - 1 (overflow processing) in CAST_RINT mode, 215 - 1 (overflow processing) in CAST_FLOOR mode, 215 - 1 (overflow processing) in CAST_CEIL mode, 215 - 1 (overflow processing) in CAST_ROUND mode, and 215 - 1 (overflow processing) in CAST_TRUNC mode. |
| float | bfloat16_t | Rounds src according to round_mode and writes the result in bfloat16_t type to dst (the overflow part is saturated). Example, for input 0.5 + 2-9 + 2-11, it is represented as 2-1 x (1 + 2-8 + 2-10) in float format. Therefore, E = –1 + 127 = 126, and M = 2-8 + 2-10.
The number of exponent bits of bfloat16_t is the same as that of float, that is, E = 126. However, bfloat16_t has only 7 mantissa bits. Therefore, the gray part needs to be rounded. In CAST_RINT mode, the result mantissa is 0000001, E = 126 and M = 2-7, indicating that the value is 0.5 + 2-8. In CAST_FLOOR mode, the result mantissa is 0000000, E = 126, and M = 0, indicating that the value is 0.5. In CAST_CEIL mode, the result mantissa is 0000001, E = 126 and M = 2-7, indicating that the value is 0.5 + 2-8. In CAST_ROUND mode, the result mantissa is 0000001, E = 126 and M = 2-7, indicating that the value is 0.5 + 2-8. In CAST_TRUNC mode, the result mantissa is 0000000, E = 126, and M = 0. The final result is 0.5. |
| half | float | Stores src in dst in float format. There is no precision conversion issue, and no rounding mode is involved. Example, for input 1.5 - 2-10, the output is 1.5 - 2-10. |
| half | int32_t | Rounds src according to round_mode and writes the result in int32_t format to dst. For example, for input –1.5: The output is –2 in CAST_RINT mode, –2 in CAST_FLOOR mode, –1 in CAST_CEIL mode, –2 in CAST_ROUND mode, and –1 in CAST_TRUNC mode. |
| half | int16_t | Rounds src according to round_mode and writes the result in int16_t format to dst (the overflow part is saturated). Example, for input 27 - 0.5: The output is 27 in CAST_RINT mode, 27 - 1 in CAST_FLOOR mode, 27 in CAST_CEIL mode, 27 in CAST_ROUND mode, and 27 - 1 in CAST_TRUNC mode. |
| half | int8_t | Rounds src according to round_mode and writes the result in int8_t format to dst (the overflow part is saturated). Example, for input 27 - 0.5: The output is 27 - 1 (overflow processing) in CAST_RINT mode, 27 - 1 in CAST_FLOOR mode, 27 - 1 (overflow processing) in CAST_CEIL mode, 27 - 1 (overflow processing) in CAST_ROUND mode, and 27 - 1 in CAST_TRUNC mode. |
| half | uint8_t | Rounds src according to round_mode and writes the result in uint8_t format to dst (the overflow part is saturated). For example, in the case of input 1.75: The output is 2 in CAST_RINT mode, 1 in CAST_FLOOR mode, 2 in CAST_CEIL mode, 2 in CAST_ROUND mode, and 1 in CAST_TRUNC mode. |
| half | int4b_t | Rounds src according to round_mode and writes the result in int4b_t format to dst (the overflow part is saturated). Example, for input 1.5: The output is 2 in CAST_RINT mode, 1 in CAST_FLOOR mode, 2 in CAST_CEIL mode, 2 in CAST_ROUND mode, and 1 in CAST_TRUNC mode. |
| bfloat16_t | float | Stores src in dst in float format. There is no precision conversion issue, and no rounding mode is involved. Example, for input 1.5 - 2-6, the output is 1.5 - 2-6. |
| bfloat16_t | int32_t | Rounds src according to round_mode and writes the result in int32_t format to dst (the overflow part is saturated). Example, for input 26 + 0.5: The output is 26 in CAST_RINT mode, 26 in CAST_FLOOR mode, 26 + 1 in CAST_CEIL mode, 26 + 1 in CAST_ROUND mode, and 26 in CAST_TRUNC mode. |
| int4b_t | half | Stores src in dst in half format. There is no precision conversion issue, and no rounding mode is involved. Example: input 1 and output 1.0 |
| uint8_t | half | Stores src in dst in half format. There is no precision conversion issue, and no rounding mode is involved. Example: input 1 and output 1.0 |
| int8_t | half | Stores src in dst in half format. There is no precision conversion issue, and no rounding mode is involved. Example: input –1 and output –1.0 |
| int16_t | half | Rounds src according to round_mode and writes the result in half format to dst. Example, for input 212 + 2, it is represented as 212 * (1 + 2-11) in half format. Therefore, E = 12 + 15 = 27, and M = 2-11:
However, half has only 10 mantissa bits. Therefore, the gray part needs to be rounded. In CAST_RINT mode, the result mantissa is 0000000000, E = 27, M = 0, indicating that the value is 212. In CAST_FLOOR mode, the result mantissa is 0000000000, E = 27, M = 0, indicating that the value is 212. In CAST_CEIL mode, the result mantissa is 0000000001, E = 27, M = 2-10, indicating that the value is 212 + 4. In CAST_ROUND mode, the result mantissa is 0000000001, E = 27, M = 2-10, indicating that the value is 212 + 4. In CAST_TRUNC mode, the result mantissa is 0000000000, E = 27, M = 0, indicating that the value is 212. |
| int16_t | float | Stores src in dst in float format. There is no precision conversion issue, and no rounding mode is involved. Example, for input 215 - 1, the output is 215 - 1. |
| int32_t | float | Rounds src according to round_mode and writes the result in float format to dst. Example, for input 225 + 3, it is represented as 225 x (1 + 2-24 + 2-25) in floating format, with E = 25 + 127 = 152 and M = 2-24 + 2-25.
However, float has only 23 mantissa bits. Therefore, the gray part needs to be rounded. In CAST_RINT mode, the result mantissa is 00000000000000000000001, E = 152, M = 2-23, indicating that the value is 225 + 4. In CAST_FLOOR mode, the result mantissa is 00000000000000000000000, E = 152, M = 0, indicating that the value is 225. In CAST_CEIL mode, the result mantissa is 00000000000000000000001, E = 152, M = 2-23, indicating that the value is 225 + 4. In CAST_ROUND mode, the result mantissa is 00000000000000000000001, E = 152, M = 2-23, indicating that the value is 225 + 4. In CAST_TRUNC mode, the result mantissa is 00000000000000000000000, E = 152, M = 0, indicating that the value is 225. |
| int32_t | int64_t | Writes src to dst in int64_t format – precision conversion is not involved. Example, for input 231 - 1, the output is 231 - 1. |
| int32_t | int16_t | Writes src to dst in int16_t (with overflow handled by saturation by default). Precision conversion is not involved. Example, for input 231 - 1, the output is 215 - 1. |
| int64_t | int32_t | Writes src to dst in int32_t format (the overflow part is saturated by default) – precision conversion is not involved. Example, for input 231, the output is 231 - 1. |
| int64_t | float | Rounds src according to round_mode and writes the result in float format to dst. Example, for input 235 + 212 + 211, it is expressed as 235 * (1 + 2-23 + 2-24) in float format, with E = 35 + 127 = 162 and M = 2-23 + 2-24.
However, float has only 23 mantissa bits. Therefore, the gray part needs to be rounded. In CAST_RINT mode, the result mantissa is 00000000000000000000010, E = 162, M = 2-22, indicating that the value is 235 + 213. In CAST_FLOOR mode, the result mantissa is 00000000000000000000001, E = 162, M = 2-23, indicating that the value is 225 + 212. In CAST_CEIL mode, the result mantissa is 00000000000000000000010, E = 162, M = 2-22, indicating that the value is 225 + 213. In CAST_ROUND mode, the result mantissa is 00000000000000000000010, E = 162, M = 2-22, indicating that the value is 225 + 213. In CAST_TRUNC mode, the result mantissa is 00000000000000000000001, E = 162, M = 2-23, indicating that the value is 225 + 212. |
[h2]Function Prototype
Computation of the first n data elements of a tensor:
template <typename T1, typename T2> __aicore__ inline void Cast(const LocalTensor<T1>& dstLocal, const LocalTensor<T2>& srcLocal, const RoundMode& round_mode, const uint32_t calCount)
[h2]Parameters
Table 2 Template parameters
| Parameter | Description |
|---|---|
| T1 | Data type of the destination operand. - For details about the data types supported by Kirin 9020, see Table 4. - For details about the data types supported by Kirin 9030, see Table 4. - For details about the data types supported by Kirin X90, see Table 4. |
| T2 | Data type of the source operands. - For details about the data types supported by Kirin 9020, see Table 4. - For details about the data types supported by Kirin 9030, see Table 4. - For details about the data types supported by Kirin X90, see Table 4. |
Table 3 Parameters
| Parameter | Input/Output | Description |
|---|---|---|
| dstLocal | Output | Destination operand. The type is LocalTensor, and the supported TPosition is VECIN, VECCALC, or VECOUT. The start address of the LocalTensor must be 32-byte aligned. |
| srcLocal | Input | Source operand. The type is LocalTensor, and the supported TPosition is VECIN, VECCALC, or VECOUT. The start address of the LocalTensor must be 32-byte aligned. |
| round_mode | Input | Precision conversion mode. The type is RoundMode. RoundMode is an enumeration class used to control the precision conversion mode. The value can be CAST_NONE, CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC, or CAST_ODD. CAST_ROUND rounds away from 0. For a positive number, x.y is rounded to (x + 1). For a negative number, –x.y is rounded to –(x + 1). |
| calCount | Input | Number of elements of the input data. |
Table 4 Parameters of the Cast instruction for Kirin series products
| src Data Type | dst Data Type | Supported roundMode |
|---|---|---|
| float | half | The overflow part is processed in saturation mode.CAST_RINT/CAST_ODD/CAST_ROUND/CAST_FLOOR/CAST_CEIL/CAST_TRUNC/CAST_NONE |
| float | int16_t | The overflow part is processed in saturation mode. CAST_RINT/CAST_ROUND/CAST_FLOOR/CAST_CEIL/CAST_TRUNC |
| float | int32_t | The overflow part is processed in saturation mode. CAST_RINT/CAST_ROUND/CAST_FLOOR/CAST_CEIL/CAST_TRUNC/CAST_NONE |
| float | float | CAST_RINT/CAST_ROUND/CAST_FLOOR/CAST_CEIL/CAST_TRUNC |
| half | int8_t/uint8_t | The overflow part is processed in saturation mode. CAST_RINT/CAST_ROUND/CAST_FLOOR/CAST_CEIL/CAST_TRUNC/CAST_NONE |
| half | int32_t | CAST_RINT/CAST_ROUND/CAST_FLOOR/CAST_CEIL/CAST_TRUNC |
| half | float | CAST_NONE |
| half | int16_t | CAST_RINT |
| int16_t | half/float | CAST_NONE |
| int32_t | float | CAST_NONE |
| int32_t | int16_t | The overflow part is processed in saturation mode.CAST_NONE |
| uint8_t | half | CAST_NONE |
| int8_t | half | CAST_NONE |
[h2]Supported Models
Kirin 9020 series processors
Kirin 9030 series processors
Kirin X90 series processors
[h2]Precautions
For details about the alignment requirements of the operand address offset, see General Constraints.
dst and src must be different tensors or the same element of the same tensor, rather than different elements of the same tensor.
When both of src and dst are of the float type, the source operand is rounded to an integer of the float type. In other cases, the source operand is rounded to a number that can be represented by the data type of dst.
[h2]Calling Example
This example shows only part of the code used in the computation process (Compute). In these examples, srcLocal is of half type and dstLocal is of int32_t type. mask is computed based on int32_t.
To run the sample code, copy the code snippet to replace parts of code of the Compute function in the sample template.
Example of computing the first n data elements of a tensor:
AscendC::Cast(dstLocal, srcLocal, AscendC::RoundMode::CAST_CEIL, 512);
An example of the result is as follows:
Input data (srcLocal): [29.5 83.6 16.75 45.1 40.62 69.06 47.6 96.5 72.7 57.56 61.25 69.7 29.27 91.2 70.1 14.484 9.625 21.58 9.336 3.125 63.72 9.9 17.28 73.2 75.7 29.81 98.8 99.06 72.94 3.785 24.94 25.56 39.1 58.94 39.6 78.4 5.43 25.48 9.58 60.8 77.56 29.7 70.3 6.312 4.047 87.1 81.6 76.56 59.28 55.66 81.75 73.56 76.9 54.38 7.254 37.84 11.08 77.6 83.6 89.2 93.06 2.96 76.56 62.16 76.25 95.44 86.6 86.75 29.83 82.2 55.03 64.9 56.44 12.89 87.06 39.34 72.25 43.06 63.4 51.72 63.9 0.703 47.84 27.73 99. 89. 97.3 1.277 58.44 14.05 78.9 98.5 28.55 44.8 41.03 40.75 74.2 74.06 10.51 69.2 25.83 35.8 85.5 25.12 82.25 95.3 36.75 55.88 90.9 57.47 7.13 18.1 40.97 31. 99.3 69.4 72.94 62.44 63.7 80. 37.94 11.11 37. 39.72 87.94 31.72 25.7 54.7 32.8 21.64 14.53 55.1 3.607 40.16 77.7 15.15 77.44 43.25 85.75 67.3 30.33 67.56 60.72 58.16 19.84 89.2 18.75 55.56 31.61 9.445 6.5 27.95 48.5 37.16 7.805 37.72 69.6 36.2 92.56 24.72 41.56 48.44 19.27 25.94 25. 8.836 55.75 77.8 25.84 46.16 71.7 63.62 33.28 3.719 55.22 45.97 35.8 27.86 42.22 3.078 92.06 0.805 51.97 76.4 32.03 74.56 28.1 91.2 35.38 0.2009 74.25 87.5 92.75 76.25 51.28 22.9 34.4 28.23 87.5 78.75 63.1 61.56 79.94 6.766 95.1 55. 56.75 39.66 94.75 24.19 29.83 72.6 99.9 12.43 46.56 51.9 92.3 42.66 91.8 95.8 35.2 13.08 60.7 22.22 6.055 2.23 13.875 71.3 99.56 91.94 92. 96.06 97.5 68.75 8.61 1.157 68.2 20.73 63.44 90. 38.78 64.4 88.9 20.75 14.03 97.06 66.8 57.9 86.94 28.5 0.2279 51.8 84.56 39.53 93. 15.66 15.23 71.75 11.44 45.28 57.38 82.5 88.7 9.74 90.4 61.56 68.56 11.22 69.3 40.28 24.78 84.44 23.92 8.4 20.88 48.2 17.42 59.84 93.2 2.191 95.94 93.06 54.53 76.5 37. 41.7 82.7 69.5 92.6 5.8 32.78 84.56 26.5 96.56 0.858 96.44 52.8 90.9 30.52 2.656 32.03 35.72 8.125 21.94 84.5 66.7 96.75 46.8 1.42 58.3 28.75 44.94 66.2 28.67 11.695 41.75 67.25 26.75 17.72 35.9 5.72 55.88 94.7 80.8 71. 86.06 36.78 81.06 56.8 61.34 11.42 74. 32.16 14.695 78.6 56.1 64.4 61.75 50.88 39.6 79.94 71.25 40.7 5.99 67.4 62.28 89.25 12.02 63.12 33.1 59.06 28.2 19.22 59.66 51.6 53.28 97.8 42.25 82. 39.7 50.6 95.06 20.64 26.62 54.9 55. 28.44 26.25 46.56 87.06 98.44 49.34 37.2 97.4 34.3 83.4 57.4 94. 29.31 79.44 19.72 54.9 50.25 58.75 92.5 17.3 17.88 44.7 6.047 50.78 75.3 21.66 71.5 97.75 35.8 93.6 4.367 31.02 66.5 48.25 34. 92.7 36.97 86.5 10.37 82. 29.39 10.63 40.72 72.5 31.56 96.5 70.44 6.074 37.34 7.58 21.72 44.97 77.6 14.22 18.62 47.97 54.6 99.56 81.7 35.75 44.22 28.64 91.56 1.005 44. 8.125 11.7 93.6 70.25 63.94 11.05 50.97 56.47 39.4 35.53 84. 10.21 42.66 62.12 87.7 71.25 87.75 56.03 60.88 31.81 68.1 91.1 67.3 53.6 96.06 43.75 27.86 46.6 87.7 29.47 2.174 88.4 49.53 63.53 84.9 91.75 48.53 91.94 88.44 58.3 88.44 23.11 91.56 71.4 59.66 93.44 28.56 93.3 59.94 90. 18.95 52.8 70.3 58. 21.47 93.7 45.03 84.25 34.06 23.86 38.4 5.566 41.5 35.1 34.8 32.8 81.44 74.75 95.9 23.56 3.562 48.72 92.7 43.88 83.75 69.06 85.8 22.84 63.78 90.94 52.78 ] Output (dstLocal): [ 30 84 17 46 41 70 48 97 73 58 62 70 30 92 71 15 10 22 10 4 64 10 18 74 76 30 99 100 73 4 25 26 40 59 40 79 6 26 10 61 78 30 71 7 5 88 82 77 60 56 82 74 77 55 8 38 12 78 84 90 94 3 77 63 77 96 87 87 30 83 56 65 57 13 88 40 73 44 64 52 64 1 48 28 99 89 98 2 59 15 79 99 29 45 42 41 75 75 11 70 26 36 86 26 83 96 37 56 91 58 8 19 41 31 100 70 73 63 64 80 38 12 37 40 88 32 26 55 33 22 15 56 4 41 78 16 78 44 86 68 31 68 61 59 20 90 19 56 32 10 7 28 49 38 8 38 70 37 93 25 42 49 20 26 25 9 56 78 26 47 72 64 34 4 56 46 36 28 43 4 93 1 52 77 33 75 29 92 36 1 75 88 93 77 52 23 35 29 88 79 64 62 80 7 96 55 57 40 95 25 30 73 100 13 47 52 93 43 92 96 36 14 61 23 7 3 14 72 100 92 92 97 98 69 9 2 69 21 64 90 39 65 89 21 15 98 67 58 87 29 1 52 85 40 93 16 16 72 12 46 58 83 89 10 91 62 69 12 70 41 25 85 24 9 21 49 18 60 94 3 96 94 55 77 37 42 83 70 93 6 33 85 27 97 1 97 53 91 31 3 33 36 9 22 85 67 97 47 2 59 29 45 67 29 12 42 68 27 18 36 6 56 95 81 71 87 37 82 57 62 12 74 33 15 79 57 65 62 51 40 80 72 41 6 68 63 90 13 64 34 60 29 20 60 52 54 98 43 82 40 51 96 21 27 55 55 29 27 47 88 99 50 38 98 35 84 58 94 30 80 20 55 51 59 93 18 18 45 7 51 76 22 72 98 36 94 5 32 67 49 34 93 37 87 11 82 30 11 41 73 32 97 71 7 38 8 22 45 78 15 19 48 55 100 82 36 45 29 92 2 44 9 12 94 71 64 12 51 57 40 36 84 11 43 63 88 72 88 57 61 32 69 92 68 54 97 44 28 47 88 30 3 89 50 64 85 92 49 92 89 59 89 24 92 72 60 94 29 94 60 90 19 53 71 58 22 94 46 85 35 24 39 6 42 36 35 33 82 75 96 24 4 49 93 44 84 70 86 23 64 91 53]
[h2]Sample Template
This section provides a template sample to help you quickly run reference instruction samples.
You can use the following template sample as the code framework and only need to copy the sample snippet in specific instructions to replace the content in bold.
#include "kernel_operator.h"
class KernelCast {
public:
__aicore__ inline KernelCast() {}
__aicore__ inline void Init(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm)
{
srcGlobal.SetGlobalBuffer((__gm__ half*)srcGm);
dstGlobal.SetGlobalBuffer((__gm__ int32_t*)dstGm);
pipe.InitBuffer(inQueueSrc, 1, 512 * sizeof(half));
pipe.InitBuffer(outQueueDst, 1, 512 * sizeof(int32_t));
}
__aicore__ inline void Process()
{
CopyIn();
Compute();
CopyOut();
}
private:
__aicore__ inline void CopyIn()
{
AscendC::LocalTensor<half> srcLocal = inQueueSrc.AllocTensor<half>();
AscendC::DataCopy(srcLocal, srcGlobal, 512);
inQueueSrc.EnQue(srcLocal);
}
__aicore__ inline void Compute()
{
AscendC::LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
AscendC::LocalTensor<int32_t> dstLocal = outQueueDst.AllocTensor<int32_t>();
**AscendC::Cast(dstLocal, srcLocal, AscendC::RoundMode::CAST_CEIL, 512);**
outQueueDst.EnQue<int32_t>(dstLocal);
inQueueSrc.FreeTensor(srcLocal);
}
__aicore__ inline void CopyOut()
{
AscendC::LocalTensor<int32_t> dstLocal = outQueueDst.DeQue<int32_t>();
AscendC::DataCopy(dstGlobal, dstLocal, 512);
outQueueDst.FreeTensor(dstLocal);
}
private:
AscendC::TPipe pipe;
AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueSrc;
AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueDst;
AscendC::GlobalTensor<half> srcGlobal;
AscendC::GlobalTensor<int32_t> dstGlobal;
};
extern "C" __global__ __aicore__ void cast_simple_kernel(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm)
{
KernelCast op;
op.Init(srcGm, dstGm);
op.Process();
}