[clang] [CIR] Wire AMDGPU into the call-convention lowering pass (PR #220197)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 16 04:06:49 PDT 2026
https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/220197
>From 45e1e3c5e49a8e4263b0bb35e6b489076be2bb75 Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Tue, 1 Sep 2026 14:31:15 +0530
Subject: [PATCH 1/3] [CIR] Wire AMDGPU into the call-convention lowering pass
---
clang/include/clang/CIR/Dialect/Passes.h | 2 +-
clang/include/clang/CIR/Dialect/Passes.td | 11 +-
.../Transforms/CallConvLoweringPass.cpp | 203 +++++++++++-------
.../TargetLowering/CIRABIRewriteContext.cpp | 13 +-
clang/lib/CIR/Lowering/CIRPasses.cpp | 6 +-
.../call-conv-lowering-amdgpu-arg-types.c | 165 ++++++++++++++
.../call-conv-lowering-amdgpu-cxx-records.cpp | 74 +++++++
.../call-conv-lowering-amdgpu-return-types.c | 158 ++++++++++++++
.../CodeGenHIP/amdgcn-buffer-rsrc-type.hip | 8 +-
.../call-conv-lowering-amdgpu-kernel-args.hip | 140 ++++++++++++
.../CodeGenHIP/cleanup-alloca-addrspace.hip | 8 +-
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 6 +-
.../abi-lowering/amdgpu-calling-conv.cir | 59 +++++
.../abi-lowering/amdgpu-kernel-hip-ptr.cir | 53 +++++
.../x86_64-lang-addrspace-nyi.cir | 2 +-
.../abi-lowering/x86_64-variadic-nyi.cir | 2 +-
16 files changed, 820 insertions(+), 90 deletions(-)
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
create mode 100644 clang/test/CIR/Transforms/abi-lowering/amdgpu-calling-conv.cir
create mode 100644 clang/test/CIR/Transforms/abi-lowering/amdgpu-kernel-hip-ptr.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..e51d119cb0d6a8 100644
--- a/clang/include/clang/CIR/Dialect/Passes.td
+++ b/clang/include/clang/CIR/Dialect/Passes.td
@@ -225,9 +225,10 @@ def CallConvLowering : Pass<"cir-call-conv-lowering", "mlir::ModuleOp"> {
Two driver modes select how each function's classification is computed:
- - `target=<name>` selects an ABI target. Currently only `"test"` (the
- MLIR test target in `mlir/lib/ABI/Targets/Test/`) is supported. Real
- targets (x86_64, AArch64, ...) will be added once the LLVM ABI library
+ - `target=<name>` selects an ABI target: `"x86_64"` (System V),
+ `"amdgpu"`, or `"test"` (the MLIR test target in
+ `mlir/lib/ABI/Targets/Test/`). The clang pipeline picks the target from
+ the module triple. Other targets will be added as the LLVM ABI library
ships them.
- `classification-attr=<name>` reads a `DictionaryAttr` named `<name>`
from each `cir.func` and parses it via the test-target injection
@@ -246,7 +247,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 aaacedeff1e74d..728f8881a4064f 100644
--- a/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp
@@ -309,8 +309,11 @@ static mlir::Type abiTypeToCIR(const llvm::abi::Type *ty, MLIRContext *ctx) {
.Case([&](const llvm::abi::FloatType *fltTy) {
return cir::getFloatingPointType(*fltTy->getSemantics(), ctx);
})
- .Case([&](const llvm::abi::PointerType *) {
- return cir::PointerType::get(cir::VoidType::get(ctx));
+ .Case([&](const llvm::abi::PointerType *ptrTy) {
+ mlir::ptr::MemorySpaceAttrInterface addrSpace;
+ if (unsigned as = ptrTy->getAddrSpace())
+ addrSpace = cir::TargetAddressSpaceAttr::get(ctx, as);
+ return cir::PointerType::get(cir::VoidType::get(ctx), addrSpace);
})
.Case([&](const llvm::abi::VectorType *vecTy) -> mlir::Type {
mlir::Type elemCIR = abiTypeToCIR(vecTy->getElementType(), ctx);
@@ -545,7 +548,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");
});
}
@@ -553,11 +556,12 @@ static const llvm::abi::Type *mapCIRType(mlir::Type type,
/// CIRABIRewriteContext.
///
/// Direct: the value passes in register(s). A coercion is forwarded in the
-/// four cases where the value has to be rebuilt on the wire: an aggregate
+/// five cases where the value has to be rebuilt on the wire: an aggregate
/// unpacked into the register(s) holding it, a scalar too wide for one register
/// split into a tuple of them, a scalar the classifier widens to fill its
-/// eightbyte, and a value whose live bytes start partway into its storage
-/// because a leading eightbyte holds no field (getDirectOffset). canFlatten
+/// eightbyte, a pointer moved into another address space, and a value whose
+/// live bytes start partway into its storage because a leading eightbyte
+/// holds no field (getDirectOffset). canFlatten
/// follows the classifier's CanBeFlattened, so the rewriter splits a
/// multi-field coerced struct into individual wire arguments unless the
/// classifier asked to keep it intact. Any other scalar passes in its
@@ -594,6 +598,26 @@ convertABIArgInfo(const llvm::abi::ArgInfo &info, MLIRContext *ctx,
// than emit a read of the wrong bytes.
if (offset && coerceIsRegisterTuple)
return std::nullopt;
+ // A pointer moved to another address space keeps its pointee, which
+ // abiTypeToCIR cannot recover.
+ if (auto origPtr = dyn_cast_if_present<cir::PointerType>(origTy)) {
+ if (const auto *coercePtr =
+ dyn_cast_if_present<llvm::abi::PointerType>(coerceAbi)) {
+ auto origAS = dyn_cast_if_present<cir::TargetAddressSpaceAttr>(
+ origPtr.getAddrSpace());
+ unsigned origASValue = origAS ? origAS.getValue() : 0;
+ if (!offset && coercePtr->getAddrSpace() != origASValue) {
+ auto coercedPtr =
+ cast<cir::PointerType>(abiTypeToCIR(coerceAbi, ctx));
+ ArgClassification classified = ArgClassification::getDirect(
+ cir::PointerType::get(origPtr.getPointee(),
+ coercedPtr.getAddrSpace()),
+ /*offset=*/0);
+ classified.canFlatten = info.getCanBeFlattened();
+ return classified;
+ }
+ }
+ }
// Compare widths rather than identity: a coerce no wider than the natural
// type carries the same value and needs no rewrite.
auto origInt = dyn_cast_if_present<cir::IntType>(origTy);
@@ -658,20 +682,30 @@ 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 calling convention to the llvm::CallingConv the ABI classifier
+/// uses. Conventions the classifier does not distinguish fall back to C.
+static llvm::CallingConv::ID convertCallingConv(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
+/// point as passed through an ellipsis. \p callConv is the calling convention
+/// the signature is classified under. 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) {
+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()) &&
@@ -681,9 +715,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))
@@ -699,15 +732,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;
};
@@ -762,19 +795,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,
+ convertCallingConv(func.getCallingConv()), modOp,
+ [&]() { return func.emitOpError(); });
}
/// Classify the single type fetched by a `cir.va_arg` as an unnamed argument
@@ -783,11 +816,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];
@@ -800,17 +833,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
@@ -969,10 +1004,25 @@ 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) {
+ // HIP passes a generic pointer kernel argument as a global one.
+ bool isHIP = false;
+ if (auto langOpts = moduleOp->getAttrOfType<cir::LoweringLangOptionsAttr>(
+ cir::CIRDialect::getLoweringLangOptionsAttrName()))
+ isHIP = langOpts.getHip();
+ amdgpuTarget = llvm::abi::createAMDGPUTargetInfo(
+ abiTypeMapper->getTypeBuilder(),
+ /*CoerceGenericPtrArgToGlobal=*/isHIP);
+ }
auto x86TargetFor =
[&](llvm::abi::X86AVXABILevel level) -> const llvm::abi::TargetInfo & {
assert(static_cast<unsigned>(level) < numAvxLevels &&
@@ -981,7 +1031,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;
};
@@ -991,6 +1041,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
@@ -1008,9 +1065,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) {
@@ -1047,11 +1104,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
@@ -1068,9 +1126,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),
+ convertCallingConv(callee.getCallingConv()), moduleOp);
if (!fc) {
anyFailed = true;
return;
@@ -1187,31 +1245,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;
@@ -1232,8 +1290,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)),
+ convertCallingConv(enclosing.getCallingConv()), moduleOp,
[&]() { return v->emitOpError(); });
if (!ac) {
signalPassFailure();
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp
index c7ca6151029eb4..04324b4f91f191 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp
@@ -486,12 +486,23 @@ mlir::Value emitCoercionToMemory(mlir::OpBuilder &builder, mlir::Location loc,
/// Coerce \p src to type \p dstTy by going through memory and load the whole
/// coerced value back out. Builds on emitCoercionToMemory, adding the final
-/// load of the destination-typed view.
+/// load of the destination-typed view. A pointer that only changes address
+/// space is cast instead.
mlir::Value emitCoercion(mlir::OpBuilder &builder, mlir::Location loc,
mlir::Type dstTy, mlir::Value src,
mlir::Block *slotBlock, const mlir::DataLayout &dl,
SmallPtrSetImpl<mlir::Operation *> &createdOps,
unsigned offset) {
+ auto srcPtrTy = dyn_cast<cir::PointerType>(src.getType());
+ auto dstPtrTy = dyn_cast<cir::PointerType>(dstTy);
+ if (!offset && srcPtrTy && dstPtrTy &&
+ srcPtrTy.getPointee() == dstPtrTy.getPointee()) {
+ auto cast = cir::CastOp::create(builder, loc, dstTy,
+ cir::CastKind::address_space, src);
+ createdOps.insert(cast);
+ return cast;
+ }
+
mlir::Value dstSlot = emitCoercionToMemory(builder, loc, dstTy, src,
slotBlock, dl, createdOps, offset);
auto load = cir::LoadOp::create(builder, loc, dstSlot);
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/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..ef65e90a528166
--- /dev/null
+++ b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-arg-types.c
@@ -0,0 +1,165 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu-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 amdgpu-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 amdgpu-amd-amdhsa -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// Checks that CallConvLowering classifies arguments of ordinary AMDGPU
+// functions, which use the default convention rather than the kernel one.
+
+// TODO(cir): Add the argument cases still NYI in CallConvLowering:
+// - Aggregates of 33 to 64 bits, which classic coerces to [2 x i32].
+// abiTypeToCIR has no array case yet.
+// - Aggregates over 64 bits that fit the 16-register budget. The classifier
+// returns Direct with no coerce type, which the bridge rejects.
+// - Aggregates past the register budget, passed byref in addrspace(5).
+// - Structs with a flexible array member and transparent unions. The bridge
+// does not set those record flags yet.
+// - Arguments passed through an ellipsis.
+
+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 { int *p; } ptr_wrapper_t;
+typedef struct { char a, b; } char2_t;
+typedef struct { char a, b, c; } small_t;
+typedef union { int i; float f; } int_float_t;
+typedef short short4 __attribute__((ext_vector_type(4)));
+typedef float float3 __attribute__((ext_vector_type(3)));
+
+void arg_void(void) {}
+
+// CIR: cir.func {{.*}}@arg_void()
+// LLVM: define {{.*}}void @arg_void()
+
+// Sub-word integers and bool carry signext or zeroext per their signedness.
+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 %{{.*}})
+
+// Word-sized and wider scalars pass through unchanged.
+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_half(_Float16 h) {}
+
+// CIR: cir.func {{.*}}@arg_half(%arg0: !cir.f16 {{.*}})
+// LLVM: define {{.*}}void @arg_half(half noundef %{{.*}})
+
+// Outside a HIP kernel a generic pointer stays generic.
+void arg_ptr(int *p) {}
+
+// CIR: cir.func {{.*}}@arg_ptr(%arg0: !cir.ptr<!s32i> {{.*}})
+// LLVM: define {{.*}}void @arg_ptr(ptr noundef %{{.*}})
+
+// A _BitInt wider than 64 bits and vectors, including 16-bit and 3-element
+// ones, pass whole.
+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 %{{.*}})
+
+void arg_float3(float3 v) {}
+
+// CIR: cir.func {{.*}}@arg_float3(%arg0: !cir.vector<3 x !cir.float> {{.*}})
+// LLVM: define {{.*}}void @arg_float3(<3 x float> 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. The callee rebuilds the
+// struct in a private (addrspace 5) slot.
+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 %{{.*}})
+
+void arg_ptr_wrapper(ptr_wrapper_t w) {}
+
+// CIR: cir.func {{.*}}@arg_ptr_wrapper(%arg0: !cir.ptr<{{[^,]*}}>{{.*}})
+// LLVM: define {{.*}}void @arg_ptr_wrapper(ptr %{{.*}})
+
+// 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 %{{.*}})
+
+void arg_union(int_float_t u) {}
+
+// CIR: cir.func {{.*}}@arg_union(%arg0: !u32i{{.*}})
+// LLVM: define {{.*}}void @arg_union(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 %{{.*}})
+
+// The call site carries the same extension as the callee.
+void call_char(char c) { arg_char(c); }
+
+// LLVM: define {{.*}}void @call_char(i8 noundef signext %{{.*}})
+// LLVM: call void @arg_char(i8 noundef signext %{{.*}})
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..e8342304b890c3
--- /dev/null
+++ b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-cxx-records.cpp
@@ -0,0 +1,74 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu-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 amdgpu-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 amdgpu-amd-amdhsa -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// Checks that CallConvLowering classifies C++ records on ordinary AMDGPU
+// functions, where C++ rules on triviality and layout change the result.
+
+// TODO(cir): Add the C++ record cases still NYI in CallConvLowering:
+// - Non-trivial records as arguments. Classic passes them indirectly as
+// ptr addrspace(5), but the bridge drops the indirect address space.
+// - Member function pointers, lowered to a 128-bit pair that the bridge
+// cannot coerce yet.
+// - Records of 33 to 64 bits, such as a base plus a field, which classic
+// coerces to [2 x i32].
+
+struct Empty {};
+
+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 %{{.*}})
+
+// An empty C++ record is one byte in size but is still dropped.
+void arg_empty(Empty e) {}
+
+// CIR: cir.func {{.*}}@_Z9arg_empty5Empty()
+// LLVM: define {{.*}}void @_Z9arg_empty5Empty()
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..0bab2eb455fe52
--- /dev/null
+++ b/clang/test/CIR/CodeGen/call-conv-lowering-amdgpu-return-types.c
@@ -0,0 +1,158 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu-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 amdgpu-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 amdgpu-amd-amdhsa -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// Checks that CallConvLowering classifies return values of ordinary AMDGPU
+// functions, which use the default convention rather than the kernel one.
+
+// TODO(cir): Add the return cases still NYI in CallConvLowering:
+// - Aggregates of 33 to 64 bits, which classic coerces to [2 x i32].
+// abiTypeToCIR has no array case yet.
+// - Aggregates over 64 bits that fit the 16-register budget. The classifier
+// returns Direct with no coerce type, which the bridge rejects.
+// - Larger aggregates, returned through sret in addrspace(5). The bridge
+// drops the indirect address space.
+// - Structs with a flexible array member. The bridge does not set that
+// record flag yet.
+// - 3-element vector returns, blocked on CIRGen's vec3 load and store NYI.
+
+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 { int *p; } ptr_wrapper_t;
+typedef struct { char a, b; } char2_t;
+typedef struct { char a, b, c; } small_t;
+typedef union { int i; float f; } int_float_t;
+
+void ret_void(void) {}
+
+// CIR: cir.func {{.*}}@ret_void()
+// LLVM: define {{.*}}void @ret_void()
+
+// Sub-word integers and bool carry signext or zeroext per their signedness.
+_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()
+
+unsigned char ret_uchar(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_uchar() -> (!u8i {{.*}}llvm.zeroext
+// LLVM: define {{.*}}zeroext i8 @ret_uchar()
+
+short ret_short(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_short() -> (!s16i {{.*}}llvm.signext
+// LLVM: define {{.*}}signext i16 @ret_short()
+
+// Word-sized and wider scalars and pointers return unchanged.
+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()
+
+double ret_double(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_double() -> {{\(?}}!cir.double
+// LLVM: define {{.*}}double @ret_double()
+
+_Float16 ret_half(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_half() -> {{\(?}}!cir.f16
+// LLVM: define {{.*}}half @ret_half()
+
+int *ret_ptr(void) { return 0; }
+
+// CIR: cir.func {{.*}}@ret_ptr() -> {{\(?}}!cir.ptr<!s32i>
+// LLVM: define {{.*}}ptr @ret_ptr()
+_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()
+
+ptr_wrapper_t ret_ptr_wrapper(void) {
+ ptr_wrapper_t w = {0};
+ return w;
+}
+
+// CIR: cir.func {{.*}}@ret_ptr_wrapper() -> !cir.ptr<{{[^>]*}}>
+// LLVM: define {{.*}}ptr @ret_ptr_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()
+
+int_float_t ret_union(void) {
+ int_float_t u = {1};
+ return u;
+}
+
+// CIR: cir.func {{.*}}@ret_union() -> !u32i
+// LLVM: define {{.*}}i32 @ret_union()
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/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..d5ddd8480ffd30
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/call-conv-lowering-amdgpu-kernel-args.hip
@@ -0,0 +1,140 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu-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 amdgpu-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 amdgpu-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// Checks that CallConvLowering classifies AMDGPU kernel arguments under the
+// kernel convention.
+
+// TODO(cir): Add the kernel argument cases still pending:
+// - Multi-field aggregates, passed byref in addrspace(4), once CallConvLowering
+// supports IndirectAliased.
+// - A multi-field struct holding a generic pointer. Classic codegen coerces
+// the pointer to global, the LLVM ABI library does not yet.
+
+#define __global__ __attribute__((global))
+#define __device__ __attribute__((device))
+
+typedef struct { int i; } i32_wrapper_t;
+typedef struct { i32_wrapper_t w; } nested_t;
+typedef struct { float f[1]; } f32_array1_t;
+typedef struct { float *x; } ptr_wrapper_t;
+typedef float float3 __attribute__((ext_vector_type(3)));
+typedef float float4 __attribute__((ext_vector_type(4)));
+typedef short short4 __attribute__((ext_vector_type(4)));
+
+// Word-sized scalars pass through unchanged.
+__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 %{{.*}})
+
+__global__ void kern_wide(long l, double d) {}
+
+// CIR: cir.func {{.*}}@_Z9kern_wideld(%arg0: !s64i {{.*}}, %arg1: !cir.double {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z9kern_wideld(i64 noundef %{{.*}}, double noundef %{{.*}})
+
+__global__ void kern_half(_Float16 h) {}
+
+// CIR: cir.func {{.*}}@_Z9kern_halfDF16_(%arg0: !cir.f16 {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z9kern_halfDF16_(half noundef %{{.*}})
+
+__global__ void kern_bitint(_BitInt(65) b) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_bitintDB65_(%arg0: !cir.int<s, 65, bitint> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_bitintDB65_(i65 noundef %{{.*}})
+
+// Vectors, including 16-bit and 3-element ones, pass whole.
+__global__ void kern_float4(float4 v) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_float4Dv4_f(%arg0: !cir.vector<4 x !cir.float> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_float4Dv4_f(<4 x float> noundef %{{.*}})
+
+__global__ void kern_short4(short4 v) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_short4Dv4_s(%arg0: !cir.vector<4 x !s16i> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_short4Dv4_s(<4 x i16> noundef %{{.*}})
+
+__global__ void kern_float3(float3 v) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_float3Dv3_f(%arg0: !cir.vector<3 x !cir.float> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_float3Dv3_f(<3 x float> noundef %{{.*}})
+
+// A pointer already in a non-generic address space keeps it.
+__global__ void kern_global_ptr(__attribute__((address_space(1))) int *p) {}
+
+// CIR: cir.func {{.*}}@_Z15kern_global_ptrPU3AS1i(%arg0: !cir.ptr<!s32i, target_address_space(1)> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z15kern_global_ptrPU3AS1i(ptr addrspace(1) noundef %{{.*}})
+
+// HIP passes a generic pointer or reference as a global pointer and casts it
+// back to generic in the kernel.
+__global__ void kern_ptr(int *p) {}
+
+// CIR: cir.func {{.*}}@_Z8kern_ptrPi(%[[P:[^:]+]]: !cir.ptr<!s32i, target_address_space(1)> {{.*}}){{.*}}cc(amdgpu_kernel)
+// CIR: cir.cast address_space %[[P]] : !cir.ptr<!s32i, target_address_space(1)> -> !cir.ptr<!s32i>
+// LLVM: define {{.*}}amdgpu_kernel void @_Z8kern_ptrPi(ptr addrspace(1) noundef %{{.*}})
+
+__global__ void kern_ref(int &r) {}
+
+// CIR: cir.func {{.*}}@_Z8kern_refRi(%{{.*}}: !cir.ptr<!s32i, target_address_space(1)> {{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z8kern_refRi(ptr addrspace(1) noundef nonnull align 4 dereferenceable(4) %{{.*}})
+
+// A struct holding only a pointer is unwrapped first, so it is coerced too.
+__global__ void kern_ptr_wrap(ptr_wrapper_t w) {}
+
+// CIR: cir.func {{.*}}@_Z13kern_ptr_wrap13ptr_wrapper_t(%{{.*}}: !cir.ptr<{{.*}}, target_address_space(1)>{{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z13kern_ptr_wrap13ptr_wrapper_t(ptr addrspace(1) %{{.*}})
+
+// A device function keeps the generic pointer.
+__device__ void dev_ptr(int *p) {}
+
+// CIR: cir.func {{.*}}@_Z7dev_ptrPi(%{{.*}}: !cir.ptr<!s32i> {{.*}})
+// LLVM: define {{.*}}void @_Z7dev_ptrPi(ptr 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 %{{.*}})
+
+// The element is found through nested structs and one-element arrays.
+__global__ void kern_nested(nested_t n) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_nested8nested_t(%arg0: !s32i{{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_nested8nested_t(i32 %{{.*}})
+
+__global__ void kern_array1(f32_array1_t a) {}
+
+// CIR: cir.func {{.*}}@_Z11kern_array112f32_array1_t(%arg0: !cir.float{{.*}}){{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z11kern_array112f32_array1_t(float %{{.*}})
+
+// Kernel arguments are read from the kernarg segment, so sub-word integers
+// and bool are not extended, unlike in a device function.
+__global__ void kern_char(char c) {}
+
+// CIR: cir.func {{.*}}@_Z9kern_charc(%arg0: !s8i {llvm.noundef}{{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z9kern_charc(i8 noundef %{{.*}})
+
+__global__ void kern_uchar(unsigned char c) {}
+
+// CIR: cir.func {{.*}}@_Z10kern_ucharh(%arg0: !u8i {llvm.noundef}{{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z10kern_ucharh(i8 noundef %{{.*}})
+
+__global__ void kern_short(short s) {}
+
+// CIR: cir.func {{.*}}@_Z10kern_shorts(%arg0: !s16i {llvm.noundef}{{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z10kern_shorts(i16 noundef %{{.*}})
+
+__global__ void kern_bool(bool b) {}
+
+// CIR: cir.func {{.*}}@_Z9kern_boolb(%arg0: !cir.bool {llvm.noundef}{{.*}}cc(amdgpu_kernel)
+// LLVM: define {{.*}}amdgpu_kernel void @_Z9kern_boolb(i1 noundef %{{.*}})
+
+// The same argument to a device function is extended.
+__device__ void dev_char(char c) {}
+
+// CIR: cir.func {{.*}}@_Z8dev_charc(%arg0: !s8i {{.*}}llvm.signext
+// LLVM: define {{.*}}void @_Z8dev_charc(i8 noundef signext %{{.*}})
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/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CIR/CodeGenSYCL/kernel-caller-entry-point.cpp
index cf9a923d466144..4b0d591bde61b1 100644
--- a/clang/test/CIR/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CIR/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -17,9 +17,11 @@
// On AMDGPU and NVPTX, the kernel caller entry point uses the target's device
// kernel calling convention.
-// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple amdgpu-amd-amdhsa -fclangir -emit-cir %s -o %t-amdgcn.cir
+// TODO(cir): drop -fno-clangir-call-conv-lowering once CallConvLowering
+// supports byref (IndirectAliased) kernel arguments on AMDGPU.
+// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple amdgpu-amd-amdhsa -fclangir -fno-clangir-call-conv-lowering -emit-cir %s -o %t-amdgcn.cir
// RUN: FileCheck --input-file=%t-amdgcn.cir %s -check-prefix=CIR-AMDGCN
-// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple amdgpu-amd-amdhsa -fclangir -emit-llvm %s -o %t-amdgcn-cir.ll
+// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple amdgpu-amd-amdhsa -fclangir -fno-clangir-call-conv-lowering -emit-llvm %s -o %t-amdgcn-cir.ll
// RUN: FileCheck --input-file=%t-amdgcn-cir.ll %s -check-prefix=LLVM-OGCG-AMDGCN
// RUN: %clang_cc1 -std=c++20 -fsycl-is-device -triple amdgpu-amd-amdhsa -emit-llvm %s -o %t-amdgcn.ll
// RUN: FileCheck --input-file=%t-amdgcn.ll %s -check-prefix=LLVM-OGCG-AMDGCN
diff --git a/clang/test/CIR/Transforms/abi-lowering/amdgpu-calling-conv.cir b/clang/test/CIR/Transforms/abi-lowering/amdgpu-calling-conv.cir
new file mode 100644
index 00000000000000..4043621d054bca
--- /dev/null
+++ b/clang/test/CIR/Transforms/abi-lowering/amdgpu-calling-conv.cir
@@ -0,0 +1,59 @@
+// RUN: cir-opt %s -cir-call-conv-lowering=target=amdgpu | FileCheck %s
+
+// Checks that CallConvLowering under target=amdgpu classifies a function by
+// its calling convention. Default functions get extension and Ignore rewrites
+// applied to definitions and call sites. Kernels go to the kernel classifier.
+
+!s8i = !cir.int<s, 8>
+!s32i = !cir.int<s, 32>
+
+!rec_E0 = !cir.struct<"E0" {}>
+
+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>>>
+} {
+
+ // Signed sub-register integer is sign-extended.
+ cir.func @take_s8(%arg0: !s8i) {
+ cir.return
+ }
+
+ // CHECK: cir.func{{.*}} @take_s8(%arg0: !s8i {llvm.signext})
+
+ // 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 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 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)
+
+ // Without HIP in the module's lang options, a kernel pointer stays generic.
+ cir.func @kernel_ptr(%arg0: !cir.ptr<!s32i>) cc(amdgpu_kernel) {
+ cir.return
+ }
+
+ // CHECK: cir.func{{.*}} @kernel_ptr(%arg0: !cir.ptr<!s32i>){{.*}}cc(amdgpu_kernel)
+}
diff --git a/clang/test/CIR/Transforms/abi-lowering/amdgpu-kernel-hip-ptr.cir b/clang/test/CIR/Transforms/abi-lowering/amdgpu-kernel-hip-ptr.cir
new file mode 100644
index 00000000000000..ba35c4bc74b475
--- /dev/null
+++ b/clang/test/CIR/Transforms/abi-lowering/amdgpu-kernel-hip-ptr.cir
@@ -0,0 +1,53 @@
+// RUN: cir-opt %s -cir-call-conv-lowering=target=amdgpu | FileCheck %s
+
+// Under HIP a generic pointer kernel argument is passed as a global pointer.
+// The HIP flag comes from the module's lowering_lang_options.
+
+!s32i = !cir.int<s, 32>
+!rec_PtrWrap = !cir.struct<"PtrWrap" {data !cir.ptr<!s32i>}>
+
+module attributes {
+ dlti.dl_spec = #dlti.dl_spec<
+ #dlti.dl_entry<i32, dense<32>: vector<2xi64>>,
+ #dlti.dl_entry<i64, dense<64>: vector<2xi64>>>,
+ cir.lowering_lang_options = #cir.lowering_lang_options<
+ exceptions = false, threadsafe_statics = true, cuda = true,
+ cuda_is_device = true, hip = true, hlsl = false, gpu_rdc = false,
+ openmp = false, openmp_is_target_device = false, clang_abi_compat = 17>
+} {
+
+ // The kernel casts the incoming global pointer back to generic before the
+ // body's spill reads it.
+ cir.func @kernel_ptr(%arg0: !cir.ptr<!s32i>) cc(amdgpu_kernel) {
+ %0 = cir.alloca "p" align(8) init : !cir.ptr<!cir.ptr<!s32i>>
+ cir.store %arg0, %0 : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
+ cir.return
+ }
+
+ // CHECK: cir.func{{.*}} @kernel_ptr(%[[ARG:[^:]+]]: !cir.ptr<!s32i, target_address_space(1)>){{.*}}cc(amdgpu_kernel)
+ // CHECK-NEXT: %[[GEN:.*]] = cir.cast address_space %[[ARG]] : !cir.ptr<!s32i, target_address_space(1)> -> !cir.ptr<!s32i>
+ // CHECK: cir.store %[[GEN]], %{{.*}} : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
+
+ // A pointer already outside the generic address space is left alone.
+ cir.func @kernel_global_ptr(%arg0: !cir.ptr<!s32i, target_address_space(1)>)
+ cc(amdgpu_kernel) {
+ cir.return
+ }
+
+ // CHECK: cir.func{{.*}} @kernel_global_ptr(%arg0: !cir.ptr<!s32i, target_address_space(1)>){{.*}}cc(amdgpu_kernel)
+ // CHECK-NOT: cir.cast
+
+ // A struct holding only a pointer is unwrapped to it, then coerced.
+ cir.func @kernel_ptr_wrap(%arg0: !rec_PtrWrap) cc(amdgpu_kernel) {
+ cir.return
+ }
+
+ // CHECK: cir.func{{.*}} @kernel_ptr_wrap(%{{.*}}: !cir.ptr<{{.*}}, target_address_space(1)>){{.*}}cc(amdgpu_kernel)
+
+ // Only kernels are coerced.
+ cir.func @device_ptr(%arg0: !cir.ptr<!s32i>) {
+ cir.return
+ }
+
+ // CHECK: cir.func{{.*}} @device_ptr(%arg0: !cir.ptr<!s32i>)
+}
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 9deae33e9246b2ebb963828a35fdb5a650e0a330 Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Thu, 15 Oct 2026 16:22:27 +0530
Subject: [PATCH 2/3] rebase and fixes per feedback
---
clang/include/clang/CIR/Dialect/Passes.td | 11 +++---
.../Transforms/CallConvLoweringPass.cpp | 38 +++++++++----------
.../TargetLowering/CIRABIRewriteContext.cpp | 38 ++++++++++++++-----
3 files changed, 53 insertions(+), 34 deletions(-)
diff --git a/clang/include/clang/CIR/Dialect/Passes.td b/clang/include/clang/CIR/Dialect/Passes.td
index e51d119cb0d6a8..6cc5733d3d9acc 100644
--- a/clang/include/clang/CIR/Dialect/Passes.td
+++ b/clang/include/clang/CIR/Dialect/Passes.td
@@ -225,11 +225,12 @@ def CallConvLowering : Pass<"cir-call-conv-lowering", "mlir::ModuleOp"> {
Two driver modes select how each function's classification is computed:
- - `target=<name>` selects an ABI target: `"x86_64"` (System V),
- `"amdgpu"`, or `"test"` (the MLIR test target in
- `mlir/lib/ABI/Targets/Test/`). The clang pipeline picks the target from
- the module triple. Other targets will be added as the LLVM ABI library
- ships them.
+ - `target=<name>` selects an ABI target. The clang pipeline picks the
+ target from the module triple. Supported targets:
+ - `"x86_64"`: System V.
+ - `"amdgpu"`: AMDGPU.
+ - `"test"`: the MLIR test target in `mlir/lib/ABI/Targets/Test/`.
+ Other targets will be added as the LLVM ABI library ships them.
- `classification-attr=<name>` reads a `DictionaryAttr` named `<name>`
from each `cir.func` and parses it via the test-target injection
helper. Used by tests to inject arbitrary classifications without
diff --git a/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp b/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp
index 728f8881a4064f..b59111b3c40c4c 100644
--- a/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp
@@ -556,9 +556,9 @@ static const llvm::abi::Type *mapCIRType(mlir::Type type,
/// CIRABIRewriteContext.
///
/// Direct: the value passes in register(s). A coercion is forwarded in the
-/// five cases where the value has to be rebuilt on the wire: an aggregate
-/// unpacked into the register(s) holding it, a scalar too wide for one register
-/// split into a tuple of them, a scalar the classifier widens to fill its
+/// cases where the value has to be rebuilt on the wire: an aggregate unpacked
+/// into the register(s) holding it, a scalar too wide for one register split
+/// into a tuple of them, a scalar the classifier widens to fill its
/// eightbyte, a pointer moved into another address space, and a value whose
/// live bytes start partway into its storage because a leading eightbyte
/// holds no field (getDirectOffset). canFlatten
@@ -600,22 +600,22 @@ convertABIArgInfo(const llvm::abi::ArgInfo &info, MLIRContext *ctx,
return std::nullopt;
// A pointer moved to another address space keeps its pointee, which
// abiTypeToCIR cannot recover.
- if (auto origPtr = dyn_cast_if_present<cir::PointerType>(origTy)) {
- if (const auto *coercePtr =
- dyn_cast_if_present<llvm::abi::PointerType>(coerceAbi)) {
- auto origAS = dyn_cast_if_present<cir::TargetAddressSpaceAttr>(
- origPtr.getAddrSpace());
- unsigned origASValue = origAS ? origAS.getValue() : 0;
- if (!offset && coercePtr->getAddrSpace() != origASValue) {
- auto coercedPtr =
- cast<cir::PointerType>(abiTypeToCIR(coerceAbi, ctx));
- ArgClassification classified = ArgClassification::getDirect(
- cir::PointerType::get(origPtr.getPointee(),
- coercedPtr.getAddrSpace()),
- /*offset=*/0);
- classified.canFlatten = info.getCanBeFlattened();
- return classified;
- }
+ auto origPtr = dyn_cast_if_present<cir::PointerType>(origTy);
+ const auto *coercePtr =
+ dyn_cast_if_present<llvm::abi::PointerType>(coerceAbi);
+ if (origPtr && coercePtr && !offset) {
+ // Compare numbers, since an explicit address_space(0) keeps a
+ // target_address_space(0) attribute that a plain pointer lacks.
+ auto origAS = dyn_cast_if_present<cir::TargetAddressSpaceAttr>(
+ origPtr.getAddrSpace());
+ if (coercePtr->getAddrSpace() != (origAS ? origAS.getValue() : 0)) {
+ auto coercedPtr = cast<cir::PointerType>(abiTypeToCIR(coerceAbi, ctx));
+ ArgClassification classified = ArgClassification::getDirect(
+ cir::PointerType::get(origPtr.getPointee(),
+ coercedPtr.getAddrSpace()),
+ /*offset=*/0);
+ classified.canFlatten = info.getCanBeFlattened();
+ return classified;
}
}
// Compare widths rather than identity: a coerce no wider than the natural
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp
index 04324b4f91f191..10863c56052bcf 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp
@@ -484,6 +484,24 @@ mlir::Value emitCoercionToMemory(mlir::OpBuilder &builder, mlir::Location loc,
dstTy == slotTy ? 0 : offset);
}
+/// Cast \p src to \p dstTy when both are pointers to the same pointee, so only
+/// the address space changes. Returns null when a real coercion is needed.
+static mlir::Value
+emitAddrSpaceCoercion(mlir::OpBuilder &builder, mlir::Location loc,
+ mlir::Type dstTy, mlir::Value src,
+ SmallPtrSetImpl<mlir::Operation *> &createdOps,
+ unsigned offset) {
+ auto srcPtrTy = dyn_cast<cir::PointerType>(src.getType());
+ auto dstPtrTy = dyn_cast<cir::PointerType>(dstTy);
+ if (offset || !srcPtrTy || !dstPtrTy ||
+ srcPtrTy.getPointee() != dstPtrTy.getPointee())
+ return {};
+ auto cast = cir::CastOp::create(builder, loc, dstTy,
+ cir::CastKind::address_space, src);
+ createdOps.insert(cast);
+ return cast;
+}
+
/// Coerce \p src to type \p dstTy by going through memory and load the whole
/// coerced value back out. Builds on emitCoercionToMemory, adding the final
/// load of the destination-typed view. A pointer that only changes address
@@ -493,15 +511,9 @@ mlir::Value emitCoercion(mlir::OpBuilder &builder, mlir::Location loc,
mlir::Block *slotBlock, const mlir::DataLayout &dl,
SmallPtrSetImpl<mlir::Operation *> &createdOps,
unsigned offset) {
- auto srcPtrTy = dyn_cast<cir::PointerType>(src.getType());
- auto dstPtrTy = dyn_cast<cir::PointerType>(dstTy);
- if (!offset && srcPtrTy && dstPtrTy &&
- srcPtrTy.getPointee() == dstPtrTy.getPointee()) {
- auto cast = cir::CastOp::create(builder, loc, dstTy,
- cir::CastKind::address_space, src);
- createdOps.insert(cast);
+ if (mlir::Value cast =
+ emitAddrSpaceCoercion(builder, loc, dstTy, src, createdOps, offset))
return cast;
- }
mlir::Value dstSlot = emitCoercionToMemory(builder, loc, dstTy, src,
slotBlock, dl, createdOps, offset);
@@ -805,14 +817,20 @@ replaceUsesFromSlot(mlir::OpBuilder &builder, mlir::Location loc,
/// \p origTy through a coercion slot at the start of \p slotBlock, and point
/// the uses of \p val outside \p createdOps at the result with
/// replaceUsesFromSlot. A direct \p offset lands on the coerced side, so the
-/// view of \p origTy starts at the slot. The coercion's operations are added
-/// to \p createdOps.
+/// view of \p origTy starts at the slot. A pointer that only changes address
+/// space is cast instead. The coercion's operations are added to
+/// \p createdOps.
static mlir::LogicalResult
emitCoercionAndReplaceUses(mlir::OpBuilder &builder, mlir::Location loc,
mlir::Type origTy, mlir::Value src, mlir::Value val,
mlir::Block *slotBlock, const mlir::DataLayout &dl,
SmallPtrSetImpl<mlir::Operation *> &createdOps,
unsigned offset) {
+ if (mlir::Value cast = emitAddrSpaceCoercion(builder, loc, origTy, src,
+ createdOps, offset)) {
+ val.replaceAllUsesExcept(cast, createdOps);
+ return mlir::success();
+ }
mlir::Value view = emitCoercionToMemory(builder, loc, origTy, src, slotBlock,
dl, createdOps, offset);
return replaceUsesFromSlot(builder, loc, val, view, dl, createdOps);
>From a8d086487c7a5deea96a4631a729f37ab5a33b70 Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Fri, 16 Oct 2026 16:35:39 +0530
Subject: [PATCH 3/3] update targetInfoFor
---
.../Transforms/CallConvLoweringPass.cpp | 36 ++++++++++++-------
.../TargetLowering/CIRABIRewriteContext.cpp | 2 ++
2 files changed, 25 insertions(+), 13 deletions(-)
diff --git a/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp b/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp
index b59111b3c40c4c..2b2890d9c85e1d 100644
--- a/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp
@@ -684,6 +684,8 @@ static llvm::abi::RequiredArgs requiredArgs(cir::FuncType fnTy) {
/// Map a CIR calling convention to the llvm::CallingConv the ABI classifier
/// uses. Conventions the classifier does not distinguish fall back to C.
+/// FIXME: Only the AMDGPU kernel convention is mapped. Every other convention
+/// is classified as C.
static llvm::CallingConv::ID convertCallingConv(cir::CallingConv cc) {
if (cc == cir::CallingConv::AMDGPUKernel)
return llvm::CallingConv::AMDGPU_KERNEL;
@@ -1013,16 +1015,6 @@ void CallConvLoweringPass::runOnOperation() {
std::unique_ptr<llvm::abi::TargetInfo> amdgpuTarget;
if (useAbiClassifier)
abiTypeMapper.emplace(dl);
- if (isAMDGPU) {
- // HIP passes a generic pointer kernel argument as a global one.
- bool isHIP = false;
- if (auto langOpts = moduleOp->getAttrOfType<cir::LoweringLangOptionsAttr>(
- cir::CIRDialect::getLoweringLangOptionsAttrName()))
- isHIP = langOpts.getHip();
- amdgpuTarget = llvm::abi::createAMDGPUTargetInfo(
- abiTypeMapper->getTypeBuilder(),
- /*CoerceGenericPtrArgToGlobal=*/isHIP);
- }
auto x86TargetFor =
[&](llvm::abi::X86AVXABILevel level) -> const llvm::abi::TargetInfo & {
assert(static_cast<unsigned>(level) < numAvxLevels &&
@@ -1041,12 +1033,30 @@ void CallConvLoweringPass::runOnOperation() {
return baseAvxLevel;
return funcAvxLevel(func, baseAvxLevel);
};
+ auto amdgpuTargetFor = [&]() -> const llvm::abi::TargetInfo & {
+ if (!amdgpuTarget) {
+ // HIP passes a generic pointer kernel argument as a global one.
+ auto langOpts = moduleOp->getAttrOfType<cir::LoweringLangOptionsAttr>(
+ cir::CIRDialect::getLoweringLangOptionsAttrName());
+ amdgpuTarget = llvm::abi::createAMDGPUTargetInfo(
+ abiTypeMapper->getTypeBuilder(),
+ /*CoerceGenericPtrArgToGlobal=*/langOpts && langOpts.getHip());
+ }
+ return *amdgpuTarget;
+ };
// 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));
+ switch (target) {
+ case cir::CallConvTarget::X86_64:
+ return x86TargetFor(avxLevelFor(func));
+ case cir::CallConvTarget::AMDGPU:
+ return amdgpuTargetFor();
+ case cir::CallConvTarget::None:
+ case cir::CallConvTarget::Test:
+ break;
+ }
+ llvm_unreachable("targetInfoFor needs an LLVM ABI library target");
};
// Classify every cir.func up front. No IR mutation happens here, so
diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp
index 10863c56052bcf..46e06b814776f3 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp
@@ -496,6 +496,8 @@ emitAddrSpaceCoercion(mlir::OpBuilder &builder, mlir::Location loc,
if (offset || !srcPtrTy || !dstPtrTy ||
srcPtrTy.getPointee() != dstPtrTy.getPointee())
return {};
+ assert(srcPtrTy.getAddrSpace() != dstPtrTy.getAddrSpace() &&
+ "pointers to the same pointee must differ in address space");
auto cast = cir::CastOp::create(builder, loc, dstTy,
cir::CastKind::address_space, src);
createdOps.insert(cast);
More information about the cfe-commits
mailing list