AMDGPU Device 函数传参详解摘要本文深入解析 AMDGPU Device 函数__device__的传参和调用约定Calling Convention。核心内容包括1区分 Kernel 与 Device 函数两条不同的 ABI 路径2详解CC_AMDGPU_Func传参规则和RetCC_AMDGPU_Func返回值约定3说明 Callee-saved 寄存器与 Caller-saved 寄存器的划分4梳理 Clang 前端AMDGPUABIInfo的分类规则5通过 HIP/IR/汇编示例展示端到端调用流程以及与 x86-64 SysV ABI 对比。适用于 LLVM/AMDGPU 后端开发、编译器优化及高性能计算开发者。适用范围llvm/lib/Target/AMDGPUGCN/SI 及以后GCNSubtarget参考来源AMDGPUCallingConv.td、AMDGPUISelLowering.cpp、SIRegisterInfo.cpp、AMDGPU.cpp等CPU 参考X86CallingConv.td中CC_X86_64_C/CSR_64_RT目录Kernel 与 Device 两条轨道Caller-save 与 Callee-saveCallingConvention 对照TableGen传参与返回C 挂钩SIRegisterInfo.cpp一次 Device Call 的端到端流程Clang 前端AMDGPUABIInfo示例HIP → IR → 汇编与 x86-64 SysV 对照调试与测试用例关键源文件索引1. Kernel 与 Device 两条路径AMDGPU Compute 程序里同时存在两类“函数”后端对它们的 ABI完全不同轨道IR CallingConv典型 C/HIP 来源参数从哪里来返回值Compute KernelAMDGPU_KERNEL(91) /SPIR_KERNEL__global__、OpenCL__kernelKernarg buffer常量段经 User SGPR 指针间接访问通常voidDevice/LeafC/Fast/Cold→CC_AMDGPU_Func__device__、kernel 内可调用的内部函数VGPR0–31/SGPR0–29inreg Scratch 栈VGPR0–31要点Kernel 形参由analyzeFormalArgumentsComputeAMDGPUISelLowering.cpp在 kernarg segment 内按字节偏移布局不占用CC_AMDGPU_Func的 VGPR/SGPR 编号槽。Device 函数是普通call/ret使用CC_AMDGPU_Func/RetCC_AMDGPU_Func与 CPU 上 SysV 的“寄存器 栈”模型类似但寄存器是 VGPR/SGPR栈是Scratch。Kernel 传参请参考AMDGPU KernelCompute传参详解2. Caller-save 与 Callee-save术语责任方含义Callee-saved被调函数若改写须在 prologue/epilogue save/restorecall 返回后 caller 可假定不变Caller-saved调用方callee 可破坏caller 若 call 后仍要旧值须 spillLLVM 两套表PEI vs Regalloc与 TableGen 一致机制API编译阶段CalleeSavedRegsgetCalleeSavedRegs()PEIcallee 若修改 CSR 须在 prologue/epilogue save/restoreRegMaskgetCallPreservedMask(C)ISelRegalloc跨 call 的 live 与 clobber 分析对C/Fast/ColdgetCallPreservedMask返回CSR_AMDGPU_RegMask或 GFX90A 变体。3. CallingConvention 对照CallingConvID典型来源入参 CC返回 CCRegMaskPEI CSRC0默认device 互调CC_AMDGPU_FuncRetCC_AMDGPU_FuncCSR_AMDGPU_RegMask条纹 VGPR/SGPRFast / Cold8 / 9LLVM 通用同 C同 C同 C同 CAMDGPU_Gfx—GfxCC_SI_GfxRetCC_SI_GfxCSR_AMDGPU_SI_Gfx_*Gfx CSR 表AMDGPU_KERNEL91Kernel见 Kernel 专篇voidN/A入口Entry 早退AMDGPU_CS_Chain—Mesh chainCC_AMDGPU_CS_CHAIN—AMDGPU_AllVGPRs链式特例CC_AMDGPU泛用在 SI 上先 delegate 到CC_SI_SHADERC再 delegate 到CC_AMDGPU_Func。4. TableGen传参与返回4.1CC_AMDGPU_Func入参inregSGPR0SGPR29非 inregVGPR0VGPR31溢出CCAssignToStack4,4byvalCCPassByVal4,4// AMDGPUCallingConv.td节选 def CC_AMDGPU_Func : CallingConv[ CCIfByValCCPassByVal4, 4, // 结构体 byval → 指针 栈布局 CCIfType[i1], CCPromoteToTypei32, CCIfType[i8, i16], CCIfExtendCCPromoteToTypei32, // inreg标量/小向量 → SGPR0..29 CCIfInRegCCIfType[f32, i32, f16, i16, v2i16, v2f16, bf16, v2bf16] , CCAssignToRegSGPR0..29, // 默认→ VGPR0..31 CCIfType[i32, f32, i16, f16, v2i16, v2f16, i1, bf16, v2bf16], CCAssignToRegVGPR0..31, // 寄存器用尽 → Scratch 栈4 字节对齐 CCIfType..., CCAssignToStack4, 4 ];4.1.1 默认路径VGPR 顺序分配无inreg时参数按32 位槽顺序占用VGPR0, VGPR1, …LLVM / C 类型占用 VGPR第 N 个i32参数i32/float1第 1 个 → v0第 2 个 → v1i64/double2v0:v12 x i32/2 x float2v0:v1i8/i16/i1提升为i32后 1带signext/zeroexthalf/bf161或按类型规则视 subtarget 与 true16寄存器上限ClangAMDGPUABIInfo中MaxNumRegsForArgsRet 1616 个 32 位槽 ≈ VGPR0–15 的“直接传参预算”超出部分间接传参或走栈见 第 7 节。4.1.2inreg路径SGPRIR 上带inreg属性的参数走SGPR0–SGPR29用于 wave-uniform 标量、部分指针、或图形/编译器显式标记的 uniform 数据。Callee 从s0, s1, …读入而不是 v0, v1。4.1.3 溢出与 byval情况行为实参超过可用 VGPR 槽CCAssignToStack4,4→Scratch经 buffer resource SGPR s32SPbyval结构体CCPassByVal4,4按 4 字节对齐在 caller 栈上布局传指针或拷贝过大 / 非平凡聚合Clang 常降为indirectprivate AS 指针等Scratch不是x86 式RSP向下增长由SIFrameLowering与s32stack pointer及 scratch resource SGPR 管理。4.2RetCC_AMDGPU_Func返回i1/i16先提升/扩展到i32。标量与小向量返回值一律VGPR0VGPR31按顺序占用。返回类型寄存器void无i32/floatv0i64/doublev0:v12 x float等连续 VGPRCallee 在ret前把结果写入约定 VGPRCaller 在s_swappc返回后从 v0等读取。4.3 Callee-savedCSR_AMDGPUdef CSR_AMDGPU_VGPRs : CalleeSavedRegs(add (sequence VGPR%u, 40, 47), (sequence VGPR%u, 56, 63), ... ); def CSR_AMDGPU_SGPRs : CalleeSavedRegs(add (sequence SGPR%u, 30, 39), (sequence SGPR%u, 48, 55), ... ); def CSR_AMDGPU : CalleeSavedRegs(add CSR_AMDGPU_VGPRs, CSR_AMDGPU_SGPRs);保留 SGPR不在 ABI 表内由SIRegisterInfo::reserveRegister处理SGPR用途s30–s31返回地址s_swappc/s_setpcs32Stack pointerScratchs33Frame pointers34Base pointerCC_SI_Gfx另保留SGPR0–3给 buffer descriptor。GFX90ACSR_AMDGPU_GFX90AInsts增加CSR_AMDGPU_AGPRsAGPR32–255。4.4 MustSpill 形式化LLVM 各目标通用与 X86 相同MustSpillCaller(CALL) { R | LiveAcrossCall(v, CALL) ∧ clobbersPhysReg(Mask, R) }VGPR0–31mask 中为clobbered→ caller 若 call 后仍需要旧值必须 spill。条纹 CSR VGPR/SGPRmask 中为preserved→ caller 通常不必为 call 单独保存callee 负责若其要改写。5. C 挂钩SIRegisterInfo.cppconstMCPhysReg*SIRegisterInfo::getCalleeSavedRegs(constMachineFunction*MF)const{switch(MF-getFunction().getCallingConv()){caseCallingConv::C:caseCallingConv::Fast:caseCallingConv::Cold:returnST.hasGFX90AInsts()?CSR_AMDGPU_GFX90AInsts_SaveList:CSR_AMDGPU_SaveList;caseCallingConv::AMDGPU_Gfx:return...CSR_AMDGPU_SI_Gfx_*_SaveList;caseCallingConv::AMDGPU_CS_ChainPreserve:returnCSR_AMDGPU_CS_ChainPreserve_SaveList;default:returnNoCalleeSavedReg;// Kernel 等}}constuint32_t*SIRegisterInfo::getCallPreservedMask(...,CallingConv::ID CC)const{switch(CC){caseCallingConv::C:caseCallingConv::Fast:caseCallingConv::Cold:returnCSR_AMDGPU_RegMask;// 或 GFX90A 变体caseCallingConv::AMDGPU_Gfx:returnCSR_AMDGPU_SI_Gfx_RegMask;caseCallingConv::AMDGPU_CS_Chain:caseCallingConv::AMDGPU_CS_ChainPreserve:returnAMDGPU_AllVGPRs_RegMask;// 假定不返回default:returnnullptr;}}6. 一次 Device Call 的端到端流程Caller例如 kernel 或其它 __device__ 函数 Callee__device__ 函数 ───────────────────────────────────────── ─────────────────────────── Clang: classifyArgumentType → 直接/间接/inreg ↓ LLVM IR: call foo(i32 %a) ; 默认 C calling conv ↓ ISel LowerCall: ① 按 CC_AMDGPU_Func 把实参放入 VGPR/SGPR 或建栈实参 ② 对 LiveAcross 且将被 clobber 的物理寄存器插入 spillRegMask ③ s_getpc 重定位 → 得到 callee 地址 ④ s_swappc_b64 s[30:31], s[addr:addr1] ; s30:s31 ← 返回 PC ↓ ↓ LowerFormalArguments: 从 v0/s0/栈 读形参 函数体... 返回值写入 v0... s_setpc_b64 s[30:31] ↓ ↓ ⑤ 从 v0 等读取返回值reload spill继续执行与 x86call压栈返回地址不同AMDGPU 用s_swappc交换s30:s31与目标地址返回时用s_setpc_b64 s[30:31]。7. Clang 前端AMDGPUABIInfo文件clang/lib/CodeGen/Targets/AMDGPU.cpp规则说明MaxNumRegsForArgsRet 16参数返回合计最多 16 个 32 位寄存器槽的直接传递预算小聚合 ≤8B可 pack 进 1–2 个 VGPRi16/i32/2×i32单元素 struct降为内部元素类型直接传递过大 structgetIndirectAliased等间接传递inreg由 Clang 根据类型/属性设置驱动 TableGen 的CCIfInReg→ SGPR 路径KernelclassifyKernelArgumentType与 Device 路径分离8. 示例HIP → IR → 汇编8.1 例 1单参数i32与返回值HIP概念代码__device__intadd_one(intx){returnx1;}__global__voidkernel(int*out){intradd_one(42);*outr;}Device 函数 IR简化define i32 add_one(i32 %x) { entry: %r add i32 %x, 1 ret i32 %r }Caller 侧典型 GCN 汇编摘自llvm/test/CodeGen/AMDGPU/call-argument-types.ll思路; 实参 42 → 第 1 个参数寄存器 VGPR0 v_mov_b32_e32 v0, 42 ; 解析 callee 地址并 call s_getpc_b64 s[4:5] s_add_u32 s4, s4, add_onerel32lo4 s_addc_u32 s5, s5, add_onerel32hi12 s_swappc_b64 s[30:31], s[4:5] ; 返回后返回值在 v0例如 43Callee 侧从v0读%x结果写v0s_setpc_b64 s[30:31]返回。8.2 例 2多参数按 VGPR 顺序HIP__device__floatmix3(floata,floatb,floatc){returnabc;}寄存器映射参数VGPRav0bv1cv2Caller概念汇编v_mov_b32_e32 v0, ... ; a v_mov_b32_e32 v1, ... ; b v_mov_b32_e32 v2, ... ; c s_swappc_b64 s[30:31], s[4:5] ; 返回v0 abc8.3 例 3inreg走 SGPRIRdefine void foo(i32 inreg %uniform_val) { ; 使用 s0 作为第 1 个 inreg 参数 ... }Caller将 uniform 值放入s0或当前 inreg 槽位对应的 SGPR再s_swappc。Callee从s0读取而非 v0。参考测试llvm/test/CodeGen/AMDGPU/call-args-inreg.ll、calling-conventions.llamdgpu_ps的inreg参数从 s0 读入后在 VGPR 上运算。8.4 例 4跨 call 的 Caller spillCSR VGPR当 caller 在v40属于条纹 callee-saved 区间上仍有 live 值且该值跨 call 存活时Lowering 会在 call 前后插入 scratch spill/reload。摘自llvm/test/CodeGen/AMDGPU/nested-calls.ll模式buffer_store_dword v40, off, s[0:3], s33 ; spill v40 v_mov_b32_e32 v0, 42 ; 设置 call 参数 s_swappc_b64 s[30:31], s[16:17] buffer_load_dword v40, off, s[0:3], s33 ; reload v40说明传参通道v0–v31与用于长期保存的条纹 VGPR职责不同RegMask live 分析决定是否在 call 点 spill。8.5 例 5与 Kernel 传参对比勿混淆Kernel IRdefine amdgpu_kernel void my_kernel(ptr addrspace(1) %out, i32 %n) { ... }项目KernelDevice 函数CallingConvamdgpu_kernel/AMDGPU_KERNELC参数布局analyzeFormalArgumentsCompute→ kernarg 字节偏移CC_AMDGPU_Func→ v0… / s0…入口加载从 User SGPRkernarg ptrs_load/buffer_load直接从 VGPR/SGPR 读能否被call否入口是Kernel 内调用__device__函数时kernel 侧仍用 Device CC 向 callee 传参如v_mov_b32 v0, 42s_swappc与 kernel 自身形参的 kernarg 布局无关。9. 与 x86-64 SysV 对照维度x86-64 SysV (CC_X86_64_C)AMDGPU Device (CC_AMDGPU_Func)标量整数/指针实参RDI, RSI, RDX, RCX, R8, R9 栈VGPR0–31默认Uniform / inreg无直接对应SGPR0–29浮点/向量实参XMM0–7 等多在VGPR返回值RAX RDX 等VGPR0–31Callee-savedRBX, RBP, R12–R15 等连续区间条纹VGPR/SGPR 多段栈RSPScratchs32Callcall 栈上返回地址s_swappcs30:s31GPU Kernel 入口无Kernarg buffer独立 ABI10. 调试与测试用例目的建议命令 / 测试单参数 / 多类型传参llvm/test/CodeGen/AMDGPU/call-argument-types.llinregllvm/test/CodeGen/AMDGPU/call-args-inreg.ll嵌套 call / spillllvm/test/CodeGen/AMDGPU/nested-calls.llCC 变体fastcc/coldcc/psllvm/test/CodeGen/AMDGPU/calling-conventions.llRegMask / preservedllvm/test/CodeGen/AMDGPU/call-preserved-registers.ll查看 IRclang -S -emit-llvm --cuda-device-only -O2 -o - file.cu查看汇编llc -mtripleamdgcn-amd-amdhsa -mcpugfx900 file.ll断点建议AMDGPUCallLowering::lowerCall/SIISelLowering::LowerCallSIRegisterInfo::getCallPreservedMaskAMDGPUABIInfo::classifyArgumentTypeClang11. 关键源文件索引文件内容AMDGPUCallingConv.tdCC_AMDGPU_Func、RetCC_AMDGPU_Func、CSR_AMDGPU*AMDGPUISelLowering.cppCCAssignFnForCall/CCAssignFnForReturnSIISelLowering.cppLowerCall、LowerFormalArgumentsSIRegisterInfo.cppgetCalleeSavedRegs、getCallPreservedMaskSIFrameLowering.cppScratch 帧、determineCalleeSavesclang/lib/CodeGen/Targets/AMDGPU.cppAMDGPUABIInfollvm/test/CodeGen/AMDGPU/*.ll汇编级回归与 FileCheck
AMDGPU Device 函数传参详解
AMDGPU Device 函数传参详解摘要本文深入解析 AMDGPU Device 函数__device__的传参和调用约定Calling Convention。核心内容包括1区分 Kernel 与 Device 函数两条不同的 ABI 路径2详解CC_AMDGPU_Func传参规则和RetCC_AMDGPU_Func返回值约定3说明 Callee-saved 寄存器与 Caller-saved 寄存器的划分4梳理 Clang 前端AMDGPUABIInfo的分类规则5通过 HIP/IR/汇编示例展示端到端调用流程以及与 x86-64 SysV ABI 对比。适用于 LLVM/AMDGPU 后端开发、编译器优化及高性能计算开发者。适用范围llvm/lib/Target/AMDGPUGCN/SI 及以后GCNSubtarget参考来源AMDGPUCallingConv.td、AMDGPUISelLowering.cpp、SIRegisterInfo.cpp、AMDGPU.cpp等CPU 参考X86CallingConv.td中CC_X86_64_C/CSR_64_RT目录Kernel 与 Device 两条轨道Caller-save 与 Callee-saveCallingConvention 对照TableGen传参与返回C 挂钩SIRegisterInfo.cpp一次 Device Call 的端到端流程Clang 前端AMDGPUABIInfo示例HIP → IR → 汇编与 x86-64 SysV 对照调试与测试用例关键源文件索引1. Kernel 与 Device 两条路径AMDGPU Compute 程序里同时存在两类“函数”后端对它们的 ABI完全不同轨道IR CallingConv典型 C/HIP 来源参数从哪里来返回值Compute KernelAMDGPU_KERNEL(91) /SPIR_KERNEL__global__、OpenCL__kernelKernarg buffer常量段经 User SGPR 指针间接访问通常voidDevice/LeafC/Fast/Cold→CC_AMDGPU_Func__device__、kernel 内可调用的内部函数VGPR0–31/SGPR0–29inreg Scratch 栈VGPR0–31要点Kernel 形参由analyzeFormalArgumentsComputeAMDGPUISelLowering.cpp在 kernarg segment 内按字节偏移布局不占用CC_AMDGPU_Func的 VGPR/SGPR 编号槽。Device 函数是普通call/ret使用CC_AMDGPU_Func/RetCC_AMDGPU_Func与 CPU 上 SysV 的“寄存器 栈”模型类似但寄存器是 VGPR/SGPR栈是Scratch。Kernel 传参请参考AMDGPU KernelCompute传参详解2. Caller-save 与 Callee-save术语责任方含义Callee-saved被调函数若改写须在 prologue/epilogue save/restorecall 返回后 caller 可假定不变Caller-saved调用方callee 可破坏caller 若 call 后仍要旧值须 spillLLVM 两套表PEI vs Regalloc与 TableGen 一致机制API编译阶段CalleeSavedRegsgetCalleeSavedRegs()PEIcallee 若修改 CSR 须在 prologue/epilogue save/restoreRegMaskgetCallPreservedMask(C)ISelRegalloc跨 call 的 live 与 clobber 分析对C/Fast/ColdgetCallPreservedMask返回CSR_AMDGPU_RegMask或 GFX90A 变体。3. CallingConvention 对照CallingConvID典型来源入参 CC返回 CCRegMaskPEI CSRC0默认device 互调CC_AMDGPU_FuncRetCC_AMDGPU_FuncCSR_AMDGPU_RegMask条纹 VGPR/SGPRFast / Cold8 / 9LLVM 通用同 C同 C同 C同 CAMDGPU_Gfx—GfxCC_SI_GfxRetCC_SI_GfxCSR_AMDGPU_SI_Gfx_*Gfx CSR 表AMDGPU_KERNEL91Kernel见 Kernel 专篇voidN/A入口Entry 早退AMDGPU_CS_Chain—Mesh chainCC_AMDGPU_CS_CHAIN—AMDGPU_AllVGPRs链式特例CC_AMDGPU泛用在 SI 上先 delegate 到CC_SI_SHADERC再 delegate 到CC_AMDGPU_Func。4. TableGen传参与返回4.1CC_AMDGPU_Func入参inregSGPR0SGPR29非 inregVGPR0VGPR31溢出CCAssignToStack4,4byvalCCPassByVal4,4// AMDGPUCallingConv.td节选 def CC_AMDGPU_Func : CallingConv[ CCIfByValCCPassByVal4, 4, // 结构体 byval → 指针 栈布局 CCIfType[i1], CCPromoteToTypei32, CCIfType[i8, i16], CCIfExtendCCPromoteToTypei32, // inreg标量/小向量 → SGPR0..29 CCIfInRegCCIfType[f32, i32, f16, i16, v2i16, v2f16, bf16, v2bf16] , CCAssignToRegSGPR0..29, // 默认→ VGPR0..31 CCIfType[i32, f32, i16, f16, v2i16, v2f16, i1, bf16, v2bf16], CCAssignToRegVGPR0..31, // 寄存器用尽 → Scratch 栈4 字节对齐 CCIfType..., CCAssignToStack4, 4 ];4.1.1 默认路径VGPR 顺序分配无inreg时参数按32 位槽顺序占用VGPR0, VGPR1, …LLVM / C 类型占用 VGPR第 N 个i32参数i32/float1第 1 个 → v0第 2 个 → v1i64/double2v0:v12 x i32/2 x float2v0:v1i8/i16/i1提升为i32后 1带signext/zeroexthalf/bf161或按类型规则视 subtarget 与 true16寄存器上限ClangAMDGPUABIInfo中MaxNumRegsForArgsRet 1616 个 32 位槽 ≈ VGPR0–15 的“直接传参预算”超出部分间接传参或走栈见 第 7 节。4.1.2inreg路径SGPRIR 上带inreg属性的参数走SGPR0–SGPR29用于 wave-uniform 标量、部分指针、或图形/编译器显式标记的 uniform 数据。Callee 从s0, s1, …读入而不是 v0, v1。4.1.3 溢出与 byval情况行为实参超过可用 VGPR 槽CCAssignToStack4,4→Scratch经 buffer resource SGPR s32SPbyval结构体CCPassByVal4,4按 4 字节对齐在 caller 栈上布局传指针或拷贝过大 / 非平凡聚合Clang 常降为indirectprivate AS 指针等Scratch不是x86 式RSP向下增长由SIFrameLowering与s32stack pointer及 scratch resource SGPR 管理。4.2RetCC_AMDGPU_Func返回i1/i16先提升/扩展到i32。标量与小向量返回值一律VGPR0VGPR31按顺序占用。返回类型寄存器void无i32/floatv0i64/doublev0:v12 x float等连续 VGPRCallee 在ret前把结果写入约定 VGPRCaller 在s_swappc返回后从 v0等读取。4.3 Callee-savedCSR_AMDGPUdef CSR_AMDGPU_VGPRs : CalleeSavedRegs(add (sequence VGPR%u, 40, 47), (sequence VGPR%u, 56, 63), ... ); def CSR_AMDGPU_SGPRs : CalleeSavedRegs(add (sequence SGPR%u, 30, 39), (sequence SGPR%u, 48, 55), ... ); def CSR_AMDGPU : CalleeSavedRegs(add CSR_AMDGPU_VGPRs, CSR_AMDGPU_SGPRs);保留 SGPR不在 ABI 表内由SIRegisterInfo::reserveRegister处理SGPR用途s30–s31返回地址s_swappc/s_setpcs32Stack pointerScratchs33Frame pointers34Base pointerCC_SI_Gfx另保留SGPR0–3给 buffer descriptor。GFX90ACSR_AMDGPU_GFX90AInsts增加CSR_AMDGPU_AGPRsAGPR32–255。4.4 MustSpill 形式化LLVM 各目标通用与 X86 相同MustSpillCaller(CALL) { R | LiveAcrossCall(v, CALL) ∧ clobbersPhysReg(Mask, R) }VGPR0–31mask 中为clobbered→ caller 若 call 后仍需要旧值必须 spill。条纹 CSR VGPR/SGPRmask 中为preserved→ caller 通常不必为 call 单独保存callee 负责若其要改写。5. C 挂钩SIRegisterInfo.cppconstMCPhysReg*SIRegisterInfo::getCalleeSavedRegs(constMachineFunction*MF)const{switch(MF-getFunction().getCallingConv()){caseCallingConv::C:caseCallingConv::Fast:caseCallingConv::Cold:returnST.hasGFX90AInsts()?CSR_AMDGPU_GFX90AInsts_SaveList:CSR_AMDGPU_SaveList;caseCallingConv::AMDGPU_Gfx:return...CSR_AMDGPU_SI_Gfx_*_SaveList;caseCallingConv::AMDGPU_CS_ChainPreserve:returnCSR_AMDGPU_CS_ChainPreserve_SaveList;default:returnNoCalleeSavedReg;// Kernel 等}}constuint32_t*SIRegisterInfo::getCallPreservedMask(...,CallingConv::ID CC)const{switch(CC){caseCallingConv::C:caseCallingConv::Fast:caseCallingConv::Cold:returnCSR_AMDGPU_RegMask;// 或 GFX90A 变体caseCallingConv::AMDGPU_Gfx:returnCSR_AMDGPU_SI_Gfx_RegMask;caseCallingConv::AMDGPU_CS_Chain:caseCallingConv::AMDGPU_CS_ChainPreserve:returnAMDGPU_AllVGPRs_RegMask;// 假定不返回default:returnnullptr;}}6. 一次 Device Call 的端到端流程Caller例如 kernel 或其它 __device__ 函数 Callee__device__ 函数 ───────────────────────────────────────── ─────────────────────────── Clang: classifyArgumentType → 直接/间接/inreg ↓ LLVM IR: call foo(i32 %a) ; 默认 C calling conv ↓ ISel LowerCall: ① 按 CC_AMDGPU_Func 把实参放入 VGPR/SGPR 或建栈实参 ② 对 LiveAcross 且将被 clobber 的物理寄存器插入 spillRegMask ③ s_getpc 重定位 → 得到 callee 地址 ④ s_swappc_b64 s[30:31], s[addr:addr1] ; s30:s31 ← 返回 PC ↓ ↓ LowerFormalArguments: 从 v0/s0/栈 读形参 函数体... 返回值写入 v0... s_setpc_b64 s[30:31] ↓ ↓ ⑤ 从 v0 等读取返回值reload spill继续执行与 x86call压栈返回地址不同AMDGPU 用s_swappc交换s30:s31与目标地址返回时用s_setpc_b64 s[30:31]。7. Clang 前端AMDGPUABIInfo文件clang/lib/CodeGen/Targets/AMDGPU.cpp规则说明MaxNumRegsForArgsRet 16参数返回合计最多 16 个 32 位寄存器槽的直接传递预算小聚合 ≤8B可 pack 进 1–2 个 VGPRi16/i32/2×i32单元素 struct降为内部元素类型直接传递过大 structgetIndirectAliased等间接传递inreg由 Clang 根据类型/属性设置驱动 TableGen 的CCIfInReg→ SGPR 路径KernelclassifyKernelArgumentType与 Device 路径分离8. 示例HIP → IR → 汇编8.1 例 1单参数i32与返回值HIP概念代码__device__intadd_one(intx){returnx1;}__global__voidkernel(int*out){intradd_one(42);*outr;}Device 函数 IR简化define i32 add_one(i32 %x) { entry: %r add i32 %x, 1 ret i32 %r }Caller 侧典型 GCN 汇编摘自llvm/test/CodeGen/AMDGPU/call-argument-types.ll思路; 实参 42 → 第 1 个参数寄存器 VGPR0 v_mov_b32_e32 v0, 42 ; 解析 callee 地址并 call s_getpc_b64 s[4:5] s_add_u32 s4, s4, add_onerel32lo4 s_addc_u32 s5, s5, add_onerel32hi12 s_swappc_b64 s[30:31], s[4:5] ; 返回后返回值在 v0例如 43Callee 侧从v0读%x结果写v0s_setpc_b64 s[30:31]返回。8.2 例 2多参数按 VGPR 顺序HIP__device__floatmix3(floata,floatb,floatc){returnabc;}寄存器映射参数VGPRav0bv1cv2Caller概念汇编v_mov_b32_e32 v0, ... ; a v_mov_b32_e32 v1, ... ; b v_mov_b32_e32 v2, ... ; c s_swappc_b64 s[30:31], s[4:5] ; 返回v0 abc8.3 例 3inreg走 SGPRIRdefine void foo(i32 inreg %uniform_val) { ; 使用 s0 作为第 1 个 inreg 参数 ... }Caller将 uniform 值放入s0或当前 inreg 槽位对应的 SGPR再s_swappc。Callee从s0读取而非 v0。参考测试llvm/test/CodeGen/AMDGPU/call-args-inreg.ll、calling-conventions.llamdgpu_ps的inreg参数从 s0 读入后在 VGPR 上运算。8.4 例 4跨 call 的 Caller spillCSR VGPR当 caller 在v40属于条纹 callee-saved 区间上仍有 live 值且该值跨 call 存活时Lowering 会在 call 前后插入 scratch spill/reload。摘自llvm/test/CodeGen/AMDGPU/nested-calls.ll模式buffer_store_dword v40, off, s[0:3], s33 ; spill v40 v_mov_b32_e32 v0, 42 ; 设置 call 参数 s_swappc_b64 s[30:31], s[16:17] buffer_load_dword v40, off, s[0:3], s33 ; reload v40说明传参通道v0–v31与用于长期保存的条纹 VGPR职责不同RegMask live 分析决定是否在 call 点 spill。8.5 例 5与 Kernel 传参对比勿混淆Kernel IRdefine amdgpu_kernel void my_kernel(ptr addrspace(1) %out, i32 %n) { ... }项目KernelDevice 函数CallingConvamdgpu_kernel/AMDGPU_KERNELC参数布局analyzeFormalArgumentsCompute→ kernarg 字节偏移CC_AMDGPU_Func→ v0… / s0…入口加载从 User SGPRkernarg ptrs_load/buffer_load直接从 VGPR/SGPR 读能否被call否入口是Kernel 内调用__device__函数时kernel 侧仍用 Device CC 向 callee 传参如v_mov_b32 v0, 42s_swappc与 kernel 自身形参的 kernarg 布局无关。9. 与 x86-64 SysV 对照维度x86-64 SysV (CC_X86_64_C)AMDGPU Device (CC_AMDGPU_Func)标量整数/指针实参RDI, RSI, RDX, RCX, R8, R9 栈VGPR0–31默认Uniform / inreg无直接对应SGPR0–29浮点/向量实参XMM0–7 等多在VGPR返回值RAX RDX 等VGPR0–31Callee-savedRBX, RBP, R12–R15 等连续区间条纹VGPR/SGPR 多段栈RSPScratchs32Callcall 栈上返回地址s_swappcs30:s31GPU Kernel 入口无Kernarg buffer独立 ABI10. 调试与测试用例目的建议命令 / 测试单参数 / 多类型传参llvm/test/CodeGen/AMDGPU/call-argument-types.llinregllvm/test/CodeGen/AMDGPU/call-args-inreg.ll嵌套 call / spillllvm/test/CodeGen/AMDGPU/nested-calls.llCC 变体fastcc/coldcc/psllvm/test/CodeGen/AMDGPU/calling-conventions.llRegMask / preservedllvm/test/CodeGen/AMDGPU/call-preserved-registers.ll查看 IRclang -S -emit-llvm --cuda-device-only -O2 -o - file.cu查看汇编llc -mtripleamdgcn-amd-amdhsa -mcpugfx900 file.ll断点建议AMDGPUCallLowering::lowerCall/SIISelLowering::LowerCallSIRegisterInfo::getCallPreservedMaskAMDGPUABIInfo::classifyArgumentTypeClang11. 关键源文件索引文件内容AMDGPUCallingConv.tdCC_AMDGPU_Func、RetCC_AMDGPU_Func、CSR_AMDGPU*AMDGPUISelLowering.cppCCAssignFnForCall/CCAssignFnForReturnSIISelLowering.cppLowerCall、LowerFormalArgumentsSIRegisterInfo.cppgetCalleeSavedRegs、getCallPreservedMaskSIFrameLowering.cppScratch 帧、determineCalleeSavesclang/lib/CodeGen/Targets/AMDGPU.cppAMDGPUABIInfollvm/test/CodeGen/AMDGPU/*.ll汇编级回归与 FileCheck
