函数功能

根据源操作数和目的操作数Tensor的数据类型进行精度转换。

在介绍不同的类型转换模式之前，先介绍下浮点数的表示方式：

half共16 bit，包括1 bit符号位（S），5 bit指数位（E）和10 bit尾数位（M）。
当E不全为0或不全为1时，表示的结果为：

(-1)^S * 2^{E - 15} * (1 + M)

当E全为0时，表示的结果为：

(-1)^S * 2^-14 * M

当E全为1时，若M全为0，表示的结果为±inf（取决于符号位）；若M不全为0，表示的结果为nan。

上图中S=0，E=15，M = 2^-1 + 2^-2，表示的结果为1.75。
float共32 bit，包括1 bit符号位（S），8 bit指数位（E）和23 bit尾数位（M）。
当E不全为0或不全为1时，表示的结果为：

(-1)^S * 2^{E - 127} * (1 + M)

当E全为0时，表示的结果为：

(-1)^S * 2^-126 * M

当E全为1时，若M全为0，表示的结果为±inf（取决于符号位）；若M不全为0，表示的结果为nan。

上图中S = 0，E = 127，M=2^-1 + 2^-2，最终表示的结果为1.75 。

表1 精度转换规则
src类型	dst类型	类型转换模式介绍
float	float	将src 按照round_mode 取整，仍以float 格式存入dst 中。示例：输入0.5 CAST_RINT 模式输出0.0，CAST_FLOOR 模式输出0.0，CAST_CEIL 模式出1.0，CAST_ROUND 模式输出1.0，CAST_TRUNC 模式输出0.0。
	half	将src 按照round_mode 取到half所能表示的数，以half格式（溢出默认按照饱和处理）存入dst中。示例：输入0.5+2^-12，写成float的表示形式：2^-1 * (1+2^-11)，因此E=-1+127=126，M=2^-11： half的指数位可以表示的出2^-1，有E=-1+15=14，但half只有10 bit尾数位，因此灰色部分要进行舍入。 CAST_RINT 模式舍入得尾数0000000000，E=14，M=0，最终表示的结果为0.5； CAST_FLOOR 模式舍入得尾数0000000000，E=14，M=0，最终表示的结果为0.5； CAST_CEIL 模式舍入得尾数0000000001，E=14，M=2^-10，最终表示的结果为0.5+2^-11； CAST_ROUND 模式舍入得尾数0000000001，E=14，M=2^-10，最终表示的结果为0.5+2^-11； CAST_TRUNC 模式舍入得尾数0000000000，E=14，M=0，最终表示的结果为0.5； CAST_ODD 模式舍入得尾数0000000001，E=14，M=2^-10，最终表示的结果为0.5+2^-11 。
	int64_t	将src按照round_mode取整，以int64_t格式（溢出默认按照饱和处理）存入dst中。示例：输入2²²+0.5 CAST_RINT 模式输出2²²，CAST_FLOOR 模式输出2²²，CAST_CEIL 模式出2²² + 1，CAST_ROUND 模式输出2²²+1，CAST_TRUNC 模式输出2²²。
	int32_t	将src按照round_mode取整，以int32_t格式（溢出默认按照饱和处理）存入dst中。示例：输入2²²+ 0.5 CAST_RINT 模式输出2²²，CAST_FLOOR 模式输出2²² ，CAST_CEIL 模式出2²²+1，CAST_ROUND 模式输出2²²+1，CAST_TRUNC 模式输出2²²。
	int16_t	将src按照round_mode取整，以int16_t格式（溢出默认按照饱和处理）存入dst中。示例：输入2²²+0.5 CAST_RINT 模式输出2¹⁵-1，CAST_FLOOR 模式输出2¹⁵-1，CAST_CEIL 模式出2¹⁵-1，CAST_ROUND 模式输出2¹⁵-1，CAST_TRUNC 模式输出2¹⁵-1（溢出处理）。
