r355421 - [CUDA][HIP][Sema] Fix template kernel with function as template parameter

[Yaxun Liu via cfe-commits]

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 at -1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
+// expected-error at -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 at -1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
+// expected-error at -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 at -1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
 
 template <typename T> __host__ __device__ void hd3() { host_fn(); }
-// expected-error at -1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
+// expected-error at -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.