asm内联语句在编译不支持的情况下的办理办法

asm内联语句在编译不支持的情况下的办理办法

本文链接
在编写 CUDA 的 C++ 程序时，NVIDIA 的库文件中包含这样一行代码，这行代码在 NVIDIA的编译器 nvcc 编译时可以正常通过，但迁移到 AMD 的编译器中将无法识别，天生报错。为了便于做迁移工作，我需要获得这行代码的执行信息。
首先观察一下这行报错代码。

1. asm volatile("cvt.rni.sat.s8.f16 %0, %1;" : "=h"(int16) : "h"(int16_in));

复制代码

asm代码信息

询问GPT可以得出下面的回答：
这行代码是使用ARM架构的NEON指令集进行编程的一个例子。
NEON是一种SIMD（单指令多数据）架构扩展，用于ARM处理器，它允许开发者通过单条指令对多个数据进行并行处理。
这行代码的具体作用是。
asm volatile: 这是GCC的扩展语法，用于在C或C++代码中嵌入汇编语言。volatile关键字告诉编译器不要对这段汇编代码进行任何优化。
"cvt.rni.sat.s8.f16 %0, %1;" 这是NEON指令的一部分。这条指令的具体作用是：
cvt: 表现转换（convert）。
rni: 表现舍入到最接近的整数，假如有两个整数同样接近，则舍入到偶数。
sat: 表现饱和（saturation），即假如结果超出了目的范例的范围，则将其限定在目的范例的最大或最小值。
s8: 表现目的数据范例是8位有符号整数。
f16: 表现源数据范例是16位浮点数。
%0 和 %1: 分别表现第一个和第二个操作数的位置。
"=h"(int16): 这部分指定了输出操作数%0。"=h"表现这是一个16位（half）的输出操作数，并且它将被存储在变量int16中。
"h"(int16_in): 这部分指定了输入操作数%1。"h"表现这是一个16位（half）的输入操作数，它来自变量int16_in。
综上所述，这行代码的作用是将16位浮点数int16_in转换为8位有符号整数，结果存储在int16中。转换过程中会进行舍入和饱和处理。
提取其中的关键信息，可以发现。

1. > rni: 表示舍入到最接近的整数，如果有两个整数同样接近，则舍入到偶数。
2. > sat: 表示饱和(saturation)，即如果结果超出了目标类型的范围，则将其限制在目标类型的最大或最小值。

复制代码

此时根据这些信息，可以联想到这是一行作饱和舍入的代码。对于饱和舍入(Saturation Rounding)，其具体定义:

• 饱和舍入（Saturation Rounding）是一种数值处理方法，常用于数字信号处理和图像处理范畴。在饱和舍入中，当一个数值需要被转换或舍入到某个特定的数值范围时，假如这个数值超出了目的范围，它不会被简单地截断或进行标准的四舍五入，而是被“饱和”到目的范围的最小值或最大值。
  
• 例如，假设我们有一个8位的有符号整数范围，这个范围是从-128到127。假如一个数值在转换过程中计算出的结果是130，那么按照饱和舍入的规则，这个数值会被饱和到127，由于这是这个范围内的最大值。同样，假如一个数值计算出的结果是-130，它会被饱和到-128，由于这是这个范围内的最小值。
  
• 饱和舍入的好处是它避免了数据溢出的问题，保持了数据的完整性，并且在某些应用中，如图像处理，它有助于防止图像质量的下降。
  

有了这些前置信息，我们就可以知道，这行代码做了两件事，就是将传入的数据做了一次舍入操作，再对数据范围做了截取。对于舍入方式，其中也有表明: 舍入到最接近的整数(rni)。
rni 是“round to nearest integer”的缩写，表现舍入到最接近的整数。
这种舍入方式遵照以下规则：
假如小数部分恰好是0.5，那么结果会舍入到最接近的偶数。这被称为“银行家舍入”或“四舍六入五成双”。
假如小数部分小于0.5，那么结果会向下舍入到更小的整数。
假如小数部分大于或等于0.5，那么结果会向上舍入到更大的整数。
例如，使用rni舍入方法：

1. > 1.5   舍入为  2
2. > 2.5   舍入为  2
3. > -1.5  舍入为  -2
4. > -2.5  舍入为  -2

复制代码