half	float	将src以float格式存入dst中，不存在精度转换问题，无舍入模式。示例：输入1.5-2^-10，输出1.5-2^-10
	int32_t	将src按照round_mode取整，以int32_t格式存入dst中。示例：输入-1.5 CAST_RINT 模式输出-2，CAST_FLOOR 模式输出-2，CAST_CEIL 模式出-1，CAST_ROUND 模式输出-2，CAST_TRUNC 模式输出-1。
	int16_t	将src按照round_mode取整，以int16_t格式（溢出默认按照饱和处理）存入dst中。示例：输入2⁷-0.5 CAST_RINT 模式输出2⁷，CAST_FLOOR 模式输出2⁷-1，CAST_CEIL 模式出2⁷，CAST_ROUND 模式输出2⁷，CAST_TRUNC 模式输出2⁷-1。
	int8_t	将src按照round_mode取整，以int8_t格式（溢出默认按照饱和处理）存入dst中。示例：输入2⁷-0.5 CAST_RINT 模式输出2⁷-1，CAST_FLOOR 模式输出2⁷-1，CAST_CEIL 模式出2⁷-1，CAST_ROUND 模式输出2⁷-1，CAST_TRUNC 模式输出2⁷-1（溢出处理）。
	uint8_t	将src按照round_mode取整，以uint8_t格式（溢出默认按照饱和处理）存入dst中。示例：输入1.75 CAST_RINT 模式输出2，CAST_FLOOR 模式输出1，CAST_CEIL 模式出2，CAST_ROUND 模式输出2，CAST_TRUNC 模式输出1。
uint8_t	half	将src以half格式存入dst中，不存在精度转换问题，无舍入模式。；示例：输入1，输出1.0
int8_t	half	将src以half格式存入dst中，不存在精度转换问题，无舍入模式。示例：输入-1，输出-1.0
int16_t	half	将src按照round_mode取到half所能表示的数，以half格式存入dst中。示例：输入2¹²+2，写成half的表示形式：2¹² * (1+2^-11)，要求E=12+15=27，M=2^-11：由于half只有10 bit尾数位，因此灰色部分要进行舍入。 CAST_RINT 模式舍入得尾数0000000000，E=27，M=0，最终表示的结果为2¹²； CAST_FLOOR 模式舍入得尾数0000000000，E=27，M=0，最终表示的结果为2¹²； CAST_CEIL 模式舍入得尾数0000000001，E=27，M=2^-10，最终表示的结果为2¹²+4； CAST_ROUND 模式舍入得尾数0000000001，E=27，M=2^-10，最终表示的结果为2¹²+4； CAST_TRUNC 模式舍入得尾数0000000000，E=27，M=0，最终表示的结果为2¹²。
int16_t	float	将src以float格式存入dst中，不存在精度转换问题，无舍入模式。示例：输入2¹⁵-1，输出2¹⁵-1
int32_t	float	将src按照round_mode取到float所能表示的数，以float格式存入dst中。示例：输入2²⁵+3，写成float的表示形式：2²⁵ * (1+2^-24+2^-25)，要求E=25+127=152， M=2^-24+2^-25：由于float只有23 bit尾数位，因此灰色部分要进行舍入。 CAST_RINT 模式舍入得尾数00000000000000000000001，E=152，M=2^-23，最终表示的结果为2²⁵+4； CAST_FLOOR 模式舍入得尾数00000000000000000000000，E=152，M=0，最终表示的结果为2²⁵； CAST_CEIL 模式舍入得尾数00000000000000000000001，E=152，M=2^-23，最终表示的结果为2²⁵+4； CAST_ROUND 模式舍入得尾数00000000000000000000001，E=152，M=2^-23，最终表示的结果为2²⁵+4； CAST_TRUNC 模式舍入得尾数00000000000000000000000，E=152，M=0，最终表示的结果为2²⁵ 。
	int64_t	将src以int64_t格式存入dst中，不存在精度转换问题，无舍入模式。示例：输入2³¹-1，输出2³¹-1
	int16_t	将src以int16_t格式（溢出默认按照饱和处理）存入dst中，不存在精度转换问题，无舍入模式。示例：输入2³¹-1，输出2¹⁵-1
int64_t	int32_t	将src以int32_t格式（溢出默认按照饱和处理）存入dst中，不存在精度转换问题，无舍入模式。示例：输入2³¹，输出2³¹-1
int64_t	float	将src按照round_mode取到float所能表示的数，以float格式存入dst中。示例：输入2³⁵+2¹²+2¹¹，写成float的表示形式：2³⁵ * (1+2^-23+2^-24)，要求E=35+127=162，M=2^-23+2^-24：由于float只有23 bit尾数位，因此灰色部分要进行舍入。 CAST_RINT 模式舍入得尾数00000000000000000000010，E=162，M=2^-22，最终表示的结果为2³⁵+2¹³； CAST_FLOOR 模式舍入得尾数00000000000000000000001，E=162，M=2^-23，最终表示的结果为2²⁵+2¹²； CAST_CEIL模式舍入得尾数00000000000000000000010，E=162，M=2^-22，最终表示的结果为2²⁵+2¹³； CAST_ROUND模式舍入得尾数00000000000000000000010，E=162，M=2^-22，最终表示的结果为2²⁵+2¹³； CAST_TRUNC模式舍入得尾数00000000000000000000001，E=162，M=2^-23，最终表示的结果为2²⁵+2¹²。

