Author: yaxunl Date: Tue Mar 5 10:19:35 2019 New Revision: 355421 URL: http://llvm.org/viewvc/llvm-project?rev=355421&view=rev Log: [CUDA][HIP][Sema] Fix template kernel with function as template parameter
If a kernel template has a function as its template parameter, a device function should be allowed as template argument since a kernel can call a device function. However, currently if the kernel template is instantiated in a host function, clang will emit an error message saying the device function is an invalid candidate for the template parameter. This happens because clang checks the reference to the device function during parsing the template arguments. At this point, the template is not instantiated yet. Clang incorrectly assumes the device function is called by the host function and emits the error message. This patch fixes the issue by disabling checking of device function during parsing template arguments and deferring the check to the instantion of the template. At that point, the template decl is already available, therefore the check can be done against the instantiated function template decl. Differential Revision: https://reviews.llvm.org/D56411 Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp cfe/trunk/lib/Sema/SemaExpr.cpp cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=355421&r1=355420&r2=355421&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaCUDA.cpp (original) +++ cfe/trunk/lib/Sema/SemaCUDA.cpp Tue Mar 5 10:19:35 2019 @@ -675,6 +675,11 @@ Sema::DeviceDiagBuilder Sema::CUDADiagIf bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) { assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); assert(Callee && "Callee may not be null."); + + auto &ExprEvalCtx = ExprEvalContexts.back(); + if (ExprEvalCtx.isUnevaluated() || ExprEvalCtx.isConstantEvaluated()) + return true; + // FIXME: Is bailing out early correct here? Should we instead assume that // the caller is a global initializer? FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext); Modified: cfe/trunk/lib/Sema/SemaExpr.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=355421&r1=355420&r2=355421&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaExpr.cpp (original) +++ cfe/trunk/lib/Sema/SemaExpr.cpp Tue Mar 5 10:19:35 2019 @@ -14799,6 +14799,9 @@ void Sema::MarkFunctionReferenced(Source if (FPT && isUnresolvedExceptionSpec(FPT->getExceptionSpecType())) ResolveExceptionSpec(Loc, FPT); + if (getLangOpts().CUDA) + CheckCUDACall(Loc, Func); + // If we don't need to mark the function as used, and we don't need to // try to provide a definition, there's nothing more to do. if ((Func->isUsed(/*CheckUsedAttr=*/false) || !OdrUse) && Modified: cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu?rev=355421&r1=355420&r2=355421&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu (original) +++ cfe/trunk/test/SemaCUDA/call-device-fn-from-host.cu Tue Mar 5 10:19:35 2019 @@ -37,7 +37,7 @@ __host__ __device__ void T::hd3() { } template <typename T> __host__ __device__ void hd2() { device_fn(); } -// expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}} +// expected-error@-1 2 {{reference to __device__ function 'device_fn' in __host__ __device__ function}} void host_fn() { hd2<int>(); } __host__ __device__ void hd() { device_fn(); } @@ -90,3 +90,8 @@ __host__ __device__ void fn_ptr_template static __host__ __device__ void hd_func() { device_fn(); } __global__ void kernel() { hd_func(); } void host_func(void) { kernel<<<1, 1>>>(); } + +// Should allow host function call kernel template with device function argument. +__device__ void f(); +template<void(*F)()> __global__ void t() { F(); } +__host__ void g() { t<f><<<1,1>>>(); } Modified: cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu?rev=355421&r1=355420&r2=355421&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu (original) +++ cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu Tue Mar 5 10:19:35 2019 @@ -56,14 +56,14 @@ __host__ __device__ void T::hd3() { } template <typename T> __host__ __device__ void hd2() { host_fn(); } -// expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} +// expected-error@-1 2 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} __global__ void kernel() { hd2<int>(); } __host__ __device__ void hd() { host_fn(); } // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} template <typename T> __host__ __device__ void hd3() { host_fn(); } -// expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} +// expected-error@-1 2 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} __device__ void device_fn() { hd3<int>(); } // No error because this is never instantiated. _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
