[PATCH] D108493: [HIP] Warn capture this pointer in device lambda

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl updated this revision to Diff 371363.
yaxunl marked an inline comment as done.
yaxunl retitled this revision from "[HIP] Allow capture this pointer in device lambda" to "[HIP] Warn capture this pointer in device lambda".
yaxunl edited the summary of this revision.

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D108493/new/

https://reviews.llvm.org/D108493

Files:
  clang/include/clang/Basic/DiagnosticSemaKinds.td
  clang/lib/Sema/SemaCUDA.cpp
  clang/test/SemaCUDA/lambda.cu


Index: clang/test/SemaCUDA/lambda.cu
===================================================================
--- clang/test/SemaCUDA/lambda.cu
+++ clang/test/SemaCUDA/lambda.cu
@@ -1,5 +1,9 @@
 // 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"
 
@@ -7,7 +11,8 @@
 
 template<class F>
 __global__ void kernel(F f) { f(); }
-// dev-note at -1 7{{called by 'kernel<(lambda}}
+// dev-note at -1 3{{called by 'kernel<(lambda}}
+// warn-note at -2 5{{called by 'kernel<(lambda}}
 
 __host__ __device__ void hd(int x);
 
@@ -22,19 +27,23 @@
     kernel<<<1,1>>>([](){ hd(0); });
 
     kernel<<<1,1>>>([=](){ hd(b); });
-    // dev-error at -1 {{capture host side class data member by this pointer in device or host device lambda function}}
+    // warn-warning at -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 at -1 {{capture host side class data member by this pointer in device or host device lambda function}}
+    // warn-warning at -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 at -1 {{capture host side class data member by this pointer in device or host device lambda function}}
+    // warn-warning at -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 at -1 {{capture host side class data member by this pointer in device or host device lambda function}}
+      // warn-warning at -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 at -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);
   }
 };
 
Index: clang/lib/Sema/SemaCUDA.cpp
===================================================================
--- clang/lib/Sema/SemaCUDA.cpp
+++ clang/lib/Sema/SemaCUDA.cpp
@@ -878,8 +878,13 @@
                           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;
 }
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8391,8 +8391,10 @@
   "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>;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D108493.371363.patch
Type: text/x-patch
Size: 4745 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210908/6cb7a0cd/attachment.bin>