[PATCH] D45223: [CUDA] Fix overloading resolution failure due to kernel calling convention
Yaxun Liu via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Tue Apr 3 12:35:05 PDT 2018
yaxunl created this revision.
yaxunl added reviewers: rjmccall, tra.
The following test causes overloading resolution failure in clang due to
missing kernel calling convention in the function pointer type.
template <typename T>
__global__ void EmptyKernel(void) {}
struct Dummy {
/// Type definition of the EmptyKernel kernel entry point
typedef void (*EmptyKernelPtr)();
EmptyKernelPtr Empty() { return EmptyKernel<void>; }
};
This happens before DRE is created. The fix is to drop the kernel
calling convention when converting function types.
https://reviews.llvm.org/D45223
Files:
lib/Sema/SemaOverload.cpp
test/CodeGenCUDA/kernel-amdgcn.cu
Index: test/CodeGenCUDA/kernel-amdgcn.cu
===================================================================
--- test/CodeGenCUDA/kernel-amdgcn.cu
+++ test/CodeGenCUDA/kernel-amdgcn.cu
@@ -15,15 +15,27 @@
non_kernel();
}
+// CHECK: define amdgpu_kernel void @_Z11EmptyKernelIvEvv
+template <typename T>
+__global__ void EmptyKernel(void) {}
+
+struct Dummy {
+ /// Type definition of the EmptyKernel kernel entry point
+ typedef void (*EmptyKernelPtr)();
+ EmptyKernelPtr Empty() { return EmptyKernel<void>; }
+};
+
// CHECK: define amdgpu_kernel void @_Z15template_kernelI1AEvT_
template<class T>
__global__ void template_kernel(T x) {}
void launch(void *f);
int main() {
+ Dummy D;
launch((void*)A::kernel);
launch((void*)kernel);
launch((void*)template_kernel<A>);
+ launch((void*)D.Empty());
return 0;
}
Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -1471,16 +1471,26 @@
Changed = true;
}
- // Drop 'noexcept' if not present in target type.
if (const auto *FromFPT = dyn_cast<FunctionProtoType>(FromFn)) {
const auto *ToFPT = cast<FunctionProtoType>(ToFn);
+
+ // Drop 'noexcept' if not present in target type.
if (FromFPT->isNothrow(Context) && !ToFPT->isNothrow(Context)) {
FromFn = cast<FunctionType>(
Context.getFunctionTypeWithExceptionSpec(QualType(FromFPT, 0),
EST_None)
.getTypePtr());
Changed = true;
}
+
+ // Drop cuda_kernel calling convention since it is invisible in AST.
+ if (FromFPT->getCallConv() == CC_CUDAKernel &&
+ FromFPT->getCallConv() != ToFPT->getCallConv()) {
+ FromFn = Context.adjustFunctionType(
+ FromFn, FromEInfo.withCallingConv(ToFPT->getCallConv()));
+ Changed = true;
+ }
+
// Convert FromFPT's ExtParameterInfo if necessary. The conversion is valid
// only if the ExtParameterInfo lists of the two function prototypes can be
// merged and the merged list is identical to ToFPT's ExtParameterInfo list.
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D45223.140844.patch
Type: text/x-patch
Size: 2171 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20180403/483a7015/attachment.bin>
More information about the cfe-commits
mailing list