函数原型

表2 0-3级接口原型定义
接口级别	原型定义
0级接口	mask参数使用逐bit模式，该模式的具体介绍请参考表3中的mask参数说明： template <typename T1, typename T2> __aicore__ inline void Cast(const LocalTensor<T1>& dstLocal, const LocalTensor<T2>& srcLocal, const RoundMode& round_mode, uint64_t mask[2], const uint8_t repeatTimes, const UnaryRepeatParams& repeatParams); mask参数使用连续模式，该模式的具体介绍请参考表3中的mask参数说明： template <typename T1, typename T2> __aicore__ inline void Cast(const LocalTensor<T1>& dstLocal, const LocalTensor<T2>& srcLocal, const RoundMode& round_mode, uint64_t mask, const uint8_t repeatTimes, const UnaryRepeatParams& repeatParams);
2级接口	template <typename T1, typename T2> __aicore__ inline void Cast(const LocalTensor<T1>& dstLocal, const LocalTensor<T2>& srcLocal, const RoundMode& round_mode, uint32_t calCount);

参数说明

表3 0级接口参数说明
参数名	输入/输出	描述
dstLocal	输出	目的操作数。类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas 训练系列产品，支持的数据类型见表5 Atlas推理系列产品AI Core，支持的数据类型见表6 Atlas A2训练系列产品，支持的数据类型见表7
srcLocal	输入	源操作数。类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas 训练系列产品，支持的数据类型见表5 Atlas推理系列产品AI Core，支持的数据类型见表6 Atlas A2训练系列产品，支持的数据类型见表7
round_mode	输入	精度转换处理模式，类型是RoundMode。 RoundMode为枚举类型，用以控制精度转换处理模式，具体定义为： enum class RoundMode { CAST_NONE = 0, // 在转换有精度损失时表示CAST_RINT模式，不涉及精度损失时表示不取整 CAST_RINT, // rint，四舍六入五成双取整 CAST_FLOOR, // floor，向负无穷取整 CAST_CEIL, // ceil，向正无穷取整 CAST_ROUND, // round，四舍五入取整 CAST_TRUNC, // trunc，向零取整 CAST_ODD, // Von Neumann rounding，最近邻奇数取整 }; 对于Atlas 训练系列产品，CAST_ROUND表示反向0取整，远离0，对正数 x.y 变成(x+1)，对负数 -x.y，变成 -(x+1)。
mask	输入	mask用于控制每次迭代内参与计算的元素。连续模式：表示前面连续的多少个元素参与计算。数据类型为uint64。取值范围和操作数的数据类型有关，数据类型不同，每次迭代内能够处理的元素个数最大值不同。当源操作数或目的操作数为16位时，mask∈[1, 128]；当源操作数或目的操作数为32位时，mask∈[1, 64]；当源操作数或目的操作数为64位时，mask∈[1, 32]。逐bit模式：可以按位控制哪些元素参与计算，bit位的值为1表示参与计算，0表示不参与。参数类型为长度为2的uint64_t类型数组。例如，mask=[8, 0]，8=0b1000，表示仅第4个元素参与计算。参数取值范围和操作数的数据类型有关，数据类型不同，每次迭代内能够处理的元素个数最大值不同。当源操作数或目的操作数为16位时，mask[0]、mask[1]∈[0, 2⁶⁴-1]并且不同时为0；当源操作数或目的操作数为32位时，mask[1]为0，mask[0]∈(0, 2⁶⁴-1]；当源操作数或目的操作数为64位时，mask[1]为0，mask[0]∈(0, 2³²-1]。当源操作数和目的操作数位数不同时，以数据类型的字节较大的为准。例如，源操作数为half类型，目的操作数为int32_t类型，计算mask时以int32_t为准。
repeatTimes	输入	重复迭代次数。矢量计算单元，每次读取连续的256 Bytes数据进行计算，为完成对输入数据的处理，必须通过多次迭代（repeat）才能完成所有数据的读取与计算。repeatTimes表示迭代的次数。关于该参数的具体描述请参考重复迭代次数-Repeat times。
UnaryRepeatParams	输入	控制操作数地址步长的数据结构。结构体内包含操作数相邻迭代间相同block的地址步长，操作数同一迭代内不同block的地址步长等参数。该数据结构的定义请参考UnaryRepeatParams。相邻迭代间相同block的地址步长参数的详细说明请参考相邻迭代间相同block的地址步长；同一迭代内不同block的地址步长参数请参考同一迭代内不同block的地址步长。

表4 2级接口参数说明
参数名	输入/输出	描述
dstLocal	输出	目的操作数。类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas 训练系列产品，支持的数据类型见表5 Atlas推理系列产品AI Core，支持的数据类型见表6 Atlas A2训练系列产品，支持的数据类型见表7
srcLocal	输入	源操作数。类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas 训练系列产品，支持的数据类型见表5 Atlas推理系列产品AI Core，支持的数据类型见表6 Atlas A2训练系列产品，支持的数据类型见表7
round_mode	输入	精度转换处理模式，类型是RoundMode。 RoundMode为枚举类型，用以控制精度转换处理模式，具体定义为： enum class RoundMode { CAST_NONE = 0, // 在转换有精度损失时表示CAST_RINT模式，不涉及精度损失时表示不取整 CAST_RINT, // rint，四舍六入五成双取整 CAST_FLOOR, // floor，向负无穷取整 CAST_CEIL, // ceil，向正无穷取整（C语言ceil） CAST_ROUND, // round，四舍五入取整 CAST_TRUNC, // trunc，向零取整 CAST_ODD, // Von Neumann rounding，最近邻奇数取整 };
calCount	输入	输入数据元素个数。

表5 Atlas 训练系列产品 Cast指令参数说明
src 数据类型	dst 数据类型	roundMode supported
half	float	CAST_NONE
	int32_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC
	int8_t	CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC, CAST_NONE
	uint8_t	CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC, CAST_NONE
float	half	CAST_NONE, CAST_ODD
float	int32_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC
uint8_t	half	CAST_NONE
int8_t	half	CAST_NONE
int32_t	float	CAST_NONE

表6 Atlas推理系列产品AI Core Cast指令参数说明
src数据类型	dst数据类型	roundMode supported	halfBlock
half	int32_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None
	int16_t	CAST_RINT	None
	float	CAST_NONE	None
	int8_t	CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC, CAST_NONE	None
	uint8_t	CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC, CAST_NONE	None
float	int32_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None
float	half	CAST_NONE, CAST_ODD	None
uint8_t	half	CAST_NONE	None
int8_t	half	CAST_NONE	None
int16_t	half	CAST_NONE	None
int32_t	float	CAST_NONE	None

表7 Atlas A2训练系列产品 Cast指令参数说明
src数据类型	dst数据类型	roundMode supported	halfBlock
half	int32_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None
	int16_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None
	float	CAST_NONE	None
	int8_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC, CAST_NONE	None
	uint8_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC, CAST_NONE	None
float	float	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None
	int32_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None
	half	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC, CAST_ODD, CAST_NONE	None
	int64_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None
	int16_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None
float	bfloat16_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None
bfloat16_t	float	CAST_NONE	None
bfloat16_t	int32_t	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None
uint8_t	half	CAST_NONE	None
int8_t	half	CAST_NONE	None
int16_t	half	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC, CAST_NONE	None
int16_t	float	CAST_NONE	None
int32_t	float	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC, CAST_NONE	None
	int16_t	CAST_NONE	None
	int64_t	CAST_NONE	None
int64_t	int32_t	CAST_NONE	None
int64_t	float	CAST_RINT, CAST_FLOOR, CAST_CEIL, CAST_ROUND, CAST_TRUNC	None

返回值

无

支持的型号

Atlas 训练系列产品

Atlas推理系列产品AI Core

Atlas A2训练系列产品

注意事项

repeatTimes∈[0,255]。
每个repeat的并行度取决于数据精度、AI处理器型号，如f32->f16转换每次迭代操作64个源/目的元素。
指令dstRepStride/srcRepStride∈[0,255]。
dst/src所支持的数据类型与AI处理器型号有关，如果不支持，工具会报错。
dst与src的应为不同Tensor，或同一Tensor的同一元素，不支持同一Tensor的不同元素。
src为float，dst为float时，取整模式表示向整数取整（仍为float类型），其他情况表示向dst 数据类型所能表示的数字取整。
为了节省地址空间，开发者可以定义一个Tensor，供源操作数与目的操作数同时使用（即地址重叠），相关约束如下：
- 对于单次repeat（repeatTimes=1），且源操作数与目的操作数之间要求100%完全重叠，不支持部分重叠。
- 对于多次repeat（repeatTimes>1），若源操作数与目的操作数之间存在依赖，即第N次迭代的目的操作数是第N+1次的源操作数，这种情况是不支持地址重叠的。
二进制下的舍入和十进制类似，具体如下：
- CAST_RINT模式下，若待舍入部分的第一位为0，则不进位；若第一位为1且后续位不全为0，则进位；若第一位为1且后续位全为0，当M的最后一位为0则不进位，当M的最后一位为1则进位。
- CAST_FLOOR模式下，若S为0，则不进位；若S为1，当待舍入部分全为0则不进位，否则，进位。
- CAST_CEIL模式下，若S为1，则不进位；若S为0，当待舍入部分全为0则不进位；否则，进位。
- CAST_ROUND模式下，若待舍入部分的第一位为0，则不进位；否则，进位。
- CAST_TRUNC模式下，总是不进位。
- CAST_ODD模式下，若待舍入部分全为0，则不进位；若待舍入部分不全为0，当M的最后一位为1则不进位，当M的最后一位为0则进位。

调用示例

本样例中只展示Compute流程中的部分代码。本样例的srcLocal为half类型，dstLocal为int32_t类型，计算mask时以int32_t为准。

如果您需要运行样例代码，请将该代码段拷贝并替换样例模板中Compute函数的部分代码即可。

0级接口样例-mask连续模式

uint64_t mask = 256 / sizeof(int32_t);
// repeatTimes = 8, 64 elements one repeat, 512 elements total
// dstBlkStride, srcBlkStride = 1, no gap between blocks in one repeat
// dstRepStride, srcRepStride = 8, no gap between repeats
Cast(dstLocal, srcLocal, RoundMode::CAST_CEIL, mask, 8, { 1, 1, 8, 8 });

0级接口样例-mask逐bit模式

uint64_t mask[2] = { 0, UINT64_MAX };
// repeatTimes = 8, 64 elements one repeat, 512 elements total
// dstBlkStride, srcBlkStride = 1, no gap between blocks in one repeat
// dstRepStride, srcRepStride = 8, no gap between repeats
Cast(dstLocal, srcLocal, RoundMode::CAST_CEIL, mask, 8, { 1, 1, 8, 8 });

2级接口样例

Cast(dstLocal, srcLocal, RoundMode::CAST_CEIL, 512);

结果示例如下：

输入数据(srcLocal): [6.938 -8.86 -0.2263 ... 1.971 1.778]
输出数据(dstLocal): 
[7 -8 0 ... 2 2]

样例模板

为了方便您快速运行指令中的参考样例，本章节提供样例模板。

您可以将以下样例模板作为代码框架，只需将具体指令中的样例片段拷贝替换下文代码段中的加粗内容即可。

#include "kernel_operator.h"
namespace AscendC {
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()
    {
        LocalTensor<half> srcLocal = inQueueSrc.AllocTensor<half>();
        DataCopy(srcLocal, srcGlobal, 512);
        inQueueSrc.EnQue(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
        LocalTensor<int32_t> dstLocal = outQueueDst.AllocTensor<int32_t>();

        Cast(dstLocal, srcLocal, RoundMode::CAST_CEIL, 512);

        outQueueDst.EnQue<int32_t>(dstLocal);
        inQueueSrc.FreeTensor(srcLocal);
    }
    __aicore__ inline void CopyOut()
    {
        LocalTensor<int32_t> dstLocal = outQueueDst.DeQue<int32_t>();
        DataCopy(dstGlobal, dstLocal, 512);
        outQueueDst.FreeTensor(dstLocal);
    }
private:
    TPipe pipe;
    TQue<QuePosition::VECIN, 1> inQueueSrc;
    TQue<QuePosition::VECOUT, 1> outQueueDst;
    GlobalTensor<half> srcGlobal;
    GlobalTensor<int32_t> dstGlobal;
};
} // namespace AscendC
extern "C" __global__ __aicore__ void cast_simple_kernel(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm)
{
    AscendC::KernelCast op;
    op.Init(srcGm, dstGm);
    op.Process();
}

Cast