首页 > 其他分享 >asm内联语句在编译不支持的情况下的解决办法

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

时间:2024-09-18 17:14:05浏览次数:1  
标签:语句 舍入 int16 数值 127 128 内联 int8 asm

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

本文链接

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

首先观察一下这行报错代码。

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中。转换过程中会进行舍入和饱和处理。

提取其中的关键信息,可以发现。

> rni: 表示舍入到最接近的整数,如果有两个整数同样接近,则舍入到偶数。
> 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.5   舍入为  2
> 2.5   舍入为  2
> -1.5  舍入为  -2
> -2.5  舍入为  -2

在 AMD 支持的内联asm汇编语句和寄存器类型中找不到上述的实现,而且寄存器类型的符号表示也有所差别。于是我采取最简单的实现方式,将这行内联汇编语句直接替换为 C 语句,实现其功能。由于已知了其功能,编写 C 程序也十分的简单。但是由于这行内联语句是直接调用寄存器,运行速度比用 C 编写的语法快,所以简单用 C 替换仅仅是实现了其正确性,性能有所不及。

但是我在 AMD 上尚未找到有对应的汇编指令完成这行代码的实现,因此目前不得不使用这种方法。

语句替换

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

> -150.0 舍入为  -128
> -128.0 舍入为  -128
> -1.0   舍入为  -1
> -1.6   舍入为  -2
> -1.5   舍入为  -2
> -1.4   舍入为  -1
> -1.0   舍入为  -1
> 0.0    舍入为  0
> 0.4    舍入为  0
> 0.5    舍入为  0
> 0.6    舍入为  1
> 1.0    舍入为  1
> 126.0  舍入为  126
> 127.0  舍入为  127
> 128.0  舍入为  127
> 200.0  舍入为  127

内联语句中规定了输入输出的操作数类型,输入是一个16位(half)的输入操作数,从我的上下文中可以得知,传入时的类型是 half 类型。输出是一个16位的操作数,并以此指定了操作16位数据的寄存器(h),但是传出的数据类型是int8_t, int8_t是8位数据。

从上面可以得知,我们需要的结果数据储存在 int8_t 类型中就已经足够,内联语句中调用的却是16位的寄存器。因此需要对产生的16位数据进行截取才能获得需要的8位数值。内联语句中的 s8 其实就表示输出的数据类型为8位,只不过借用了16位的寄存器而已。

直接对 halfint8_t 类型之间做转换会产生错误,因为它们不仅数据存储长度不同,表示数值的方式也是不一样的。为了保险起见,可以用 floatint 类型的局部变量储存住数值,作为中间变量,将 half 类型的浮点数转换为期望得到的整数数值。

至于将数值截取到 -128 ~ 127 之间,可以直接将超过范围的数值置为端点值。

__device__ int8_t cvt_f16_to_s8(half val)
{
    float float32 = (float)val;
    int int32 = 0;
    if (float32 > 0)
    {
        if (float32 > 127)
            int32 = 127;
        else
            int32 = (int)(float32 + 0.5);   // 强制数据类型转换
    }
    else if (float32 < 0)
    {
        if (float32 < -128)
            int32 = -128;
        else
            int32 = (int)(float32 - 0.5);  // 强制数据类型转换
    }
    return *((int8_t *)&int32);
}

我的舍入操作中,强制数据类型转换发生在 floatint 类型之间,这样可以保证数值截取时得到预期数值大小。返回值通过得到的 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)); 即可。

__device__ inline int8_t cuda_cast<int8_t, half>(half val)
{
    union
    {
        int8_t int8[2];
        int16_t int16;
    };
    union
    {
        half fp16;
        int16_t int16_in;
    };
    fp16 = val;

    //asm volatile("cvt.rni.sat.s8.f16 %0, %1;" : "=h"(int16) : "h"(int16_in));
    int8_t res = cvt_f16_to_s8(val);  // 通过 C 的语法,用函数实现
    
    return res;
}

__device__ 是运行在 GPU 上的函数 kernel 声明方式,在这里不用在意。 如果对GPU编程感兴趣,可以移步我CUDA入门的教程文档。

CUDA入门必看,如何高效地编写并行程序

