https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/220197
>From 42a2d361d217336db99ba5b14cc8291eec214889 Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Tue, 1 Sep 2026 14:31:15 +0530 Subject: [PATCH 1/2] [CIR] Wire AMDGPU into the call-convention lowering pass --- clang/include/clang/CIR/Dialect/Passes.h | 2 +- clang/include/clang/CIR/Dialect/Passes.td | 4 +- .../Transforms/CallConvLoweringPass.cpp | 169 ++++++++++-------- clang/lib/CIR/Lowering/CIRPasses.cpp | 6 +- .../CodeGenHIP/amdgcn-buffer-rsrc-type.hip | 8 +- .../CodeGenHIP/cleanup-alloca-addrspace.hip | 8 +- .../abi-lowering/amdgpu-scalars.cir | 105 +++++++++++ .../x86_64-lang-addrspace-nyi.cir | 2 +- .../abi-lowering/x86_64-variadic-nyi.cir | 2 +- 9 files changed, 224 insertions(+), 82 deletions(-) create mode 100644 clang/test/CIR/Transforms/abi-lowering/amdgpu-scalars.cir diff --git a/clang/include/clang/CIR/Dialect/Passes.h b/clang/include/clang/CIR/Dialect/Passes.h index dedfc9fa351796..117822486cc3a2 100644 --- a/clang/include/clang/CIR/Dialect/Passes.h +++ b/clang/include/clang/CIR/Dialect/Passes.h @@ -20,7 +20,7 @@ namespace cir { /// The ABI target whose calling-convention rules drive CallConvLowering. /// None is the unset state used when the pass runs in classification-attr /// mode instead of selecting a target. -enum class CallConvTarget { None, Test, X86_64 }; +enum class CallConvTarget { None, Test, X86_64, AMDGPU }; } // namespace cir namespace mlir { diff --git a/clang/include/clang/CIR/Dialect/Passes.td b/clang/include/clang/CIR/Dialect/Passes.td index 6994aa8a4b18ad..843e217e7eef88 100644 --- a/clang/include/clang/CIR/Dialect/Passes.td +++ b/clang/include/clang/CIR/Dialect/Passes.td @@ -246,7 +246,9 @@ def CallConvLowering : Pass<"cir-call-conv-lowering", "mlir::ModuleOp"> { clEnumValN(cir::CallConvTarget::Test, "test", "MLIR test ABI target"), clEnumValN(cir::CallConvTarget::X86_64, "x86_64", - "x86_64 System V") + "x86_64 System V"), + clEnumValN(cir::CallConvTarget::AMDGPU, "amdgpu", + "AMDGPU") )}]>, Option<"classificationAttr", "classification-attr", "std::string", /*default=*/"\"\"", diff --git a/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp b/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp index 4c543191f036d8..42efe5358b5dda 100644 --- a/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp +++ b/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp @@ -338,7 +338,7 @@ static mlir::Type abiTypeToCIR(const llvm::abi::Type *ty, MLIRContext *ctx) { .Default([](const llvm::abi::Type *) -> mlir::Type { return nullptr; }); } -/// Map a CIR type to an llvm::abi::Type. classifyX86_64Function pre-filters +/// Map a CIR type to an llvm::abi::Type. classifyAbiSignature pre-filters /// the signature, so only the scalar and struct/array types handled here can /// reach this function. static const llvm::abi::Type *mapCIRType(mlir::Type type, @@ -533,7 +533,7 @@ static const llvm::abi::Type *mapCIRType(mlir::Type type, }) .Default([](mlir::Type) -> const llvm::abi::Type * { llvm_unreachable( - "mapCIRType: type not pre-filtered by classifyX86_64Function"); + "mapCIRType: type not pre-filtered by classifyAbiSignature"); }); } @@ -641,20 +641,31 @@ static llvm::abi::RequiredArgs requiredArgs(cir::FuncType fnTy) { return llvm::abi::RequiredArgs(fnTy.getNumInputs()); } -/// Classify an x86_64 SysV signature (return type + argument types) using the -/// LLVM ABI library. Shared by the cir.func path, the variadic-call path and -/// the indirect-call path (the latter classifies from the callee function +/// Map a cir.func calling convention to the llvm::CallingConv the ABI +/// classifier keys on. Only the AMDGPU kernel convention changes +/// classification today; every other convention is classified as C. +static llvm::CallingConv::ID abiCallingConv(cir::CallingConv cc) { + if (cc == cir::CallingConv::AMDGPUKernel) + return llvm::CallingConv::AMDGPU_KERNEL; + return llvm::CallingConv::C; +} + +/// Classify a signature (return type + argument types) for \p targetInfo using +/// the LLVM ABI library. Shared by the cir.func path, the variadic-call path +/// and the indirect-call path (the latter classifies from the callee function /// pointer's pointee FuncType). \p required marks where the declared /// parameters in \p inputs end. The classifier treats every argument past that -/// point as passed through an ellipsis. Returns std::nullopt and emits an NYI -/// error via \p emitError if the signature uses a type the bridge does not -/// handle yet. -static std::optional<FunctionClassification> classifyX86_64Signature( - mlir::Type retCIR, mlir::TypeRange inputs, llvm::abi::RequiredArgs required, - MLIRContext *ctx, const DataLayout &dl, - mlir::abi::ABITypeMapper &typeMapper, - const llvm::abi::TargetInfo &targetInfo, ModuleOp modOp, - llvm::function_ref<mlir::InFlightDiagnostic()> emitError) { +/// point as passed through an ellipsis. \p callConv selects the ABI convention +/// (e.g. AMDGPU_KERNEL), which some targets classify differently. Returns +/// std::nullopt and emits an NYI error via \p emitError if the signature uses a +/// type the bridge does not handle yet. +static std::optional<FunctionClassification> +classifyAbiSignature(mlir::Type retCIR, mlir::TypeRange inputs, + llvm::abi::RequiredArgs required, MLIRContext *ctx, + const DataLayout &dl, mlir::abi::ABITypeMapper &typeMapper, + const llvm::abi::TargetInfo &targetInfo, + llvm::CallingConv::ID callConv, ModuleOp modOp, + llvm::function_ref<mlir::InFlightDiagnostic()> emitError) { assert(retCIR && "signature return type must be non-null"); assert((!required.allowsOptionalArgs() || required.getNumRequiredArgs() <= inputs.size()) && @@ -664,9 +675,8 @@ static std::optional<FunctionClassification> classifyX86_64Signature( auto reject = [&](mlir::Type t) -> bool { if (isSupportedType(t, dl)) return false; - emitError() - << "x86_64 calling-convention lowering not yet implemented for type " - << t; + emitError() << "calling-convention lowering not yet implemented for type " + << t; return true; }; if (!voidRet && reject(retCIR)) @@ -682,15 +692,15 @@ static std::optional<FunctionClassification> classifyX86_64Signature( for (mlir::Type a : inputs) argAbi.push_back(mapCIRType(a, typeMapper, dl, modOp)); - std::unique_ptr<llvm::abi::FunctionInfo> fi = llvm::abi::FunctionInfo::create( - llvm::CallingConv::C, retAbi, argAbi, required); + std::unique_ptr<llvm::abi::FunctionInfo> fi = + llvm::abi::FunctionInfo::create(callConv, retAbi, argAbi, required); targetInfo.computeInfo(*fi); // convertABIArgInfo returns nullopt when the classifier picks a coercion this // bridge cannot represent. auto nyiCoercion = [&](mlir::Type t) { - emitError() << "x86_64 calling-convention lowering not yet " - "implemented for the ABI coercion of type " + emitError() << "calling-convention lowering not yet implemented for the " + "ABI coercion of type " << t; }; @@ -745,19 +755,19 @@ static llvm::abi::X86AVXABILevel funcAvxLevel(cir::FuncOp func, return avx ? std::max(base, llvm::abi::X86AVXABILevel::AVX) : base; } -/// Classify a cir.func for x86_64 SysV using the LLVM ABI library. Returns +/// Classify a cir.func for \p targetInfo using the LLVM ABI library. Returns /// std::nullopt and emits an NYI error if the signature uses a type the bridge /// does not handle yet. static std::optional<FunctionClassification> -classifyX86_64Function(cir::FuncOp func, const DataLayout &dl, - mlir::abi::ABITypeMapper &typeMapper, - const llvm::abi::TargetInfo &targetInfo, - ModuleOp modOp) { +classifyAbiFunction(cir::FuncOp func, const DataLayout &dl, + mlir::abi::ABITypeMapper &typeMapper, + const llvm::abi::TargetInfo &targetInfo, ModuleOp modOp) { cir::FuncType fnTy = func.getFunctionType(); - return classifyX86_64Signature(fnTy.getReturnType(), fnTy.getInputs(), - requiredArgs(fnTy), func->getContext(), dl, - typeMapper, targetInfo, modOp, - [&]() { return func.emitOpError(); }); + return classifyAbiSignature(fnTy.getReturnType(), fnTy.getInputs(), + requiredArgs(fnTy), func->getContext(), dl, + typeMapper, targetInfo, + abiCallingConv(func.getCallingConv()), modOp, + [&]() { return func.emitOpError(); }); } /// Classify the single type fetched by a `cir.va_arg` as an unnamed argument @@ -766,11 +776,11 @@ classifyX86_64Function(cir::FuncOp func, const DataLayout &dl, static std::optional<ArgClassification> classifyX86_64VarArgType( mlir::Type ty, MLIRContext *ctx, const DataLayout &dl, mlir::abi::ABITypeMapper &typeMapper, - const llvm::abi::TargetInfo &targetInfo, ModuleOp modOp, - llvm::function_ref<mlir::InFlightDiagnostic()> emitError) { - std::optional<FunctionClassification> fc = classifyX86_64Signature( + const llvm::abi::TargetInfo &targetInfo, llvm::CallingConv::ID callConv, + ModuleOp modOp, llvm::function_ref<mlir::InFlightDiagnostic()> emitError) { + std::optional<FunctionClassification> fc = classifyAbiSignature( cir::VoidType::get(ctx), mlir::TypeRange(ty), llvm::abi::RequiredArgs(0), - ctx, dl, typeMapper, targetInfo, modOp, emitError); + ctx, dl, typeMapper, targetInfo, callConv, modOp, emitError); if (!fc) return std::nullopt; return fc->argInfos[0]; @@ -783,17 +793,19 @@ static std::optional<ArgClassification> classifyX86_64VarArgType( /// small struct is passed in registers early in the list and in memory once /// the integer registers are gone. Classifying from the call's operands /// rather than the callee's signature is what makes that accounting right. -static std::optional<FunctionClassification> classifyX86_64VariadicCall( - cir::CIRCallOpInterface call, cir::FuncType calleeTy, const DataLayout &dl, - mlir::abi::ABITypeMapper &typeMapper, - const llvm::abi::TargetInfo &targetInfo, ModuleOp modOp) { +static std::optional<FunctionClassification> +classifyAbiVariadicCall(cir::CIRCallOpInterface call, cir::FuncType calleeTy, + const DataLayout &dl, + mlir::abi::ABITypeMapper &typeMapper, + const llvm::abi::TargetInfo &targetInfo, + llvm::CallingConv::ID callConv, ModuleOp modOp) { assert(calleeTy.isVarArg() && "only a variadic callee can take more operands than it declares"); Operation *op = call.getOperation(); - return classifyX86_64Signature( + return classifyAbiSignature( calleeTy.getReturnType(), call.getArgOperands().getTypes(), requiredArgs(calleeTy), op->getContext(), dl, typeMapper, targetInfo, - modOp, [&]() { return op->emitOpError(); }); + callConv, modOp, [&]() { return op->emitOpError(); }); } /// Whether \p fc gives the callee access to memory through a pointer the ABI @@ -952,10 +964,18 @@ void CallConvLoweringPass::runOnOperation() { static constexpr unsigned numAvxLevels = static_cast<unsigned>(llvm::abi::X86AVXABILevel::Last) + 1; bool isX86 = target == cir::CallConvTarget::X86_64; - std::optional<mlir::abi::ABITypeMapper> x86TypeMapper; + bool isAMDGPU = target == cir::CallConvTarget::AMDGPU; + // The x86_64 and AMDGPU drivers classify with the LLVM ABI library and share + // the type mapper; the test and classification-attr drivers do neither. + bool useAbiClassifier = isX86 || isAMDGPU; + std::optional<mlir::abi::ABITypeMapper> abiTypeMapper; std::array<std::unique_ptr<llvm::abi::TargetInfo>, numAvxLevels> x86Targets; - if (isX86) - x86TypeMapper.emplace(dl); + std::unique_ptr<llvm::abi::TargetInfo> amdgpuTarget; + if (useAbiClassifier) + abiTypeMapper.emplace(dl); + if (isAMDGPU) + amdgpuTarget = + llvm::abi::createAMDGPUTargetInfo(abiTypeMapper->getTypeBuilder()); auto x86TargetFor = [&](llvm::abi::X86AVXABILevel level) -> const llvm::abi::TargetInfo & { assert(static_cast<unsigned>(level) < numAvxLevels && @@ -964,7 +984,7 @@ void CallConvLoweringPass::runOnOperation() { x86Targets[static_cast<unsigned>(level)]; if (!slot) slot = llvm::abi::createX86_64TargetInfo( - x86TypeMapper->getTypeBuilder(), level, + abiTypeMapper->getTypeBuilder(), level, /*Has64BitPointers=*/true, x86AbiCompat); return *slot; }; @@ -974,6 +994,13 @@ void CallConvLoweringPass::runOnOperation() { return baseAvxLevel; return funcAvxLevel(func, baseAvxLevel); }; + // The classifier to use for a function/call reached from \p func. AMDGPU has + // a single classifier; x86_64 picks the one matching func's AVX level. + auto targetInfoFor = [&](cir::FuncOp func) -> const llvm::abi::TargetInfo & { + if (isAMDGPU) + return *amdgpuTarget; + return x86TargetFor(avxLevelFor(func)); + }; // Classify every cir.func up front. No IR mutation happens here, so // later walks can consult any function's classification regardless of @@ -991,9 +1018,9 @@ void CallConvLoweringPass::runOnOperation() { llvm::any_of(fnTy.getInputs(), hasIncompleteRecordByValue))) return; std::optional<FunctionClassification> fc; - if (isX86) - fc = classifyX86_64Function(f, dl, *x86TypeMapper, - x86TargetFor(avxLevelFor(f)), moduleOp); + if (useAbiClassifier) + fc = classifyAbiFunction(f, dl, *abiTypeMapper, targetInfoFor(f), + moduleOp); else fc = classifyFunction(f, dl, target, classificationAttr); if (!fc) { @@ -1030,11 +1057,12 @@ void CallConvLoweringPass::runOnOperation() { return; callers[callee].push_back(op); - // Only the x86_64 driver classifies per call site. Under the other + // Only the ABI-library drivers classify per call site. Under the other // drivers the classification comes from a fixed per-function source, so // such a call stays short a classification and rewriteCallSite reports it. cir::FuncType calleeTy = callee.getFunctionType(); - if (!isX86 || call.getNumArgOperands() <= calleeTy.getNumInputs()) + if (!useAbiClassifier || + call.getNumArgOperands() <= calleeTy.getNumInputs()) return; // A callee declared without a prototype also takes more operands than it // declares, and the verifier allows it. Those extra arguments are named @@ -1051,9 +1079,9 @@ void CallConvLoweringPass::runOnOperation() { // Classic instead arranges every call site from the caller and reports a // caller whose level disagrees with its callee in checkFunctionCallABI, // which has no equivalent here yet. - std::optional<FunctionClassification> fc = - classifyX86_64VariadicCall(call, calleeTy, dl, *x86TypeMapper, - x86TargetFor(avxLevelFor(callee)), moduleOp); + std::optional<FunctionClassification> fc = classifyAbiVariadicCall( + call, calleeTy, dl, *abiTypeMapper, targetInfoFor(callee), + abiCallingConv(callee.getCallingConv()), moduleOp); if (!fc) { anyFailed = true; return; @@ -1170,31 +1198,31 @@ void CallConvLoweringPass::runOnOperation() { [&](mlir::TypeRange argTypes) -> std::optional<FunctionClassification> { // A callee resolved at run time carries no features of its own, so the // level comes from the function containing the call, which is the - // declaration classic arranges every call site from. - if (isX86) - return classifyX86_64Signature( + // declaration classic arranges every call site from. An indirect callee + // is an ordinary function pointer, so it uses the C convention. + if (useAbiClassifier) + return classifyAbiSignature( funcTy.getReturnType(), argTypes, requiredArgs(funcTy), ctx, dl, - *x86TypeMapper, - x86TargetFor(avxLevelFor(c->getParentOfType<cir::FuncOp>())), - moduleOp, [&]() { return c->emitOpError(); }); + *abiTypeMapper, targetInfoFor(c->getParentOfType<cir::FuncOp>()), + llvm::CallingConv::C, moduleOp, [&]() { return c->emitOpError(); }); return withReturnVoidness( mlir::abi::test::classify(argTypes, funcTy.getReturnType(), dl), funcTy.getReturnType()); }; // An argument passed through an ellipsis has no counterpart in the - // pointee's parameter list. On x86_64 such a call is classified from its - // own operands, as a direct one is. Under target=test it is classified - // from the pointee alone, and rewriteCallSite reports the call. + // pointee's parameter list. Under the ABI-library drivers such a call is + // classified from its own operands, as a direct one is. Under target=test + // it is classified from the pointee alone, and rewriteCallSite reports the + // call. bool classifyCallSite = - isX86 && c.getNumArgOperands() > funcTy.getNumInputs(); + useAbiClassifier && c.getNumArgOperands() > funcTy.getNumInputs(); std::optional<FunctionClassification> fc = - classifyCallSite - ? classifyX86_64VariadicCall( - c, funcTy, dl, *x86TypeMapper, - x86TargetFor(avxLevelFor(c->getParentOfType<cir::FuncOp>())), - moduleOp) - : classifySignature(funcTy.getInputs()); + classifyCallSite ? classifyAbiVariadicCall( + c, funcTy, dl, *abiTypeMapper, + targetInfoFor(c->getParentOfType<cir::FuncOp>()), + llvm::CallingConv::C, moduleOp) + : classifySignature(funcTy.getInputs()); if (!fc) { signalPassFailure(); return; @@ -1215,8 +1243,9 @@ void CallConvLoweringPass::runOnOperation() { for (cir::VAArgOp v : vaArgs) { cir::FuncOp enclosing = v->getParentOfType<cir::FuncOp>(); std::optional<ArgClassification> ac = classifyX86_64VarArgType( - v.getType(), ctx, dl, *x86TypeMapper, - x86TargetFor(avxLevelFor(enclosing)), moduleOp, + v.getType(), ctx, dl, *abiTypeMapper, + x86TargetFor(avxLevelFor(enclosing)), + abiCallingConv(enclosing.getCallingConv()), moduleOp, [&]() { return v->emitOpError(); }); if (!ac) { signalPassFailure(); diff --git a/clang/lib/CIR/Lowering/CIRPasses.cpp b/clang/lib/CIR/Lowering/CIRPasses.cpp index 6e75ce03d14ffd..5e49a3f5601870 100644 --- a/clang/lib/CIR/Lowering/CIRPasses.cpp +++ b/clang/lib/CIR/Lowering/CIRPasses.cpp @@ -26,6 +26,8 @@ static CallConvTarget getCallConvTarget(const llvm::Triple &triple) { // Windows is not supported. UEFI shares its convention. if (triple.getArch() == llvm::Triple::x86_64 && !triple.isOSWindowsOrUEFI()) return CallConvTarget::X86_64; + if (triple.isAMDGPU()) + return CallConvTarget::AMDGPU; return CallConvTarget::None; } @@ -126,8 +128,8 @@ runCIRToCIRPasses(mlir::ModuleOp theModule, mlir::MLIRContext &mlirContext, if (enableCallConvLowering) { // CallConvLowering rewrites signatures and call sites using the classifier, // so it must run after CXXABILowering has lowered C++ ABI types to plain - // records the classifier can handle. Only the x86_64 System V classifier - // is implemented; other targets are left unchanged. + // records the classifier can handle. Only the x86_64 System V and AMDGPU + // classifiers are implemented. Other targets are left unchanged. CallConvTarget target = getCallConvTarget(triple); if (target != CallConvTarget::None) { // Source the ABI-compatibility version from the module's serialized diff --git a/clang/test/CIR/CodeGenHIP/amdgcn-buffer-rsrc-type.hip b/clang/test/CIR/CodeGenHIP/amdgcn-buffer-rsrc-type.hip index e03fa4db5c20df..0045a731c579f6 100644 --- a/clang/test/CIR/CodeGenHIP/amdgcn-buffer-rsrc-type.hip +++ b/clang/test/CIR/CodeGenHIP/amdgcn-buffer-rsrc-type.hip @@ -1,12 +1,16 @@ #include "../CodeGenCUDA/Inputs/cuda.h" // REQUIRES: amdgpu-registered-target + +// TODO(cir): drop -fno-clangir-call-conv-lowering once CallConvLowering +// supports the AMDGPU direct-in-registers aggregate return of a buffer +// resource struct. // RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -x hip -std=c++11 -fclangir \ -// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: -fcuda-is-device -fno-clangir-call-conv-lowering -emit-cir %s -o %t.cir // RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s // RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -x hip -std=c++11 -fclangir \ -// RUN: -fcuda-is-device -emit-llvm %s -o %t-cir.ll +// RUN: -fcuda-is-device -fno-clangir-call-conv-lowering -emit-llvm %s -o %t-cir.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s // RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -x hip -std=c++11 \ diff --git a/clang/test/CIR/CodeGenHIP/cleanup-alloca-addrspace.hip b/clang/test/CIR/CodeGenHIP/cleanup-alloca-addrspace.hip index c55cd55a98d675..7cbb8ff8301ec0 100644 --- a/clang/test/CIR/CodeGenHIP/cleanup-alloca-addrspace.hip +++ b/clang/test/CIR/CodeGenHIP/cleanup-alloca-addrspace.hip @@ -7,9 +7,10 @@ // correct address space after flatten cfg. This is important for address-space // aware targets like amdgpu. -// CIR-FLAT-LABEL: cir.func {{.*}} @_Z1fv +// S is not trivially destructible, so f returns it through a generic sret +// pointer rather than a __retval alloca. +// CIR-FLAT-LABEL: cir.func {{.*}} @_Z1fv(%arg0: !cir.ptr<!rec_S> {{.*}}llvm.sret = !rec_S // CIR-FLAT: cir.alloca "__cleanup_dest_slot" {{.*}} : !cir.ptr<!s32i, target_address_space(5)> -// CIR-FLAT: cir.alloca "__retval" {{.*}} : !cir.ptr<!rec_S, target_address_space(5)> // CIR-FLAT: cir.alloca "nrvo" {{.*}} : !cir.ptr<!cir.bool, target_address_space(5)> // CIR-FLAT-LABEL: cir.func {{.*}} @_ZN1SD1Ev @@ -18,9 +19,8 @@ // CIR-FLAT-LABEL: cir.func {{.*}} @_ZN1SD2Ev // CIR-FLAT: cir.alloca "this" {{.*}} : !cir.ptr<!cir.ptr<!rec_S>, target_address_space(5)> -// LLVM-LABEL: define {{.*}} @_Z1fv +// LLVM-LABEL: define {{.*}} void @_Z1fv(ptr {{.*}}sret(%struct.S) // LLVM: alloca i32, align 4, addrspace(5) -// LLVM: alloca %struct.S, align 1, addrspace(5) // LLVM: alloca i8, align 1, addrspace(5) // LLVM-LABEL: define {{.*}} @_ZN1SD1Ev diff --git a/clang/test/CIR/Transforms/abi-lowering/amdgpu-scalars.cir b/clang/test/CIR/Transforms/abi-lowering/amdgpu-scalars.cir new file mode 100644 index 00000000000000..5b9049e1197575 --- /dev/null +++ b/clang/test/CIR/Transforms/abi-lowering/amdgpu-scalars.cir @@ -0,0 +1,105 @@ +// RUN: cir-opt %s -cir-call-conv-lowering=target=amdgpu | FileCheck %s + +!s8i = !cir.int<s, 8> +!s16i = !cir.int<s, 16> +!u16i = !cir.int<u, 16> +!s32i = !cir.int<s, 32> +!s64i = !cir.int<s, 64> + +!rec_E0 = !cir.struct<"E0" {}> +!rec_Pair16 = !cir.struct<"Pair16" {data !s16i, data !s16i}> + +module attributes { + dlti.dl_spec = #dlti.dl_spec< + #dlti.dl_entry<i1, dense<8>: vector<2xi64>>, + #dlti.dl_entry<i8, dense<8>: vector<2xi64>>, + #dlti.dl_entry<i16, dense<16>: vector<2xi64>>, + #dlti.dl_entry<i32, dense<32>: vector<2xi64>>, + #dlti.dl_entry<i64, dense<64>: vector<2xi64>>> +} { + + // Register-sized integers are Direct: signature and call sites unchanged. + cir.func @passthrough(%arg0: !s32i, %arg1: !s64i) -> !s32i { + cir.return %arg0 : !s32i + } + + // CHECK: cir.func{{.*}} @passthrough(%arg0: !s32i, %arg1: !s64i) -> !s32i + // CHECK-NEXT: cir.return %arg0 : !s32i + + // Floating-point scalars are Direct. + cir.func @floats(%arg0: !cir.float, %arg1: !cir.double) -> !cir.double { + cir.return %arg1 : !cir.double + } + + // CHECK: cir.func{{.*}} @floats(%arg0: !cir.float, %arg1: !cir.double) -> !cir.double + + // Pointers are Direct. + cir.func @take_ptr(%arg0: !cir.ptr<!s32i>) -> !cir.ptr<!s32i> { + cir.return %arg0 : !cir.ptr<!s32i> + } + + // CHECK: cir.func{{.*}} @take_ptr(%arg0: !cir.ptr<!s32i>) -> !cir.ptr<!s32i> + + // Signed sub-register integer is sign-extended. + cir.func @take_s8(%arg0: !s8i) { + cir.return + } + + // CHECK: cir.func{{.*}} @take_s8(%arg0: !s8i {llvm.signext}) + + // Unsigned sub-register integer is zero-extended. + cir.func @take_u16(%arg0: !u16i) { + cir.return + } + + // CHECK: cir.func{{.*}} @take_u16(%arg0: !u16i {llvm.zeroext}) + + // bool is zero-extended. + cir.func @take_bool(%arg0: !cir.bool) { + cir.return + } + + // CHECK: cir.func{{.*}} @take_bool(%arg0: !cir.bool {llvm.zeroext}) + + // Call site picks up the same extension attribute on the operand. + cir.func @call_s8(%arg0: !s8i) { + cir.call @take_s8(%arg0) : (!s8i) -> () + cir.return + } + + // CHECK: cir.call @take_s8(%arg0) : (!s8i {llvm.signext}) -> () + + // A sub-register integer return is also extended, not just arguments. + cir.func @ret_s8() -> !s8i { + %0 = cir.const #cir.int<0> : !s8i + cir.return %0 : !s8i + } + + // CHECK: cir.func{{.*}} @ret_s8() -> (!s8i {llvm.signext}) + + // A zero-field record classifies as Ignore: the empty argument is dropped + // and the real argument shifts down to the first slot. + cir.func @take_empty(%arg0: !rec_E0, %arg1: !s32i) -> !s32i { + cir.return %arg1 : !s32i + } + + // CHECK: cir.func{{.*}} @take_empty(%arg0: !s32i) -> !s32i + // CHECK-NEXT: cir.return %arg0 : !s32i + + // A small aggregate (<= 32 bits) is packed into a single i32 register. + cir.func @takes_small(%arg0: !rec_Pair16) { + %0 = cir.alloca "p" align(4) : !cir.ptr<!rec_Pair16> + cir.store %arg0, %0 : !rec_Pair16, !cir.ptr<!rec_Pair16> + cir.return + } + + // CHECK: cir.func{{.*}} @takes_small(%{{.*}}: !u32i) + + // A kernel scalar argument stays Direct: the AMDGPU_KERNEL convention is + // routed to the kernel classifier and the signature is unchanged. + cir.func @kernel_scalar(%arg0: !s32i) cc(amdgpu_kernel) { + cir.return + } + + // CHECK: cir.func{{.*}} @kernel_scalar(%arg0: !s32i){{.*}}cc(amdgpu_kernel) +} diff --git a/clang/test/CIR/Transforms/abi-lowering/x86_64-lang-addrspace-nyi.cir b/clang/test/CIR/Transforms/abi-lowering/x86_64-lang-addrspace-nyi.cir index a8350b421757dc..285df67e53c838 100644 --- a/clang/test/CIR/Transforms/abi-lowering/x86_64-lang-addrspace-nyi.cir +++ b/clang/test/CIR/Transforms/abi-lowering/x86_64-lang-addrspace-nyi.cir @@ -15,4 +15,4 @@ module attributes { } -// CHECK: op x86_64 calling-convention lowering not yet implemented for type +// CHECK: op calling-convention lowering not yet implemented for type diff --git a/clang/test/CIR/Transforms/abi-lowering/x86_64-variadic-nyi.cir b/clang/test/CIR/Transforms/abi-lowering/x86_64-variadic-nyi.cir index 52a74f07fd2380..a2aab9c43318b8 100644 --- a/clang/test/CIR/Transforms/abi-lowering/x86_64-variadic-nyi.cir +++ b/clang/test/CIR/Transforms/abi-lowering/x86_64-variadic-nyi.cir @@ -44,7 +44,7 @@ module attributes { cir.func @indirect_nyi(%arg0: !cir.ptr<!cir.func<(!cir.ptr<!s8i>, ...) -> !s32i>>, %arg1: !cir.ptr<!s8i>, %arg2: !cir.ptr<!i96>) { %v = cir.load %arg2 : !cir.ptr<!i96>, !i96 - // CHECK: :[[@LINE+1]]:10: error: 'cir.call' op x86_64 calling-convention lowering not yet implemented for type '!cir.int<s, 96>' + // CHECK: :[[@LINE+1]]:10: error: 'cir.call' op calling-convention lowering not yet implemented for type '!cir.int<s, 96>' %0 = cir.call %arg0(%arg1, %v) : (!cir.ptr<!cir.func<(!cir.ptr<!s8i>, ...) -> !s32i>>, !cir.ptr<!s8i>, !i96) -> !s32i cir.return >From a240d40a17bc7d28752ba6088f05c13b297512ad Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Wed, 7 Oct 2026 19:34:21 +0530 Subject: [PATCH 2/2] test update --- .../call-conv-lowering-amdgpu-arg-types.c | 117 ++++++++++++++++++ .../call-conv-lowering-amdgpu-cxx-records.cpp | 55 ++++++++ .../call-conv-lowering-amdgpu-return-types.c | 105 ++++++++++++++++ .../call-conv-lowering-amdgpu-kernel-args.hip | 22 ++++ 4 files changed, 299 insertions(+) create mode 100644 clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-arg-types.c create mode 100644 clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-cxx-records.cpp create mode 100644 clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-return-types.c create mode 100644 clang/test/CIR/CodeGenHIP/call-conv-lowering-amdgpu-kernel-args.hip diff --git a/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-arg-types.c b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-arg-types.c new file mode 100644 index 00000000000000..6bddbb9003f756 --- /dev/null +++ b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-arg-types.c @@ -0,0 +1,117 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +typedef struct {} empty_t; +typedef struct { char c; } i8_wrapper_t; +typedef struct { int i; } i32_wrapper_t; +typedef struct { float f; } f32_wrapper_t; +typedef struct { char a, b; } char2_t; +typedef struct { char a, b, c; } small_t; +typedef short short4 __attribute__((ext_vector_type(4))); + +void arg_void(void) {} + +// CIR: cir.func {{.*}}@arg_void() +// LLVM: define {{.*}}void @arg_void() + +void arg_bool(_Bool b) {} + +// CIR: cir.func {{.*}}@arg_bool(%arg0: !cir.bool {{.*}}llvm.zeroext +// LLVM: define {{.*}}void @arg_bool(i1 noundef zeroext %{{.*}}) + +void arg_char(char c) {} + +// CIR: cir.func {{.*}}@arg_char(%arg0: !s8i {{.*}}llvm.signext +// LLVM: define {{.*}}void @arg_char(i8 noundef signext %{{.*}}) + +void arg_uchar(unsigned char c) {} + +// CIR: cir.func {{.*}}@arg_uchar(%arg0: !u8i {{.*}}llvm.zeroext +// LLVM: define {{.*}}void @arg_uchar(i8 noundef zeroext %{{.*}}) + +void arg_short(short s) {} + +// CIR: cir.func {{.*}}@arg_short(%arg0: !s16i {{.*}}llvm.signext +// LLVM: define {{.*}}void @arg_short(i16 noundef signext %{{.*}}) + +void arg_int(int i) {} + +// CIR: cir.func {{.*}}@arg_int(%arg0: !s32i {{.*}}) +// LLVM: define {{.*}}void @arg_int(i32 noundef %{{.*}}) + +void arg_long(long l) {} + +// CIR: cir.func {{.*}}@arg_long(%arg0: !s64i {{.*}}) +// LLVM: define {{.*}}void @arg_long(i64 noundef %{{.*}}) + +void arg_float(float f) {} + +// CIR: cir.func {{.*}}@arg_float(%arg0: !cir.float {{.*}}) +// LLVM: define {{.*}}void @arg_float(float noundef %{{.*}}) + +void arg_double(double d) {} + +// CIR: cir.func {{.*}}@arg_double(%arg0: !cir.double {{.*}}) +// LLVM: define {{.*}}void @arg_double(double noundef %{{.*}}) + +void arg_ptr(int *p) {} + +// CIR: cir.func {{.*}}@arg_ptr(%arg0: !cir.ptr<!s32i> {{.*}}) +// LLVM: define {{.*}}void @arg_ptr(ptr noundef %{{.*}}) + +void arg_bitint65(_BitInt(65) i) {} + +// CIR: cir.func {{.*}}@arg_bitint65(%arg0: !cir.int<s, 65, bitint> {{.*}}) +// LLVM: define {{.*}}void @arg_bitint65(i65 noundef %{{.*}}) + +void arg_short4(short4 v) {} + +// CIR: cir.func {{.*}}@arg_short4(%arg0: !cir.vector<4 x !s16i> {{.*}}) +// LLVM: define {{.*}}void @arg_short4(<4 x i16> noundef %{{.*}}) + +// An empty struct is dropped from the signature. +void arg_empty(empty_t e) {} + +// CIR: cir.func {{.*}}@arg_empty() +// LLVM: define {{.*}}void @arg_empty() + +// Single-element structs pass as their element. +void arg_i8_wrapper(i8_wrapper_t w) {} + +// CIR: cir.func {{.*}}@arg_i8_wrapper(%arg0: !s8i{{.*}}) +// CIR: cir.alloca "coerce" {{.*}} : !cir.ptr<{{.*}}, target_address_space(5)> +// LLVM: define {{.*}}void @arg_i8_wrapper(i8 %{{.*}}) + +void arg_i32_wrapper(i32_wrapper_t w) {} + +// CIR: cir.func {{.*}}@arg_i32_wrapper(%arg0: !s32i{{.*}}) +// LLVM: define {{.*}}void @arg_i32_wrapper(i32 %{{.*}}) + +void arg_f32_wrapper(f32_wrapper_t w) {} + +// CIR: cir.func {{.*}}@arg_f32_wrapper(%arg0: !cir.float{{.*}}) +// LLVM: define {{.*}}void @arg_f32_wrapper(float %{{.*}}) + +// Aggregates up to 32 bits are packed into one integer register. +void arg_char2(char2_t s) {} + +// CIR: cir.func {{.*}}@arg_char2(%arg0: !u16i{{.*}}) +// LLVM: define {{.*}}void @arg_char2(i16 %{{.*}}) + +void arg_small(small_t s) {} + +// CIR: cir.func {{.*}}@arg_small(%arg0: !u32i{{.*}}) +// LLVM: define {{.*}}void @arg_small(i32 %{{.*}}) + +// The call site is coerced the same way as the callee. +void call_small(small_t s) { arg_small(s); } + +// CIR: cir.func {{.*}}@call_small(%arg0: !u32i{{.*}}) +// CIR: cir.call @arg_small(%{{.*}}){{.*}}: (!u32i{{.*}}) -> () +// LLVM: define {{.*}}void @call_small(i32 %{{.*}}) +// LLVM: call void @arg_small(i32 %{{.*}}) diff --git a/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-cxx-records.cpp b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-cxx-records.cpp new file mode 100644 index 00000000000000..0f85621ebc371b --- /dev/null +++ b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-cxx-records.cpp @@ -0,0 +1,55 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +struct NonTrivialDtor { + int i; + ~NonTrivialDtor(); +}; + +struct NonCopyable { + int i; + NonCopyable(); + NonCopyable(const NonCopyable &) = delete; +}; + +struct TrivialBase { int i; }; +struct DerivedSingle : TrivialBase {}; +struct MemberPtrHolder { int TrivialBase::*p; }; + +// A record that cannot be passed in registers returns through a generic sret. +NonTrivialDtor ret_non_trivial_dtor() { return {1}; } + +// CIR: cir.func {{.*}}@_Z20ret_non_trivial_dtorv(%arg0: !cir.ptr<!rec_NonTrivialDtor> {{.*}}llvm.sret = !rec_NonTrivialDtor +// LLVM: define {{.*}}void @_Z20ret_non_trivial_dtorv(ptr dead_on_unwind noalias writable sret(%struct.NonTrivialDtor) align 4 %{{.*}}) + +NonCopyable ret_non_copyable() { return NonCopyable(); } + +// CIR: cir.func {{.*}}@_Z16ret_non_copyablev(%arg0: !cir.ptr<!rec_NonCopyable> {{.*}}llvm.sret = !rec_NonCopyable +// LLVM: define {{.*}}void @_Z16ret_non_copyablev(ptr dead_on_unwind noalias writable sret(%struct.NonCopyable) align 4 %{{.*}}) + +// A base subobject holding the only field still counts as a single element. +DerivedSingle ret_derived_single() { return {}; } + +// CIR: cir.func {{.*}}@_Z18ret_derived_singlev() -> !s32i +// LLVM: define {{.*}}i32 @_Z18ret_derived_singlev() + +void arg_derived_single(DerivedSingle d) {} + +// CIR: cir.func {{.*}}@_Z18arg_derived_single13DerivedSingle(%arg0: !s32i{{.*}}) +// LLVM: define {{.*}}void @_Z18arg_derived_single13DerivedSingle(i32 %{{.*}}) + +// A data member pointer is lowered to an i64 offset before classification. +MemberPtrHolder ret_member_ptr() { return {nullptr}; } + +// CIR: cir.func {{.*}}@_Z14ret_member_ptrv() -> !s64i +// LLVM: define {{.*}}i64 @_Z14ret_member_ptrv() + +void arg_member_ptr(MemberPtrHolder h) {} + +// CIR: cir.func {{.*}}@_Z14arg_member_ptr15MemberPtrHolder(%arg0: !s64i{{.*}}) +// LLVM: define {{.*}}void @_Z14arg_member_ptr15MemberPtrHolder(i64 %{{.*}}) diff --git a/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-return-types.c b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-return-types.c new file mode 100644 index 00000000000000..6c353f00f94fc4 --- /dev/null +++ b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-return-types.c @@ -0,0 +1,105 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fclangir -fclangir-call-conv-lowering -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +typedef struct {} empty_t; +typedef struct { char c; } i8_wrapper_t; +typedef struct { int i; } i32_wrapper_t; +typedef struct { float f; } f32_wrapper_t; +typedef struct { char a, b; } char2_t; +typedef struct { char a, b, c; } small_t; + +void ret_void(void) {} + +// CIR: cir.func {{.*}}@ret_void() +// LLVM: define {{.*}}void @ret_void() + +_Bool ret_bool(void) { return 0; } + +// CIR: cir.func {{.*}}@ret_bool() -> (!cir.bool {{.*}}llvm.zeroext +// LLVM: define {{.*}}zeroext i1 @ret_bool() + +char ret_char(void) { return 0; } + +// CIR: cir.func {{.*}}@ret_char() -> (!s8i {{.*}}llvm.signext +// LLVM: define {{.*}}signext i8 @ret_char() + +short ret_short(void) { return 0; } + +// CIR: cir.func {{.*}}@ret_short() -> (!s16i {{.*}}llvm.signext +// LLVM: define {{.*}}signext i16 @ret_short() + +int ret_int(void) { return 0; } + +// CIR: cir.func {{.*}}@ret_int() -> {{\(?}}!s32i +// LLVM: define {{.*}}i32 @ret_int() + +long ret_long(void) { return 0; } + +// CIR: cir.func {{.*}}@ret_long() -> {{\(?}}!s64i +// LLVM: define {{.*}}i64 @ret_long() + +float ret_float(void) { return 0; } + +// CIR: cir.func {{.*}}@ret_float() -> {{\(?}}!cir.float +// LLVM: define {{.*}}float @ret_float() + +_BitInt(65) ret_bitint65(void) { return 0; } + +// CIR: cir.func {{.*}}@ret_bitint65() -> {{\(?}}!cir.int<s, 65, bitint> +// LLVM: define {{.*}}i65 @ret_bitint65() + +// An empty struct return becomes void. +empty_t ret_empty(void) { + empty_t e; + return e; +} + +// CIR: cir.func {{.*}}@ret_empty() +// LLVM: define {{.*}}void @ret_empty() + +// Single-element structs return as their element. +i8_wrapper_t ret_i8_wrapper(void) { + i8_wrapper_t w = {1}; + return w; +} + +// CIR: cir.func {{.*}}@ret_i8_wrapper() -> !s8i +// LLVM: define {{.*}}i8 @ret_i8_wrapper() + +i32_wrapper_t ret_i32_wrapper(void) { + i32_wrapper_t w = {1}; + return w; +} + +// CIR: cir.func {{.*}}@ret_i32_wrapper() -> !s32i +// LLVM: define {{.*}}i32 @ret_i32_wrapper() + +f32_wrapper_t ret_f32_wrapper(void) { + f32_wrapper_t w = {1.0f}; + return w; +} + +// CIR: cir.func {{.*}}@ret_f32_wrapper() -> !cir.float +// LLVM: define {{.*}}float @ret_f32_wrapper() + +// Aggregates up to 32 bits are packed into one integer register. +char2_t ret_char2(void) { + char2_t s = {1, 2}; + return s; +} + +// CIR: cir.func {{.*}}@ret_char2() -> !u16i +// LLVM: define {{.*}}i16 @ret_char2() + +small_t ret_small(void) { + small_t s = {1, 2, 3}; + return s; +} + +// CIR: cir.func {{.*}}@ret_small() -> !u32i +// LLVM: define {{.*}}i32 @ret_small() diff --git a/clang/test/CIR/CodeGenHIP/call-conv-lowering-amdgpu-kernel-args.hip b/clang/test/CIR/CodeGenHIP/call-conv-lowering-amdgpu-kernel-args.hip new file mode 100644 index 00000000000000..4422e7329644b1 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/call-conv-lowering-amdgpu-kernel-args.hip @@ -0,0 +1,22 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fclangir -fclangir-call-conv-lowering -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fclangir -fclangir-call-conv-lowering -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +#define __global__ __attribute__((global)) + +typedef struct { int i; } i32_wrapper_t; + +__global__ void kern_scalar(int i, float f) {} + +// CIR: cir.func {{.*}}@_Z11kern_scalarif(%arg0: !s32i {{.*}}, %arg1: !cir.float {{.*}}){{.*}}cc(amdgpu_kernel) +// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_scalarif(i32 noundef %{{.*}}, float noundef %{{.*}}) + +// A single-element struct kernel argument is passed as its element. +__global__ void kern_wrapper(i32_wrapper_t w) {} + +// CIR: cir.func {{.*}}@_Z12kern_wrapper13i32_wrapper_t(%arg0: !s32i{{.*}}){{.*}}cc(amdgpu_kernel) +// LLVM: define {{.*}}amdgpu_kernel void @_Z12kern_wrapper13i32_wrapper_t(i32 %{{.*}}) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
