[CUDA] Only allow __global__ on free functions and static member functions.

[jlebar]

Summary

Warn for NVCC compatibility if you declare a static member function or
inline function as global.

[tra]

include/clang/Basic/DiagnosticSemaKinds.td

6418 ↗

(On Diff #45080)

There's an Extension<> tablegen class used to report various C and C++ extensions.
This looks like a good use case for it.

6421 ↗

(On Diff #45080)

Perhaps we should use the same message as nvcc here:
'inline qualifier ignored for "global" function'

[jlebar]

Address tra's review.

[jlebar]

Thank you for the review.

include/clang/Basic/DiagnosticSemaKinds.td

6418 ↗

(On Diff #45080)

I'll also check that the rest of the CudaCompat warnings here shouldn't be Extension<>s.

6421 ↗

(On Diff #45080)

Ah, yes, given that it's not an nvcc error, we should just warn. (Or even not say anything at all; it seems like a silly warning.)

[jlebar]

Checked Thrust as we discussed. It doesn't seem to be getting hung up here.

[tra]

Small diags not. LGTM otherwise.

lib/Sema/SemaDeclAttr.cpp

3620–3629 ↗

(On Diff #45297)

Perhaps we should emit diagnostics on device side only.

[jlebar]

lib/Sema/SemaDeclAttr.cpp

3620–3629 ↗

(On Diff #45297)

Hm. I don't feel strongly, but since we're talking about global functions, this seems sort of relevant for both host and device compilation?

[tra]

lib/Sema/SemaDeclAttr.cpp

3620–3629 ↗

(On Diff #45297)

Host-only works, too. Printing this on both host and device is a bit too much.

cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td

 def err_nontemporal_builtin_must_be_pointer_intfltptr_or_vector : Error<
   "address argument to nontemporal builtin must be a pointer to integer, float, "
   "pointer, or a vector of such types (%0 invalid)">;
 
 def err_deleted_function_use : Error<"attempt to use a deleted function">;
 
 def err_kern_type_not_void_return : Error<
   "kernel function type %0 must have void return type">;
+def err_kern_is_nonstatic_method : Error<
+  "kernel function %0 must be a free function or static member function">;
 def err_config_scalar_return : Error<
   "CUDA special function 'cudaConfigureCall' must have scalar return type">;
 def err_kern_call_not_global_function : Error<
   "kernel call to non-global function %0">;
 def err_global_call_not_config : Error<
   "call to global function %0 not configured">;
 def err_ref_bad_target : Error<
   "reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
   "function %1 in %select{__device__|__global__|__host__|__host__ __device__}2 function">;
 def warn_host_calls_from_host_device : Warning<
   "calling __host__ function %0 from __host__ __device__ function %1 can lead to runtime errors">,
   InGroup<CudaCompat>;
+def warn_kern_is_method : Extension<
+  "kernel function %0 is a member function; this may not be accepted by nvcc">,
+  InGroup<CudaCompat>;
+def warn_kern_is_inline : Warning<
+  "ignored 'inline' attribute on kernel function %0">,
+  InGroup<CudaCompat>;
 
 def warn_non_pod_vararg_with_format_string : Warning<
   "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
   "%select{function|block|method|constructor}2; expected type from format "
   "string was %3">, InGroup<NonPODVarargs>, DefaultError;
 // The arguments to this diagnostic should match the warning above.
 def err_cannot_pass_objc_interface_to_vararg_format : Error<
   "cannot pass object with interface type %1 by value to variadic "

cfe/trunk/lib/Sema/SemaCUDA.cpp

 /// This function figures out if the given targets \param Target1 and
 /// \param Target2 conflict, and if they do not it fills in
 /// \param ResolvedTarget with a target that resolves for both calls.
 /// \return true if there's a conflict, false otherwise.
 static bool
 resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1,
                                 Sema::CUDAFunctionTarget Target2,
                                 Sema::CUDAFunctionTarget *ResolvedTarget) {
-  if (Target1 == Sema::CFT_Global && Target2 == Sema::CFT_Global) {
-    // TODO: this shouldn't happen, really. Methods cannot be marked __global__.
-    // Clang should detect this earlier and produce an error. Then this
-    // condition can be changed to an assertion.
-    return true;
-  }
+  // Only free functions and static member functions may be global.
+  assert(Target1 != Sema::CFT_Global);
+  assert(Target2 != Sema::CFT_Global);
 
   if (Target1 == Sema::CFT_HostDevice) {
     *ResolvedTarget = Target2;
   } else if (Target2 == Sema::CFT_HostDevice) {
     *ResolvedTarget = Target1;
   } else if (Target1 != Target2) {
     return true;
   } else {

cfe/trunk/lib/Sema/SemaDeclAttr.cpp

static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) {
   if (!FD->getReturnType()->isVoidType()) {
     SourceRange RTRange = FD->getReturnTypeSourceRange();
     S.Diag(FD->getTypeSpecStartLoc(), diag::err_kern_type_not_void_return)
         << FD->getType()
         << (RTRange.isValid() ? FixItHint::CreateReplacement(RTRange, "void")
                               : FixItHint());
     return;
   }
+  if (const auto *Method = dyn_cast<CXXMethodDecl>(FD)) {
+    if (Method->isInstance()) {
+      S.Diag(Method->getLocStart(), diag::err_kern_is_nonstatic_method)
+          << Method;
+      return;
+    }
+    S.Diag(Method->getLocStart(), diag::warn_kern_is_method) << Method;
+  }
+  // Only warn for "inline" when compiling for host, to cut down on noise.
+  if (FD->isInlineSpecified() && !S.getLangOpts().CUDAIsDevice)
+    S.Diag(FD->getLocStart(), diag::warn_kern_is_inline) << FD;
 
   D->addAttr(::new (S.Context)
               CUDAGlobalAttr(Attr.getRange(), S.Context,
                              Attr.getAttributeSpellingListIndex()));
-
 }
 
 static void handleGNUInlineAttr(Sema &S, Decl *D, const AttributeList &Attr) {
   FunctionDecl *Fn = cast<FunctionDecl>(D);
   if (!Fn->isInlineSpecified()) {
     S.Diag(Attr.getLoc(), diag::warn_gnu_inline_attribute_requires_inline);
     return;
   }

cfe/trunk/test/SemaCUDA/bad-attributes.cu

 // Tests handling of CUDA attributes that are bad either because they're
 // applied to the wrong sort of thing, or because they're given in illegal
 // combinations.
 //
 // You should be able to run this file through nvcc for compatibility testing.
 //
-// RUN: %clang_cc1 -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsyntax-only -Wcuda-compat -verify -DEXPECT_INLINE_WARNING %s
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -Wcuda-compat -verify %s
 
 #include "Inputs/cuda.h"
 
 // Try applying attributes to functions and variables.  Some should generate
 // warnings; others not.
 __device__ int a1;
 __device__ void a2();
 __host__ int b1; // expected-warning {{attribute only applies to functions}}

 __shared__ __device__ int z9;
 __shared__ __constant__ int z10;  // expected-error {{attributes are not compatible}}
 // expected-note@-1 {{conflicting attribute is here}}
 
 __global__ __device__ void z11();  // expected-error {{attributes are not compatible}}
 // expected-note@-1 {{conflicting attribute is here}}
 __global__ __host__ void z12();  // expected-error {{attributes are not compatible}}
 // expected-note@-1 {{conflicting attribute is here}}
+
+struct S {
+  __global__ void foo() {};  // expected-error {{must be a free function or static member function}}
+  __global__ static void bar(); // expected-warning {{kernel function 'bar' is a member function}}
+  // Although this is implicitly inline, we shouldn't warn.
+  __global__ static void baz() {}; // expected-warning {{kernel function 'baz' is a member function}}
+};
+
+__global__ static inline void foobar() {};
+#ifdef EXPECT_INLINE_WARNING
+// expected-warning@-2 {{ignored 'inline' attribute on kernel function 'foobar'}}
+#endif