MUSA Compiler API 完整参考手册

目录

1. 标准 C 库函数
2. 双精度浮点运算（带舍入控制）
3. 单精度浮点运算（带舍入控制）
4. 类型转换函数
5. 数学函数
6. 位操作函数
7. SIMD 向量操作
8. 同步与线程束操作
9. 原子操作
10. 半精度与 BFloat16 操作
11. 内存访问优化
12. 纹理与表面访问
13. 向量类型构造函数
14. 其他内建函数

---

1. 标准 C 库函数

__assert_fail

• 原型: void __assert_fail(const char* expr, const char* file, unsigned int line, const char* func)
• 域: 主机端（glibc 内部）
• 说明: 断言失败时由运行时调用，输出错误信息并终止程序
• 注意: 通常使用 assert() 宏而非直接调用

clock

• 原型: clock_t clock(void)
• 域: 主机端 / 设备端
• 说明:
  • 主机端：返回进程 CPU 时间
  • 设备端：返回 SM 硬件周期计数
• 相关: clock64() - 64位版本（设备端）

内存管理函数

• malloc: void* malloc(size_t size) - 主机堆内存分配
• free: void free(void* ptr) - 释放内存
• memcpy: void* memcpy(void* dst, const void* src, size_t n) - 内存拷贝
• memset: void* memset(void* s, int c, size_t n) - 内存初始化

printf

• 原型: int printf(const char* fmt, ...)
• 域: 主机端 / 设备端
• 说明: 格式化输出。设备端输出在内核结束后刷新
• 注意: 设备端使用影响性能，仅用于调试

---

2. 双精度浮点运算（带舍入控制）

舍入模式后缀：

• rn: round to nearest even (就近取偶)
• rz: round toward zero (向零)
• ru: round toward +∞ (向上)
• rd: round toward -∞ (向下)

加法

• double __dadd_rn(double a, double b)
• double __dadd_rz(double a, double b)
• double __dadd_ru(double a, double b)
• double __dadd_rd(double a, double b)

减法

• double __dsub_rn(double a, double b)
• double __dsub_rz(double a, double b)
• double __dsub_ru(double a, double b)
• double __dsub_rd(double a, double b)

乘法

• double __dmul_rn(double a, double b)
• double __dmul_rz(double a, double b)
• double __dmul_ru(double a, double b)
• double __dmul_rd(double a, double b)

除法

• double __ddiv_rn(double a, double b)
• double __ddiv_rz(double a, double b)
• double __ddiv_ru(double a, double b)
• double __ddiv_rd(double a, double b)

融合乘加 (FMA)

• double __fma_rn(double a, double b, double c) - 计算 a×b+c
• double __fma_rz(double a, double b, double c)
• double __fma_ru(double a, double b, double c)
• double __fma_rd(double a, double b, double c)

平方根

• double __dsqrt_rn(double x) - 默认舍入
• double __dsqrt_ru(double x)
• double __dsqrt_rd(double x)
• double __dsqrt_rz(double x)

---

3. 单精度浮点运算（带舍入控制）

加法

• float __fadd_rn(float a, float b)
• float __fadd_rz(float a, float b)
• float __fadd_ru(float a, float b)
• float __fadd_rd(float a, float b)

减法

• float __fsub_rn(float a, float b)
• float __fsub_rz(float a, float b)
• float __fsub_ru(float a, float b)
• float __fsub_rd(float a, float b)

乘法

• float __fmul_rn(float a, float b)
• float __fmul_rz(float a, float b)
• float __fmul_ru(float a, float b)
• float __fmul_rd(float a, float b)

除法

• float __fdiv_rn(float a, float b)
• float __fdiv_rz(float a, float b)
• float __fdiv_ru(float a, float b)
• float __fdiv_rd(float a, float b)
• float __fdividef(float a, float b) - 快速近似除法

融合乘加

• float __fmaf_rn(float a, float b, float c)
• float __fmaf_rz(float a, float b, float c)
• float __fmaf_ru(float a, float b, float c)
• float __fmaf_rd(float a, float b, float c)

平方根

• float __fsqrt_rn(float x)
• float __fsqrt_rz(float x)
• float __fsqrt_ru(float x)
• float __fsqrt_rd(float x)

倒数与倒数平方根

• float __frcp_rn(float x) - 倒数
• float __frcp_ru(float x)
• float __frsqrt_rn(float x) - 倒数平方根

