https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/220197
>From ec3d17efe802456adb382ecb70fd8c8738dcfc99 Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Tue, 1 Sep 2026 14:31:15 +0530 Subject: [PATCH] [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 | 205 +++++++++++------- .../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, 821 insertions(+), 91 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 4c543191f036d8..08795417c65f64 100644 --- a/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp +++ b/clang/lib/CIR/Dialect/Transforms/CallConvLoweringPass.cpp @@ -311,8 +311,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); @@ -338,7 +341,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 +536,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"); }); } @@ -541,11 +544,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 @@ -582,6 +586,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); @@ -641,20 +665,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()) && @@ -664,9 +698,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 +715,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 +778,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 @@ -766,11 +799,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 +816,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 +987,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 && @@ -964,7 +1014,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 +1024,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 +1048,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 +1087,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 +1109,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; @@ -1170,31 +1228,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 +1273,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 1f69ae89d47665..b994c006f6b375 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp @@ -442,12 +442,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..9d2341273e150e --- /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, 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
