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入門必看，如何高效地編寫並行程式