數學函式庫與 Intrinsic 函式 (Mathematical Functions and Intrinsics)
重點總覽
CUDA 提供兩條數學函式路徑：標準路徑（C++ 標準庫 cuda::std::，對映 CUDA Math API，host/device 皆可用、精度有 ULP 保證概念）與 intrinsic 路徑（__sinf 等 device-only 函式，映射較少原生指令、更快但較不精確）。除了標準庫外另有一批 Non-Standard 擴充函式（如 rsqrtf、norm3df、Bessel 函式）。--use_fast_math 可一鍵把一組 Math API 換成 intrinsic。
| 項目 | 重點 |
|---|---|
| 標準庫進入點 | cuda::std::，header <cuda/std/cmath>，host + device |
| Max ULP 定義 | 與「round-to-nearest ties-to-even 正確捨入結果」在 ULP 上差的最大絕對值；測試得出、不保證 |
| 0 ULP 類別 | Basic Ops、Nearest Integer、FP Manipulation、Classification/Comparison、單雙精度基本 intrinsic |
| intrinsic 三特性 | __ 前綴、只在 device code、映射較少指令（快但不準） |
| 捨入後綴 | _rn(最近偶數) _rz(趨零) _ru(向上+∞) _rd(向下−∞) |
| 取整建議 | 用 rint/rintf（單指令），勿用 round/roundf（多指令） |
| 快速數學 | --use_fast_math 把 12 組 Math API（含 cuda::std）換成 intrinsic |
| 型別前綴 | __half/__nv_bfloat16 用 h*；__float128 用 __nv_fp128_* |
CUDA C++ 數學標準函式庫（cuda::std）
CUDA 透過 cuda::std:: 命名空間提供完整的 C++ 標準庫數學函式，包含於 <cuda/std/cmath>，host 與 device 皆可用。每個函式都對映到一組 CUDA Math API（依型別不同），並在 device 上有最大 ULP 誤差界。
可用性其實依類別細分：exponential / power / trigonometric / hyperbolic / nearest-integer 的 Math API 雖 host + device，但僅支援 float 與 double；basic operations、FP manipulation、classification/comparison 則 host + device 但不含 __float128（即 __float128 版為 device-only）。整體而言 __half/__nv_bfloat16 與 __float128 版多屬 device-only，host + device 路徑主要適用 float/double。
函式回傳值與「對應精度下、依 round-to-nearest ties-to-even 取得的正確捨入結果」之間，在 ULP 上差的最大觀測絕對值。誤差界由廣泛但非窮盡的測試得出，因此不保證。
五種浮點型別各有命名慣例：
| 型別 | 命名前綴/範例 | 備註 |
|---|---|---|
__nv_bfloat16 / __half |
hexp、hsin、hsqrt… |
部分 device-only；無原生者以轉 float 再轉回模擬 |
float |
expf、sinf、fabsf… |
|
double |
exp、sin、fabs… |
|
__float128 |
__nv_fp128_exp、__nv_fp128_fma… |
超越函式多為 1 ULP |
代表該函式對 CUDA-extended 型別（__half、__nv_bfloat16）無原生實作；此時以「轉成 float → 計算 → 轉回」的方式模擬（error/gamma 與 FP manipulation 類即如此）。
一律 0 ULP 的四大類
下列類別全部函式 max ULP = 0（含 __float128 對應函式）；1 ULP 主要出現在超越函式（exp/log/三角等），不屬此四類：
| 類別 | 代表函式 | 說明 |
|---|---|---|
| Basic Operations | fabs fmod remainder remquo fma fmax fmin fdim nan |
fma(x,y,z)=x·y+z；fdim=max(x−y,0)；half 用 __habs/__hfma(device-only)/__hmax/__hmin |
| Nearest Integer | ceil floor trunc round rint nearbyint lrint llrint lround llround |
見下方效能提示 |
| FP Manipulation | frexp ldexp modf scalbn scalbln ilogb logb nextafter copysign |
ldexp(x,n)=x·2^n |
| Classification/Comparison | fpclassify isfinite isinf isnan isnormal signbit isgreater isless isunordered… |
half 對映 __hisinf/__hisnan/__hgt/__hge/__hlt/__hle/__hne |
把浮點數取整成整數時建議用 rintf()/rint()，不要用 roundf()/round()：round 在 device code 會映射成多條指令，而 rint 只映射成單一指令。truncf/trunc/ceilf/ceil/floorf/floor 同樣是單指令。
語意差別（易考、易混淆）：round/lround/llround 是 ties away from zero（中間值遠離零）；rint/nearbyint/lrint/llrint 是 ties to even（中間值取偶）。兩者不只差在指令數，捨入規則也不同。
各類函式的精度概覽（代表性 ULP）
float 0–6 ULP、double 1–10 ULP、__float128 超越函式 1 ULP」。Exponential / Logarithmic（float / double / fp128）
| 函式 | float | double | __float128 |
|---|---|---|---|
exp |
expf 0 |
exp 2 |
1 |
exp2 |
exp2f 2 |
exp2 1 |
1 |
expm1 |
expm1f 1 |
expm1 1 |
1 |
log |
logf 0 |
log 1 |
1 |
log2 |
log2f 1 |
log2 1 |
1 |
log10 |
log10f 2 |
log10 1 |
1 |
log1p |
log1pf 1 |
log1p 1 |
1 |
Power（float / double / fp128）
| 函式 | float | double | __float128 |
|---|---|---|---|
pow |
powf 4 |
pow 2 |
1 |
sqrt |
sqrtf 0（--use_fast_math 時 1） |
sqrt 0 |
0 |
cbrt |
cbrtf 1 |
cbrt 1 |
N/A |
hypot |
hypotf 3 |
hypot 2 |
1 |
sqrtf 0 ULP，但加上 --use_fast_math 後變 1 ULP（half hsqrt 0 ULP）。Trigonometric / Hyperbolic（代表性）
sinf/cosf2 ULP；sin/cos2 ULP；fp128 三角函式皆 1 ULP；halfhsin/hcos0 ULP。tan一族 4 ULP 等級；反三角asinf/acosf/atanf約 2 ULP；atan2f3 ULP、atan22 ULP。sinhf3 /sinh2；coshf2 /cosh1；tanhf2 /tanh1（halfhtanh0 ULP）。- 反雙曲
asinhf/acoshf/atanhf約 3–4 ULP（float）、2–3 ULP（double）；fp128 皆 1 ULP。
Error / Gamma（float / double）
| 函式 | float | double |
|---|---|---|
erf |
erff 2 |
erf 2 |
erfc |
erfcf 4 |
erfc 5 |
tgamma |
tgammaf 5 |
tgamma 10 |
lgamma |
6 ULP（x ∉ [−10.001, −2.264]），其餘更大 |
4 ULP（x ∉ [−23.0001, −2.2637]），其餘更大 |
__half/__nv_bfloat16 無原生實作，以轉 float 再轉回模擬。Non-Standard CUDA 數學函式
非 C/C++ 標準庫、以擴充形式提供。單/雙精度逐函式決定 host/device 可用性；對 __half/__nv_bfloat16/__float128 則只在 device code。
| 函式 | 意義 | float | double |
|---|---|---|---|
fdividef(x,y) |
x/y（device-only） |
0 ULP，等同 x/y |
N/A |
exp10f/exp10 |
10^x |
2 | 1 |
rsqrtf/rsqrt |
1/√x |
2 | 1 |
rcbrtf/rcbrt |
1/∛x |
1 | 1 |
rhypotf/rhypot |
1/√(x²+y²)（device-only） |
2 | 1 |
norm3df/norm3d |
√(x²+y²+z²)（device-only） |
3 | 2 |
norm4df/norm4d |
√(x²+y²+z²+t²)（device-only） |
3 | 2 |
rnorm3df/rnorm3d |
上者倒數（device-only） | 2 | 1 |
rnorm4df/rnorm4d |
上者倒數（device-only） | 2 | 1 |
sinpif/sinpi |
sin(πx) |
1 | 2 |
cospif/cospi |
cos(πx) |
1 | 2 |
sincospif/sincospi |
sin(πx),cos(πx) |
1 | 2 |
normcdff/normcdf |
Φ(x) |
5 | 5 |
normcdfinvf/normcdfinv |
Φ⁻¹(x) |
5 | 8 |
erfinvf/erfinv |
erf⁻¹(x) |
2 | 5 |
erfcinvf/erfcinv |
erfc⁻¹(x) |
4 | 6 |
erfcxf/erfcx |
e^{x²}·erfc(x) |
4 | 4 |
cyl_bessel_i0f/i1f（及 double） |
I₀(x)/I₁(x)（device-only） |
6 | 6 |
j0f/j1f/y0f/y1f |
Bessel J/Y |
|x| < 8 時 9 ULP；否則 abs error 2.2·10⁻⁶ | |x| < 8 時 7 ULP；否則 abs error 5·10⁻¹² |
jnf(n,x)/jn(n,x) |
Jₙ(x) |
n=128 abs error 2.2·10⁻⁶ |
n=128 abs error 5·10⁻¹² |
norm(dim,p)/rnorm(dim,p)（device-only）無法提供誤差界：採用快速演算法，round-off 造成精度損失。fdividef ≠ __fdividef
非標準函式 fdividef(x,y) 是 0 ULP、等同 x/y；而下節的 intrinsic __fdividef(x,y) 對 |y| ∈ [2⁻¹²⁶, 2¹²⁶] 是 2 ULP。兩者名字差兩個底線，精度不同。
__half/__nv_bfloat16/__float128 的非標準函式（device-only）：hrcp(x) 1/x 0 ULP、hexp10(x) 10^x 0 ULP、hrsqrt(x) 1/√x 0 ULP、htanh_approx(x) tanh 近似 1 ULP；__nv_fp128_exp10 1 ULP。
Intrinsic 函式與捨入模式
Intrinsic 是對應 CUDA C 標準庫數學函式的更快、更不精確版本：
- 名稱為原函式加
__前綴，如__sinf(x)。 - 只能在 device code 使用。
- 較快是因為映射到較少的原生指令。
--use_fast_math會自動把對應的 Math API 函式翻成 intrinsic。
精度高 ←──────────────────────────────→ 速度快
cuda::std / Math API intrinsic __sinf/__expf...
(較多指令, 低 ULP, host+device) (較少指令, 較高誤差, device-only)
└──── --use_fast_math 一鍵切換（含 cuda::std）────┘
捨入模式後綴
| 後綴 | 捨入模式 |
|---|---|
_rn |
round to nearest even（最近偶數） |
_rz |
round towards zero（趨零） |
_ru |
round up（向 +∞） |
_rd |
round down（向 −∞） |
基本 intrinsic（單/雙精度，全部 0 ULP 且 IEEE 相容）
| 運算 | float | double |
|---|---|---|
x+y |
__fadd_[rn,rz,ru,rd] |
__dadd_[rn,rz,ru,rd] |
x−y |
__fsub_[…] |
__dsub_[…] |
x·y |
__fmul_[…] |
__dmul_[…] |
x·y+z |
__fmaf_[…] |
__fma_[…] |
x/y |
__fdiv_[…] |
__ddiv_[…] |
1/x |
__frcp_[…] |
__drcp_[…] |
√x |
__fsqrt_[…] |
__dsqrt_[…] |
__fadd_/__dadd_/__fmul_/__dmul_ 映射到的加法與乘法，編譯器絕不會合併成 FFMA/DFMA 指令；相對地，由 + 與 * 運算子產生的加乘常被合併成 FFMA/DFMA。需要阻止合併時用這些 intrinsic。
單精度專屬 intrinsic（含誤差）
| 函式 | 意義 | 最大 ULP / 誤差 |
|---|---|---|
__fdividef(x,y) |
x/y |
2（|y| ∈ [2⁻¹²⁶, 2¹²⁶]） |
__frsqrt_rn(x) |
1/√x |
0 ULP |
__expf(x) |
e^x |
2 + ⌊|1.173·x|⌋ |
__exp10f(x) |
10^x |
2 + ⌊|2.97·x|⌋ |
__powf(x,y) |
x^y |
由 exp2f(y * __log2f(x)) 推導 |
__logf(x) |
ln(x) |
x∈[0.5,2] abs error 2⁻²¹·⁴¹；否則 3 ULP |
__log2f(x) |
log2(x) |
x∈[0.5,2] abs error 2⁻²²；否則 2 ULP |
__log10f(x) |
log10(x) |
x∈[0.5,2] abs error 2⁻²⁴；否則 3 ULP |
__sinf(x) |
sin(x) |
x∈[−π,π] abs error 2⁻²¹·⁴¹；範圍外更大 |
__cosf(x) |
cos(x) |
x∈[−π,π] abs error 2⁻²¹·⁴¹；範圍外更大 |
__sincosf(x,sptr,cptr) |
sin,cos |
分量同 __sinf/__cosf |
__tanf(x) |
tan(x) |
由 __sinf(x) * (1/__cosf(x)) 推導 |
__tanhf(x) |
tanh(x) |
最大相對誤差 2⁻¹¹；即使 -ftz=true 也不會把 subnormal 結果 flush 到 0 |
CUDA 另有整數位元 intrinsic，如 __ffs（find first set bit）、__popc（population count，計 set bit 數）、__clz（count leading zeros）、__brev（bit reverse）。它們屬整數 intrinsic，誤差/規格不在本浮點段落，使用時請查對應章節。
--use_fast_math 的影響
nvcc 旗標 --use_fast_math 把一組在 device code 呼叫的 Math API 函式翻成 intrinsic 對應版（CUDA C++ 標準庫函式同受影響）：
| Device Function | → Intrinsic |
|---|---|
x/y, fdividef(x,y) |
__fdividef(x,y) |
sinf / cosf / tanf |
__sinf / __cosf / __tanf |
sincosf |
__sincosf |
logf / log2f / log10f |
__logf / __log2f / __log10f |
expf / exp10f |
__expf / __exp10f |
powf |
__powf |
tanhf |
__tanhf |
與其全域開 --use_fast_math，更建議只在效能增益值得、且可接受降低精度與不同特例處理之處，選擇性地把個別函式換成 intrinsic 版本。
IEEE 754-2019、Muller《On the definition of ulp(x)》、Whitehead & Fit-Florea《Precision & Performance: Floating Point and IEEE 754 Compliance for NVIDIA GPUs》、Goldberg《What every computer scientist should know about floating-point arithmetic》、Monniaux《The pitfalls of verifying floating-point computations》、Dinda & Hetland《Do Developers Understand IEEE Floating Point?》。
考試/測驗重點
| 題型 | 關鍵答案 |
|---|---|
cuda::std 數學函式 header？ |
<cuda/std/cmath>，host 與 device 皆可用 |
| Max ULP error 怎麼定義？ | 與 round-to-nearest ties-to-even 正確捨入結果的 ULP 最大絕對差；測試得出、不保證 |
| 哪些類別全 0 ULP？ | Basic Ops、Nearest Integer、FP Manipulation、Classification/Comparison、單雙精度基本 intrinsic |
| 取整該用哪個函式、為何？ | 用 rint/rintf（單指令）；勿用 round/roundf（多指令）；trunc/ceil/floor 也是單指令 |
round vs rint 差別？ |
指令數：rint 單指令、round 多指令；tie-break：round/lround/llround 為 ties away from zero，rint/nearbyint/lrint/llrint 為 ties to even |
fdividef vs __fdividef？ |
fdividef 0 ULP 等同 x/y；__fdividef 對 |y| ∈ [2⁻¹²⁶, 2¹²⁶] 為 2 ULP |
| 表中「N/A」意思？ | __half/__nv_bfloat16 無原生實作，轉 float 再轉回模擬 |
powf vs pow ULP？ |
powf 4、pow 2、fp128 pow 1 |
sqrtf ULP？ |
0；但加 --use_fast_math 變 1 |
__expf 誤差？ |
2 + ⌊|1.173·x|⌋（__exp10f 為 2 + ⌊|2.97·x|⌋） |
| intrinsic 三特性？ | __ 前綴、只在 device code、映射較少指令（快但不準） |
_rn/_rz/_ru/_rd？ |
最近偶數 / 趨零 / 向上(+∞) / 向下(−∞) |
__fadd/__fmul 與 +/* 差別？ |
__fadd_/__dadd_/__fmul_/__dmul_ 不會被合併成 FFMA/DFMA；+、* 常被合併 |
| 單雙精度基本 intrinsic 精度？ | max ULP = 0 且 IEEE 相容 |
__sinf/__cosf 精度範圍？ |
x∈[−π,π] abs error 2⁻²¹·⁴¹，範圍外更大 |
__tanhf 特點？ |
最大相對誤差 2⁻¹¹；即使 -ftz=true 也不 flush subnormal |
--use_fast_math 做什麼？ |
把一組 device Math API（含 cuda::std）自動換成 intrinsic |
norm/rnorm(dim,p) 為何無誤差界？ |
用快速演算法，round-off 造成精度損失 |
tgamma ULP？ |
tgammaf 5、tgamma 10 |
| half/bf16 的 error/gamma 怎麼算？ | 無原生，轉 float 再轉回模擬 |