标签:语句,舍入,int16,数值,127,128,内联,int8,asm
From: https://www.cnblogs.com/beiwei31/p/18418909

相关文章

  • 函数新特性、内联函数、const详解
    函数新特性、内联函数、const详解函数新特性函数回顾与后置返回类型//1.前面放auto,标识函数返回类型放到参数列表之后//2.而放在参数列表之后的返回类型是通过->开始的autofunc(inta,intb)->void内联函数函数定义前加incline,这个普通函数就变成了内联函数......
  • python 提取出sql语句中where的值
    使用正则表达式来提取SQL语句中的WHERE条件的值。假设你的SQL语句格式比较标准,你可以使用以下Python代码来提取WHERE子句中的值。importredefextract_where_clause(sql_query):#使用正则表达式提取WHERE子句中的条件where_clause_pattern=re.compile(r'......
  • Python中的“with”语句:解锁优雅资源管理的秘密
    引言在日常开发过程中,我们经常会遇到需要打开文件、建立网络连接或数据库会话的情况。传统的做法是手动打开资源,在使用完毕后记得关闭。但这种方法存在明显不足:一旦在使用过程中发生异常,很容易忘记关闭资源,导致内存泄漏等问题。而“with”语句提供了一种更安全、简洁的方式来管理......
  • 【检索稳定,JPCS出版】第二届应用统计、建模与先进算法国际学术会议(ASMA2024,9月27日-29
    大会简介由哈尔滨理工大学主办的第二届应用统计、建模与先进算法国际学术会议(ASMA2024)将于2024年9月27日-29日于中国哈尔滨召开。会议将围绕应用统计、建模及先进算法等在数学领域中的最新研究成果,为来自国内外高等院校、科学研究所、企事业单位的专家、教授、学者、工......
  • Python中的`raise`语句:如何优雅地控制错误
    引言在软件开发过程中,错误处理是不可避免的话题。良好的错误处理不仅能提高程序的稳定性,还能显著提升用户体验。Python通过异常机制提供了一种优雅的方式来管理这些意外情况。其中,raise语句作为触发异常的关键角色,其重要性不言而喻。本文将从基础语法入手,逐步深入探讨raise语句的......
  • 【GESP】C++一级练习BCQM3005,基本输出语句printf
    一道基础练习题,练习基本输出语句printf。BCQM3005题目要求描述输出表达式1234∗5678的结果。输入无输出1234∗5678=7006652输入样例无输出样例1234*5678=7006652全文详见个人独立博客:https://www.coderli.com/gesp-1-bcqm3005/【GESP】C++一级练习B......
  • Verilog - ASCII码与16进制相互转换(Task语句,多个ASCII码转换)
    编程思想:1.使用case语句,将Ascii码与Hex对应关系连接;2.使用Task语句将Ascii码转Hex作为一个任务3.调用Task语句,将8bit Ascii码转换为4bitHex数据4.将n个8bitASCII转为n个4bitHex数据进行数据拼接,输出n*4bitHEX数据moduleascii_to_hex(input......
  • 鹏哥C语言39---分支/循环语句练习:猜数字游戏
    #define_CRT_SECURE_NO_WARNINGS#include<stdio.h>#include<stdlib.h>#include<time.h>//voidfun(inta[]) //因为传过来的是地址,所以应该用一个指针变量来接收,故这里的a本质上是个指针变量//{//   printf("%zu",sizeof(a));//输出8 在x64下,指针大小是......
  • SQL语句逻辑执行过程
    1.1SQL语句的逻辑处理顺序SQL语句的逻辑处理顺序,指的是SQL语句按照一定的规则,一整条语句应该如何执行,每一个关键字、子句部分在什么时刻执行。除了逻辑顺序,还有物理执行顺序。物理顺序是SQL语句真正被执行时的顺序(执行计划),它是由各数据库系统的关系引擎中的语句分析器、优化......
  • 循环语句与条件语句的细节与思想 --进阶C语言
    目录if-else组合if的执行顺序操作符的执行顺序测试方法C语言的布尔类型switchcase组合(补充)屏蔽警告的方法在case中执行多条语句,建议case后都带上花括号.多个case执行同样语句do、while、for循环的基本结构continue跳转的位置循环设计的思想推荐推荐使用for的前闭后开写法null......