[HIP] Warn capture this pointer in device lambda

[yaxunl]

Summary

HIP currently diagnose capture of this pointer in device lambda in
host member functions. If this pointer points to managed memory,
it can be used in both device and host functions. Under this
situation, capturing this pointer in device lambda functions
in host member functions is valid usage. Change the diagnostic
about capturing this pointer to warning.

[tra]

clang/lib/Sema/SemaCUDA.cpp

881–882

I assume there's no easy way to tell if it's a managed pointer or not.
Capturing a non-managed pointer would still be bad.
Should we make it a warning instead?

[yaxunl]

clang/lib/Sema/SemaCUDA.cpp

881–882

warning makes sense. will do.

Also this should apply to CUDA too. Although clang does not support __managed__ keyword for CUDA, this pointer may be allocated in managed memory through host API's.

[tra]

clang/lib/Sema/SemaCUDA.cpp

881–882

Also this should apply to CUDA too.

I'm not sure I understand what you mean.

Do you suggest that the warning should be issued for both HIP and CUDA? If so, the diag already applies to CUDA.

Or that we should allow capturing wrong-side this w/o diagnostics for both CUDA and HIP?

I'm fine with the former. But I do not think we should silently allow capturing wrong-side this for CUDA, because it will be an error most of the time.

[tra]

clang/lib/Sema/SemaCUDA.cpp

881–882

Ah. Never mind. I've replied before your latest patch update.
LGTM

clang/include/clang/Basic/DiagnosticSemaKinds.td

def note_cuda_const_var_unpromoted : Note<
   "const variable cannot be emitted on device side due to dynamic initialization">;
 def note_cuda_host_var : Note<
   "host variable declared here">;
 def err_ref_bad_target_global_initializer : Error<
   "reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
   "function %1 in global initializer">;
 def err_capture_bad_target : Error<
   "capture host variable %0 by reference in device or host device lambda function">;
-def err_capture_bad_target_this_ptr : Error<
-  "capture host side class data member by this pointer in device or host device lambda function">;
+def warn_maybe_capture_bad_target_this_ptr : Warning<
+  "capture host side class data member by this pointer in device or host device lambda function "
+  "may result in invalid memory access if this pointer is not accessible on device side">,
+  InGroup<DiagGroup<"gpu-maybe-wrong-side">>;
 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 err_variadic_device_fn : Error<
   "CUDA device code does not support variadic functions">;

clang/lib/Sema/SemaCUDA.cpp

void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
   if (!ShouldCheck || !Capture.isReferenceCapture())
     return;
   auto DiagKind = SemaDiagnosticBuilder::K_Deferred;
   if (Capture.isVariableCapture()) {
     SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
                           diag::err_capture_bad_target, Callee, *this)
         << Capture.getVariable();
   } else if (Capture.isThisCapture()) {
+    // Capture of this pointer is allowed since this pointer may be pointing to
+    // managed memory which is accessible on both device and host sides. It only
+    // results in invalid memory access if this pointer points to memory not
+    // accessible on device side.
     SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
-                          diag::err_capture_bad_target_this_ptr, Callee, *this);
+                          diag::warn_maybe_capture_bad_target_this_ptr, Callee,
+                          *this);
   }
   return;
 }
 
 void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   if (Method->hasAttr<CUDAHostAttr>() || Method->hasAttr<CUDADeviceAttr>())
     return;

clang/test/SemaCUDA/lambda.cu

 // RUN: %clang_cc1 -std=c++17 -fsyntax-only -verify=com %s
-// RUN: %clang_cc1 -std=c++17 -fsyntax-only -fcuda-is-device -verify=com,dev %s
+// RUN: %clang_cc1 -std=c++17 -fsyntax-only -fcuda-is-device -verify=com,dev,warn %s
+// RUN: %clang_cc1 -std=c++17 -fsyntax-only -fcuda-is-device -verify=com,dev,warn \
+// RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip %s
+// RUN: %clang_cc1 -std=c++17 -fsyntax-only -fcuda-is-device -verify=com,dev \
+// RUN:   -Wno-gpu-maybe-wrong-side %s
 
 #include "Inputs/cuda.h"
 
 auto global_lambda = [] () { return 123; };
 
 template<class F>
 __global__ void kernel(F f) { f(); }
-// dev-note@-1 7{{called by 'kernel<(lambda}}
+// dev-note@-1 3{{called by 'kernel<(lambda}}
+// warn-note@-2 5{{called by 'kernel<(lambda}}
 
 __host__ __device__ void hd(int x);
 
 class A {
   int b;
 public:
   void test() {
     [=](){ hd(b); }();
 
     [&](){ hd(b); }();
 
     kernel<<<1,1>>>([](){ hd(0); });
 
     kernel<<<1,1>>>([=](){ hd(b); });
-    // dev-error@-1 {{capture host side class data member by this pointer in device or host device lambda function}}
+    // warn-warning@-1 {{capture host side class data member by this pointer in device or host device lambda function may result in invalid memory access if this pointer is not accessible on device side}}
 
     kernel<<<1,1>>>([&](){ hd(b); });
-    // dev-error@-1 {{capture host side class data member by this pointer in device or host device lambda function}}
+    // warn-warning@-1 {{capture host side class data member by this pointer in device or host device lambda function may result in invalid memory access if this pointer is not accessible on device side}}
 
     kernel<<<1,1>>>([&] __device__ (){ hd(b); });
-    // dev-error@-1 {{capture host side class data member by this pointer in device or host device lambda function}}
+    // warn-warning@-1 {{capture host side class data member by this pointer in device or host device lambda function may result in invalid memory access if this pointer is not accessible on device side}}
 
     kernel<<<1,1>>>([&](){
       auto f = [&]{ hd(b); };
-      // dev-error@-1 {{capture host side class data member by this pointer in device or host device lambda function}}
+      // warn-warning@-1 {{capture host side class data member by this pointer in device or host device lambda function may result in invalid memory access if this pointer is not accessible on device side}}
       f();
     });
+
+    auto lambda1 = [this] __device__ { hd(this->b); };
+    // warn-warning@-1 {{capture host side class data member by this pointer in device or host device lambda function may result in invalid memory access if this pointer is not accessible on device side}}
+    kernel<<<1,1>>>(lambda1);
   }
 };
 
 int main(void) {
   auto lambda_kernel = [&]__global__(){};
   // com-error@-1 {{kernel function 'operator()' must be a free function or static member function}}
 
   int b;