kerbowa created this revision.
kerbowa added reviewers: yaxunl, hliao.
Herald added subscribers: cfe-commits, nhaehnle, jvesely.
Herald added a project: clang.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D79213
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
@@ -2259,6 +2259,18 @@
return CGF.Builder.CreateFPCast(value, varType, "arg.unpromote");
}
+/// Returns true if the argument is a generic HIP pointer that was coerced to a
+/// global pointer.
+bool isCoercedHIPGlobalPointer(CodeGenFunction &CGF,
+ const LangOptions &LangOpts,
+ const ABIArgInfo &ArgI, const QualType &Ty) {
+ return LangOpts.HIP && isa<llvm::PointerType>(ArgI.getCoerceToType()) &&
+ ArgI.getCoerceToType()->getPointerAddressSpace() == 1 &&
+ CGF.ConvertType(Ty)->getPointerAddressSpace() == 0 &&
+ ArgI.getCoerceToType()->getPointerElementType() ==
+ CGF.ConvertType(Ty)->getPointerElementType();
+}
+
/// Returns the attribute (either parameter attribute, or function
/// attribute), which declares argument ArgNo to be non-null.
static const NonNullAttr *getNonNullAttr(const Decl *FD, const ParmVarDecl
*PVD,
@@ -2541,6 +2553,14 @@
// Pointer to store into.
Address Ptr = emitAddressAtOffset(*this, Alloca, ArgI);
+ // Restrict qualified HIP pointers that were coerced to global pointers
+ // can be marked with the noalias attribute.
+ if (isCoercedHIPGlobalPointer(*this, getLangOpts(), ArgI, Ty) &&
+ Arg->getType().isRestrictQualified()) {
+ auto AI = cast<llvm::Argument>(FnArgs[FirstIRArg]);
+ AI->addAttr(llvm::Attribute::NoAlias);
+ }
+
// Fast-isel and the optimizer generally like scalar values better than
// FCAs, so we flatten them if this is safe to do for this argument.
llvm::StructType *STy =
dyn_cast<llvm::StructType>(ArgI.getCoerceToType());
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
@@ -2259,6 +2259,18 @@
return CGF.Builder.CreateFPCast(value, varType, "arg.unpromote");
}
+/// Returns true if the argument is a generic HIP pointer that was coerced to a
+/// global pointer.
+bool isCoercedHIPGlobalPointer(CodeGenFunction &CGF,
+ const LangOptions &LangOpts,
+ const ABIArgInfo &ArgI, const QualType &Ty) {
+ return LangOpts.HIP && isa<llvm::PointerType>(ArgI.getCoerceToType()) &&
+ ArgI.getCoerceToType()->getPointerAddressSpace() == 1 &&
+ CGF.ConvertType(Ty)->getPointerAddressSpace() == 0 &&
+ ArgI.getCoerceToType()->getPointerElementType() ==
+ CGF.ConvertType(Ty)->getPointerElementType();
+}
+
/// Returns the attribute (either parameter attribute, or function
/// attribute), which declares argument ArgNo to be non-null.
static const NonNullAttr *getNonNullAttr(const Decl *FD, const ParmVarDecl *PVD,
@@ -2541,6 +2553,14 @@
// Pointer to store into.
Address Ptr = emitAddressAtOffset(*this, Alloca, ArgI);
+ // Restrict qualified HIP pointers that were coerced to global pointers
+ // can be marked with the noalias attribute.
+ if (isCoercedHIPGlobalPointer(*this, getLangOpts(), ArgI, Ty) &&
+ Arg->getType().isRestrictQualified()) {
+ auto AI = cast<llvm::Argument>(FnArgs[FirstIRArg]);
+ AI->addAttr(llvm::Attribute::NoAlias);
+ }
+
// Fast-isel and the optimizer generally like scalar values better than
// FCAs, so we flatten them if this is safe to do for this argument.
llvm::StructType *STy = dyn_cast<llvm::StructType>(ArgI.getCoerceToType());
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
