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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added a reviewer: tra.
yaxunl requested review of this revision.

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.


https://reviews.llvm.org/D108493

Files:
  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,7 @@
 // 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,cuda %s
+// RUN: %clang_cc1 -std=c++17 -fsyntax-only -fcuda-is-device -verify=com,dev \
+// RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip %s
 
 #include "Inputs/cuda.h"
 
@@ -7,7 +9,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}}
+// cuda-note at -2 5{{called by 'kernel<(lambda}}
 
 __host__ __device__ void hd(int x);
 
@@ -22,19 +25,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}}
+    // cuda-error at -1 {{capture host side class data member by this pointer in device or host device lambda function}}
 
     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}}
+    // cuda-error at -1 {{capture host side class data member by this pointer in device or host device lambda function}}
 
     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}}
+    // cuda-error at -1 {{capture host side class data member by this pointer in device or host device lambda function}}
 
     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}}
+      // cuda-error at -1 {{capture host side class data member by this pointer in device or host device lambda function}}
       f();
     });
+
+    auto lambda1 = [this] __device__ { hd(this->b); };
+    // cuda-error at -1 {{capture host side class data member by this pointer in device or host device lambda function}}
+    kernel<<<1,1>>>(lambda1);
   }
 };
 
Index: clang/lib/Sema/SemaCUDA.cpp
===================================================================
--- clang/lib/Sema/SemaCUDA.cpp
+++ clang/lib/Sema/SemaCUDA.cpp
@@ -877,7 +877,10 @@
     SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
                           diag::err_capture_bad_target, Callee, *this)
         << Capture.getVariable();
-  } else if (Capture.isThisCapture()) {
+  } else if (Capture.isThisCapture() && !LangOpts.HIP) {
+    // Capture of this pointer is allowed for HIP since this pointer may be
+    // pointing to managed memory which is accessible on both device and
+    // host sides.
     SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
                           diag::err_capture_bad_target_this_ptr, Callee, *this);
   }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D108493.367912.patch
Type: text/x-patch
Size: 3068 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210820/1b9e56fa/attachment.bin>