---

4. 类型转换函数

Double to Float

• float __double2float_rn(double x)
• float __double2float_rz(double x)
• float __double2float_ru(double x)
• float __double2float_rd(double x)

Double to Integer

• int __double2int_rn(double x)
• int __double2int_rz(double x)
• int __double2int_ru(double x)
• int __double2int_rd(double x)

Double to Long Long

• long long __double2ll_rn(double x)
• long long __double2ll_rz(double x)
• long long __double2ll_ru(double x)
• long long __double2ll_rd(double x)

Double to Unsigned

• unsigned int __double2uint_rn(double x)
• unsigned int __double2uint_rz(double x)
• unsigned int __double2uint_ru(double x)
• unsigned int __double2uint_rd(double x)

Float to Integer

• int __float2int_rn(float x)
• int __float2int_rz(float x)
• int __float2int_ru(float x)
• int __float2int_rd(float x)

Integer/Long to Float/Double

• float __int2float_rn(int x)
• double __int2double_rn(int x)
• float __ll2float_rn(long long x)
• double __ll2double_rn(long long x)

位重解释

• int __float_as_int(float x)
• float __int_as_float(int x)
• long long __double_as_longlong(double x)
• double __longlong_as_double(long long x)

Double 分解与组合

• int __double2hiint(double x) - 获取高32位
• int __double2loint(double x) - 获取低32位
• double __hiloint2double(int hi, int lo) - 从高低32位组合

---

5. 数学函数

标准数学函数（支持 float/double）

• 三角函数: sin, cos, tan, asin, acos, atan, atan2
• 双曲函数: sinh, cosh, tanh, asinh, acosh, atanh
• 指数对数: exp, exp2, expm1, log, log2, log10, log1p
• 幂函数: pow, sqrt, cbrt, hypot
• 取整函数: ceil, floor, round, trunc, rint, nearbyint
• 其他: fabs, fmod, remainder, remquo, fdim, fmax, fmin

快速数学函数（设备端，单精度）

• float __sinf(float x) - 快速正弦
• float __cosf(float x) - 快速余弦
• float __tanf(float x) - 快速正切
• float __expf(float x) - 快速指数
• float __logf(float x) - 快速对数
• float __powf(float a, float b) - 快速幂
• void __sincosf(float x, float* s, float* c) - 同时计算sin和cos

特殊函数

• 误差函数: erf, erfc, erfcinv, erfinv
• 伽马函数: lgamma, tgamma
• 贝塞尔函数: cyl_bessel_i0
• 归一化函数: normcdf, normcdfinv
• 向量范数: norm3df, rnorm3df

浮点特殊值检测

• bool isfinite(float/double x)
• bool isinf(float/double x)
• bool isnan(float/double x)
• bool signbit(float/double x)

---

6. 位操作函数

位反转

• unsigned int __brev(unsigned int x) - 32位反转
• unsigned long long __brevll(unsigned long long x) - 64位反转

字节排列

• unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s) - 字节重排

前导零计数

• int __clz(int x) - 32位前导零
• int __clzll(long long x) - 64位前导零

第一个置位位

• int __ffs(int x) - 32位
• int __ffsll(long long x) - 64位

漏斗移位

• unsigned int __funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift)
• unsigned int __funnelshift_r(unsigned int lo, unsigned int hi, unsigned int shift)

半字加法

• int __hadd(int a, int b) - (a+b)>>1，带符号
• unsigned int __uhadd(unsigned int a, unsigned int b) - 无符号
• int __rhadd(int a, int b) - (a+b+1)>>1，带符号
• unsigned int __urhadd(unsigned int a, unsigned int b) - 无符号

乘法高位

• int __mulhi(int a, int b) - 32位乘法高32位
• unsigned int __umulhi(unsigned int a, unsigned int b)
• long long __mul64hi(long long a, long long b) - 64位
• unsigned long long __umul64hi(unsigned long long a, unsigned long long b)

24位乘法

• int __mul24(int a, int b) - 24位有符号乘法
• unsigned int __umul24(unsigned int a, unsigned int b) - 无符号

SAD运算

• unsigned int __sad(int a, int b, unsigned int c) - |a-b|+c
• unsigned int __usad(unsigned int a, unsigned int b, unsigned int c)

饱和运算

• float __saturatef(float x) - 钳制到[0,1]

