hliao created this revision.
hliao added reviewers: rjmccall, kerbowa, yaxunl.
Herald added subscribers: cfe-commits, nhaehnle, jvesely.
Herald added a project: clang.
hliao updated this revision to Diff 262011.
hliao added a comment.
hliao added a parent revision: D79394: [clang][codegen] Refactor argument
loading in function prolog. NFC..
Add dependency.
- 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.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D79395
Files:
clang/lib/CodeGen/CGCall.cpp
clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
Index: clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
===================================================================
--- clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
+++ 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]++;
+}
Index: clang/lib/CodeGen/CGCall.cpp
===================================================================
--- clang/lib/CodeGen/CGCall.cpp
+++ 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<llvm::StructType>(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<ParmVarDecl>(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<AlignValueAttr>();
if (!AVAttr)
if (const auto *TOTy = dyn_cast<TypedefType>(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<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.
Index: clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
===================================================================
--- clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
+++ 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]++;
+}
Index: clang/lib/CodeGen/CGCall.cpp
===================================================================
--- clang/lib/CodeGen/CGCall.cpp
+++ 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<llvm::StructType>(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<ParmVarDecl>(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<AlignValueAttr>();
if (!AVAttr)
if (const auto *TOTy = dyn_cast<TypedefType>(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<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.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
