[clang] 26e492e - [HIP] Warn capture this pointer in device lambda

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2021-09-08T13:45:26-04:00
New Revision: 26e492e134c006c63b3d9f9f9eabdeba014b1d2c

URL: https://github.com/llvm/llvm-project/commit/26e492e134c006c63b3d9f9f9eabdeba014b1d2c
DIFF: https://github.com/llvm/llvm-project/commit/26e492e134c006c63b3d9f9f9eabdeba014b1d2c.diff

LOG: [HIP] Warn capture this pointer in device lambda

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.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D108493

Added: 
    

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

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index bebaf8fc9f0bb..f8e89549f0503 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8391,8 +8391,10 @@ def err_ref_bad_target_global_initializer : Error<
   "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>;

diff  --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 75364c10c154b..840b3daae63cf 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -878,8 +878,13 @@ void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
                           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;
 }

diff  --git a/clang/test/SemaCUDA/lambda.cu b/clang/test/SemaCUDA/lambda.cu
index 6f305a683c009..524cdd429bf53 100644
--- a/clang/test/SemaCUDA/lambda.cu
+++ b/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 @@ auto global_lambda = [] () { return 123; };
 
 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 @@ public:
     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);
   }
 };