---

7. SIMD 向量操作

SIMD操作对打包在32位整数中的2个16位或4个8位值进行并行运算。

向量加法

• unsigned int __vadd2(unsigned int a, unsigned int b) - 2x16位
• unsigned int __vadd4(unsigned int a, unsigned int b) - 4x8位
• unsigned int __vaddss2(unsigned int a, unsigned int b) - 饱和有符号
• unsigned int __vaddus2(unsigned int a, unsigned int b) - 饱和无符号

向量减法

• unsigned int __vsub2(unsigned int a, unsigned int b)
• unsigned int __vsub4(unsigned int a, unsigned int b)
• unsigned int __vsubss2(unsigned int a, unsigned int b) - 饱和
• unsigned int __vsubus2(unsigned int a, unsigned int b) - 饱和

向量绝对值

• unsigned int __vabs2(unsigned int a) - 2x16位
• unsigned int __vabs4(unsigned int a) - 4x8位
• unsigned int __vabsss2(unsigned int a) - 饱和

向量绝对差

• unsigned int __vabsdiffs2(unsigned int a, unsigned int b) - 有符号
• unsigned int __vabsdiffu2(unsigned int a, unsigned int b) - 无符号

向量比较

• unsigned int __vcmpeq2(unsigned int a, unsigned int b) - 相等
• unsigned int __vcmpge2(unsigned int a, unsigned int b) - 大于等于
• unsigned int __vcmpgt2(unsigned int a, unsigned int b) - 大于
• unsigned int __vcmple2(unsigned int a, unsigned int b) - 小于等于
• unsigned int __vcmplt2(unsigned int a, unsigned int b) - 小于
• unsigned int __vcmpne2(unsigned int a, unsigned int b) - 不等

向量最值

• unsigned int __vmax2(unsigned int a, unsigned int b) - 最大值
• unsigned int __vmin2(unsigned int a, unsigned int b) - 最小值

向量平均

• unsigned int __vavg2(unsigned int a, unsigned int b) - 平均值
• unsigned int __vhadd2(unsigned int a, unsigned int b) - 半加

向量设置比较结果

• unsigned int __vset*2/4(unsigned int a, unsigned int b) - 设置比较掩码

向量取反

• unsigned int __vneg2(unsigned int a)
• unsigned int __vnegss2(unsigned int a) - 饱和

向量SAD

• unsigned int __vsad2(unsigned int a, unsigned int b) - 绝对差之和

---

8. 同步与线程束操作

基本同步

• void __syncthreads(void) - 块内所有线程同步
• int __syncthreads_and(int predicate) - 同步并返回AND结果
• int __syncthreads_or(int predicate) - 同步并返回OR结果
• int __syncthreads_count(int predicate) - 统计为真的线程数

内存栅栏

• void __threadfence(void) - 设备级内存栅栏
• void __threadfence_block(void) - 块级内存栅栏
• void __threadfence_system(void) - 系统级内存栅栏

Warp级同步

• void __syncwarp(unsigned mask=0xFFFFFFFF) - warp内线程同步
• unsigned __activemask(void) - 获取活动线程掩码
• void __barrier_sync(unsigned mask) - 带掩码的栅栏同步

Warp投票函数

• int __all_sync(unsigned mask, int predicate) - 所有线程为真
• int __any_sync(unsigned mask, int predicate) - 任一线程为真
• unsigned __ballot_sync(unsigned mask, int predicate) - 投票位图
• int __uni_sync(unsigned mask, int predicate) - 所有线程值相同

Warp Shuffle操作

• int __shfl_sync(unsigned mask, int var, int srcLane, int width=32) - 直接交换
• int __shfl_up_sync(unsigned mask, int var, unsigned delta, int width=32) - 向上
• int __shfl_down_sync(unsigned mask, int var, unsigned delta, int width=32) - 向下
• int __shfl_xor_sync(unsigned mask, int var, int laneMask, int width=32) - XOR交换

Warp匹配

• unsigned __match_any_sync(unsigned mask, T value) - 查找相同值的线程
• unsigned __match_all_sync(unsigned mask, T value, int* pred) - 检查所有线程值相同

Warp归约

