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位的暫存器而已。
直接對 half
和 int8_t
型別之間做轉換會產生錯誤,因為它們不僅資料儲存長度不同,表示數值的方式也是不一樣的。為了保險起見,可以用 float
和 int
型別的區域性變數儲存住數值,作為中間變數,將 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);
}
我的舍入操作中,強制資料型別轉換髮生在 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));
即可。
__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入門必看,如何高效地編寫並行程式