在 AMD 支持的内联asm汇编语句和寄存器范例中找不到上述的实现，而且寄存器范例的符号表现也有所差别。于是我接纳最简单的实现方式，将这行内联汇编语句直接替换为 C 语句，实现其功能。由于已知了其功能，编写 C 程序也十分的简单。但是由于这行内联语句是直接调用寄存器，运行速度比用 C 编写的语法快，所以简单用 C 替换仅仅是实现了其正确性，性能有所不及。
但是我在 AMD 上尚未找到有对应的汇编指令完成这行代码的实现，因此现在不得不使用这种方法。
语句替换

在不考虑数据范例转换的情况下，我们先来看舍入的规则。以下的数据是实际在 NVIDIA 编译器上调用asm上述代码所实现的结果。可以看到，当数值超过 127 或者小于 -128 的时候，会将数据截断在 127 和 -128 处。这也是8位有符号整数int8_t所能表现的范围(-128 ~ 127)。

1. > -150.0 舍入为  -128
2. > -128.0 舍入为  -128
3. > -1.0   舍入为  -1
4. > -1.6   舍入为  -2
5. > -1.5   舍入为  -2
6. > -1.4   舍入为  -1
7. > -1.0   舍入为  -1
8. > 0.0    舍入为  0
9. > 0.4    舍入为  0
10. > 0.5    舍入为  0
11. > 0.6    舍入为  1
12. > 1.0    舍入为  1
13. > 126.0  舍入为  126
14. > 127.0  舍入为  127
15. > 128.0  舍入为  127
16. > 200.0  舍入为  127

复制代码

内联语句中规定了输入输出的操作数范例，输入是一个16位(half)的输入操作数，从我的上下文中可以得知，传入时的范例是 half 范例。输出是一个16位的操作数，并以此指定了操作16位数据的寄存器(h)，但是传出的数据范例是int8_t, int8_t是8位数据。
从上面可以得知，我们需要的结果数据储存在 int8_t 范例中就已经充足，内联语句中调用的却是16位的寄存器。因此需要对产生的16位数据进行截取才能获得需要的8位数值。内联语句中的 s8 实在就表现输出的数据范例为8位，只不过借用了16位的寄存器而已。
直接对 half 和 int8_t 范例之间做转换会产生错误，由于它们不仅数据存储长度不同，表现数值的方式也是不一样的。为了保险起见，可以用 float 和 int 范例的局部变量储存住数值，作为中心变量，将 half 范例的浮点数转换为期望得到的整数数值。
至于将数值截取到 -128 ~ 127 之间，可以直接将超过范围的数值置为端点值。

1. __device__ int8_t cvt_f16_to_s8(half val)
2. {
3.     float float32 = (float)val;
4.     int int32 = 0;
5.     if (float32 > 0)
6.     {
7.         if (float32 > 127)
8.             int32 = 127;
9.         else
10.             int32 = (int)(float32 + 0.5);   // 强制数据类型转换
11.     }
12.     else if (float32 < 0)
13.     {
14.         if (float32 < -128)
15.             int32 = -128;
16.         else
17.             int32 = (int)(float32 - 0.5);  // 强制数据类型转换
18.     }
19.     return *((int8_t *)&int32);
20. }

复制代码

我的舍入操作中，强制数据范例转换发生在 float 和 int 范例之间，这样可以保证数值截取时得到预期数值大小。返回值通过得到的 int 范例数值地址，转换为 int8_t * 的指针，并取这个 int8_t 的值返回，这样可以保证返回值是 int8_t 范例。至此完成了上述内联汇编语句的全部功能。将这个 cvt_f16_to_s8(half val); 函数替换掉 asm volatile("cvt.rni.sat.s8.f16 %0, %1;" : "=h"(int16) : "h"(int16_in)); 即可。

1. __device__ inline int8_t cuda_cast<int8_t, half>(half val)
2. {
3.     union
4.     {
5.         int8_t int8[2];
6.         int16_t int16;
7.     };
8.     union
9.     {
10.         half fp16;
11.         int16_t int16_in;
12.     };
13.     fp16 = val;
15.     //asm volatile("cvt.rni.sat.s8.f16 %0, %1;" : "=h"(int16) : "h"(int16_in));
16.     int8_t res = cvt_f16_to_s8(val);  // 通过 C 的语法，用函数实现
17.    
18.     return res;
19. }

复制代码

__device__ 是运行在 GPU 上的函数 kernel 声明方式，在这里不用在意。 假如对GPU编程感兴趣，可以移步我CUDA入门的教程文档。
CUDA入门必看，如何高效地编写并行程序

免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！更多信息从访问主页：qidao123.com:ToB企服之家，中国第一个企服评测及商务社交产业平台。