• unsigned __reduce_add_sync(unsigned mask, unsigned value) - 加法归约
• unsigned __reduce_and_sync(unsigned mask, unsigned value) - AND归约
• unsigned __reduce_or_sync(unsigned mask, unsigned value) - OR归约
• unsigned __reduce_xor_sync(unsigned mask, unsigned value) - XOR归约
• unsigned __reduce_max_sync(unsigned mask, unsigned value) - 最大值归约
• unsigned __reduce_min_sync(unsigned mask, unsigned value) - 最小值归约

其他

• void __trap(void) - 触发陷阱，终止内核
• void __nanosleep(unsigned int ns) - 纳秒级睡眠

---

9. 原子操作

基础原子操作（支持 int/unsigned/float/double）

• T atomicAdd(T* addr, T val) - 原子加
• T atomicSub(T* addr, T val) - 原子减
• T atomicExch(T* addr, T val) - 原子交换
• T atomicMin(T* addr, T val) - 原子最小
• T atomicMax(T* addr, T val) - 原子最大
• T atomicCAS(T* addr, T compare, T val) - 比较并交换

位运算原子操作（整数类型）

• T atomicAnd(T* addr, T val) - 原子AND
• T atomicOr(T* addr, T val) - 原子OR
• T atomicXor(T* addr, T val) - 原子XOR

特殊原子操作

• unsigned atomicInc(unsigned* addr, unsigned val) - 原子递增（模val+1）
• unsigned atomicDec(unsigned* addr, unsigned val) - 原子递减（模val+1）

作用域变体

每个原子操作都有以下作用域变体：

• atomicAdd_block() - 块级作用域
• atomicAdd_system() - 系统级作用域

示例：

int atomicAdd_block(int* addr, int val);
int atomicAdd_system(int* addr, int val);

---

10. 半精度与 BFloat16 操作

Half精度（__half）

算术运算

• __half __hadd(__half a, __half b) - 加法
• __half __hsub(__half a, __half b) - 减法
• __half __hmul(__half a, __half b) - 乘法
• __half __hdiv(__half a, __half b) - 除法
• __half __hfma(__half a, __half b, __half c) - FMA
• __half __hneg(__half a) - 取负

饱和运算

• __half __hadd_sat(__half a, __half b)
• __half __hsub_sat(__half a, __half b)
• __half __hmul_sat(__half a, __half b)

比较运算

• bool __heq(__half a, __half b) - 相等
• bool __hne(__half a, __half b) - 不等
• bool __hlt(__half a, __half b) - 小于
• bool __hle(__half a, __half b) - 小于等于
• bool __hgt(__half a, __half b) - 大于
• bool __hge(__half a, __half b) - 大于等于

特殊函数

• __half hceil(__half x) - 向上取整
• __half hfloor(__half x) - 向下取整
• __half hrint(__half x) - 四舍五入
• __half htrunc(__half x) - 截断
• __half hsin(__half x) - 正弦
• __half hcos(__half x) - 余弦
• __half hsqrt(__half x) - 平方根
• __half hrsqrt(__half x) - 倒数平方根
• __half hexp2(__half x) - 2的幂

类型转换

• float __half2float(__half x)
• __half __float2half(float x)
• int __half2int_rn(__half x)
• __half __int2half_rn(int x)

Half2向量操作

• __half2 __hadd2(__half2 a, __half2 b) - 向量加法
• __half2 __hsub2(__half2 a, __half2 b) - 向量减法
• __half2 __hmul2(__half2 a, __half2 b) - 向量乘法
• __half2 __hfma2(__half2 a, __half2 b, __half2 c) - 向量FMA

BFloat16（__nv_bfloat16）

算术运算

• __nv_bfloat16 __hadd(__nv_bfloat16 a, __nv_bfloat16 b)
• __nv_bfloat16 __hsub(__nv_bfloat16 a, __nv_bfloat16 b)
• __nv_bfloat16 __hmul(__nv_bfloat16 a, __nv_bfloat16 b)
• __nv_bfloat16 __hdiv(__nv_bfloat16 a, __nv_bfloat16 b)
• __nv_bfloat16 __hfma(__nv_bfloat16 a, __nv_bfloat16 b, __nv_bfloat16 c)

比较运算

• bool __heq(__nv_bfloat16 a, __nv_bfloat16 b)
• bool __hne(__nv_bfloat16 a, __nv_bfloat16 b)
• bool __hlt(__nv_bfloat16 a, __nv_bfloat16 b)
• bool __hle(__nv_bfloat16 a, __nv_bfloat16 b)
• bool __hgt(__nv_bfloat16 a, __nv_bfloat16 b)
• bool __hge(__nv_bfloat16 a, __nv_bfloat16 b)

