maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[clang][codegen] Hoist parameter attribute setting in function prolog. 

Summary:
- If the coerced type is still a pointer, it should be set with proper
  parameter attributes, such as `noalias`, `nonnull`, and etc. Hoist
  that (pointer) parameter attribute setting so that the coerced pointer
  parameter could be marked properly.

Depends on D79394

Reviewers: rjmccall, kerbowa, yaxunl

Subscribers: jvesely, nhaehnle, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D79395
This commit is contained in:
Michael Liao 2020-05-05 00:55:13 -04:00
parent 276c8dde0b
commit 9142c0b46b
2 changed files with 25 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

23
clang/lib/CodeGen/CGCall.cpp
View File

@ -2425,15 +2425,18 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
case ABIArgInfo::Extend:
case ABIArgInfo::Direct: {
auto AI = Fn->getArg(FirstIRArg);
llvm::Type *LTy = ConvertType(Arg->getType());
// If we have the trivial case, handle it with no muss and fuss.
if (!isa<llvm::StructType>(ArgI.getCoerceToType()) &&
ArgI.getCoerceToType() == ConvertType(Ty) &&
ArgI.getDirectOffset() == 0) {
// 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<ParmVarDecl>(Arg)) {
// Set `nonnull` attribute if any.
if (getNonNullAttr(CurCodeDecl, PVD, PVD->getType(),
PVD->getFunctionScopeIndex()) &&
!CGM.getCodeGenOpts().NullPointerIsValid)
@ -2471,6 +2474,7 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
AI->addAttr(llvm::Attribute::NonNull);
}
// Set `align` attribute if any.
const auto *AVAttr = PVD->getAttr<AlignValueAttr>();
if (!AVAttr)
if (const auto *TOTy = dyn_cast<TypedefType>(OTy))
@ -2488,8 +2492,17 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
}
}
// 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<llvm::StructType>(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.

7
clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
View File

@ -67,3 +67,10 @@ __global__ void kernel6(struct T t) {
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]++;
}