diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -2425,15 +2425,18 @@ case ABIArgInfo::Extend: case ABIArgInfo::Direct: { - - // If we have the trivial case, handle it with no muss and fuss. - if (!isa(ArgI.getCoerceToType()) && - ArgI.getCoerceToType() == ConvertType(Ty) && - ArgI.getDirectOffset() == 0) { + auto AI = Fn->getArg(FirstIRArg); + llvm::Type *LTy = ConvertType(Arg->getType()); + + // Prepare parameter attributes. So far, only attributes for pointer + // parameters are prepared. See + // http://llvm.org/docs/LangRef.html#paramattrs. + if (ArgI.getDirectOffset() == 0 && LTy->isPointerTy() && + ArgI.getCoerceToType()->isPointerTy()) { assert(NumIRArgs == 1); - auto AI = Fn->getArg(FirstIRArg); if (const ParmVarDecl *PVD = dyn_cast(Arg)) { + // Set `nonnull` attribute if any. if (getNonNullAttr(CurCodeDecl, PVD, PVD->getType(), PVD->getFunctionScopeIndex()) && !CGM.getCodeGenOpts().NullPointerIsValid) @@ -2471,6 +2474,7 @@ AI->addAttr(llvm::Attribute::NonNull); } + // Set `align` attribute if any. const auto *AVAttr = PVD->getAttr(); if (!AVAttr) if (const auto *TOTy = dyn_cast(OTy)) @@ -2488,8 +2492,17 @@ } } + // Set 'noalias' if an argument type has the `restrict` qualifier. if (Arg->getType().isRestrictQualified()) AI->addAttr(llvm::Attribute::NoAlias); + } + + // Prepare the argument value. If we have the trivial case, handle it + // with no muss and fuss. + if (!isa(ArgI.getCoerceToType()) && + ArgI.getCoerceToType() == ConvertType(Ty) && + ArgI.getDirectOffset() == 0) { + assert(NumIRArgs == 1); // LLVM expects swifterror parameters to be used in very restricted // ways. Copy the value into a less-restricted temporary. diff --git a/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu b/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu --- a/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu +++ b/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu @@ -67,3 +67,10 @@ t.x[0][0] += 1.f; t.x[1][0] += 2.f; } + +// Check that coerced pointers retain the noalias attribute when qualified with __restrict. +// CHECK: define amdgpu_kernel void @_Z7kernel7Pi(i32 addrspace(1)* noalias %x.coerce) +// HOST: define void @_Z22__device_stub__kernel7Pi(i32* noalias %x) +__global__ void kernel7(int *__restrict x) { + x[0]++; +}