类型转换

• float __bfloat162float(__nv_bfloat16 x)
• __nv_bfloat16 __float2bfloat16(float x)
• __nv_bfloat16 __double2bfloat16(double x)
• int __bfloat162int_rn(__nv_bfloat16 x)
• __nv_bfloat16 __int2bfloat16_rn(int x)

BFloat162向量操作

• __nv_bfloat162 __hadd2(__nv_bfloat162 a, __nv_bfloat162 b)
• __nv_bfloat162 __hmul2(__nv_bfloat162 a, __nv_bfloat162 b)
• __nv_bfloat162 make_bfloat162(__nv_bfloat16 x, __nv_bfloat16 y)

---

11. 内存访问优化

缓存加载指令

这些函数通过不同的缓存层次结构加载数据：

• T __ldg(const T* ptr) - 通过只读数据缓存加载
• T __ldcg(const T* ptr) - 通过L2缓存加载（缓存全局）
• T __ldca(const T* ptr) - 通过L2缓存加载（缓存所有）
• T __ldcs(const T* ptr) - 通过L2缓存加载（流式）
• T __ldlu(const T* ptr) - 通过L2缓存加载（最后使用）
• T __ldcv(const T* ptr) - 通过L2缓存加载（易失性）

支持的类型包括基本类型和向量类型。

缓存存储指令

• void __stcg(T* ptr, T val) - 通过L2缓存存储（缓存全局）
• void __stcs(T* ptr, T val) - 通过L2缓存存储（流式）
• void __stwb(T* ptr, T val) - 通过L2缓存存储（写回）
• void __stwt(T* ptr, T val) - 通过L2缓存存储（直写）

地址空间检测

• unsigned __isGlobal(const void* ptr) - 检查是否为全局内存
• unsigned __isShared(const void* ptr) - 检查是否为共享内存
• unsigned __isConstant(const void* ptr) - 检查是否为常量内存
• unsigned __isLocal(const void* ptr) - 检查是否为本地内存

---

12. 纹理与表面访问

纹理采样

1D纹理

• T tex1D(musaTextureObject_t texObj, float x) - 基础采样
• T tex1DLod(musaTextureObject_t texObj, float x, float level) - 带LOD
• T tex1DGrad(musaTextureObject_t texObj, float x, float dx, float dy) - 带梯度
• T tex1Dfetch(musaTextureObject_t texObj, int x) - 直接获取

2D纹理

• T tex2D(musaTextureObject_t texObj, float x, float y) - 基础采样
• T tex2DLod(musaTextureObject_t texObj, float x, float y, float level) - 带LOD
• T tex2DGrad(musaTextureObject_t texObj, float x, float y, float2 dx, float2 dy) - 带梯度
• T tex2Dgather(musaTextureObject_t texObj, float x, float y, int comp) - Gather操作

3D纹理

• T tex3D(musaTextureObject_t texObj, float x, float y, float z) - 基础采样
• T tex3DLod(musaTextureObject_t texObj, float x, float y, float z, float level) - 带LOD
• T tex3DGrad(musaTextureObject_t texObj, float x, float y, float z, float4 dx, float4 dy) - 带梯度

立方体贴图

• T texCubemap(musaTextureObject_t texObj, float x, float y, float z) - 基础采样
• T texCubemapLod(musaTextureObject_t texObj, float x, float y, float z, float level) - 带LOD

分层纹理

• T tex1DLayered(musaTextureObject_t texObj, float x, int layer) - 1D分层
• T tex2DLayered(musaTextureObject_t texObj, float x, float y, int layer) - 2D分层
• T texCubemapLayered(musaTextureObject_t texObj, float x, float y, float z, int layer) - 立方体分层

表面访问

读取

• T surf1Dread(musaSurfaceObject_t surfObj, int x, musaSurfaceBoundaryMode mode) - 1D
• T surf2Dread(musaSurfaceObject_t surfObj, int x, int y, musaSurfaceBoundaryMode mode) - 2D
• T surf3Dread(musaSurfaceObject_t surfObj, int x, int y, int z, musaSurfaceBoundaryMode mode) - 3D

写入

• void surf1Dwrite(T val, musaSurfaceObject_t surfObj, int x, musaSurfaceBoundaryMode mode) - 1D
• void surf2Dwrite(T val, musaSurfaceObject_t surfObj, int x, int y, musaSurfaceBoundaryMode mode) - 2D
• void surf3Dwrite(T val, musaSurfaceObject_t surfObj, int x, int y, int z, musaSurfaceBoundaryMode mode) - 3D

分层表面

• T surf1DLayeredread(musaSurfaceObject_t surfObj, int x, int layer, musaSurfaceBoundaryMode mode)
• T surf2DLayeredread(musaSurfaceObject_t surfObj, int x, int y, int layer, musaSurfaceBoundaryMode mode)
• T surfCubemapread(musaSurfaceObject_t surfObj, int x, int y, int face, musaSurfaceBoundaryMode mode)
• T surfCubemapLayeredread(musaSurfaceObject_t surfObj, int x, int y, int layerFace, musaSurfaceBoundaryMode mode)

---

13. 向量类型构造函数

基本类型向量构造

每种基本类型都有对应的make函数创建向量：

char向量

• char1 make_char1(signed char x)
• char2 make_char2(signed char x, signed char y)
• char3 make_char3(signed char x, signed char y, signed char z)
• char4 make_char4(signed char x, signed char y, signed char z, signed char w)

int向量

• int1 make_int1(int x)
• int2 make_int2(int x, int y)
• int3 make_int3(int x, int y, int z)
• int4 make_int4(int x, int y, int z, int w)

float向量

• float1 make_float1(float x)
• float2 make_float2(float x, float y)
• float3 make_float3(float x, float y, float z)
• float4 make_float4(float x, float y, float z, float w)

double向量

• double1 make_double1(double x)
• double2 make_double2(double x, double y)
• double3 make_double3(double x, double y, double z)
• double4 make_double4(double x, double y, double z, double w)

其他类型（short, long, longlong, uchar, ushort, uint, ulong, ulonglong）均有对应的make函数。

dim3构造

• dim3(unsigned int vx = 1, unsigned int vy = 1, unsigned int vz = 1)
• dim3(uint3 v)

---

14. 其他内建函数

性能计数器

• unsigned int __pm0(void) - 读取性能计数器0
• unsigned int __pm1(void) - 读取性能计数器1
• unsigned int __pm2(void) - 读取性能计数器2
• unsigned int __pm3(void) - 读取性能计数器3

特殊数学函数

• double fdivide(double a, double b) - IEEE兼容除法
• float fdividef(float a, float b) - IEEE兼容除法（单精度）
• double dadd(double a, double b, musaRoundMode mode) - 带舍入模式的加法
• double dmul(double a, double b, musaRoundMode mode) - 带舍入模式的乘法
• double fma(double a, double b, double c, musaRoundMode mode) - 带舍入模式的FMA

最值函数

• int/float/double min(T a, T b) - 最小值
• int/float/double max(T a, T b) - 最大值
• long long llmin(long long a, long long b)
• long long llmax(long long a, long long b)
• unsigned long long ullmin(unsigned long long a, unsigned long long b)
• unsigned long long ullmax(unsigned long long a, unsigned long long b)

绝对值

• int/long/long long/float/double abs(T x)
• long labs(long x)
• long long llabs(long long x)

Pi相关三角函数

• float cospi(float x) - cos(π*x)
• float sinpi(float x) - sin(π*x)
• void sincospi(double x, double* s, double* c) - 同时计算sin(πx)和cos(πx)

---

使用注意事项

1. 设备端限制: 大多数内建函数仅在设备代码中可用，需要在 __device__ 或 __global__ 函数中调用。

2. 精度与性能权衡:

   • 带 f 后缀的快速数学函数（如 __sinf）牺牲精度换取性能
   • 带舍入控制的函数提供精确的IEEE-754兼容结果

3. 编译选项影响:

   • -use_fast_math 会影响数学函数的精度和性能
   • -fmad=false 可禁用融合乘加优化

4. 内存访问优化:

   • 使用 __ldg() 读取只读数据可提高性能
   • 缓存提示函数需要适当的硬件支持

5. 原子操作作用域:

   • 默认原子操作作用于设备级
   • _block 后缀限制在线程块内
   • _system 后缀扩展到系统级（多GPU）

6. Warp操作要求:

   • Warp级函数要求参与线程在同一warp内
   • 使用正确的mask参数避免未定义行为