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); } 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 __global__ void kernel(F f) { f(); } -// dev-note@-1 7{{called by 'kernel<(lambda}} +// dev-note@-1 3{{called by 'kernel<(lambda}} +// cuda-note@-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@-1 {{capture host side class data member by this pointer in device or host device lambda function}} + // cuda-error@-1 {{capture host side class data member by this pointer in device or host device lambda function}} kernel<<<1,1>>>([&](){ hd(b); }); - // dev-error@-1 {{capture host side class data member by this pointer in device or host device lambda function}} + // cuda-error@-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@-1 {{capture host side class data member by this pointer in device or host device lambda function}} + // cuda-error@-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@-1 {{capture host side class data member by this pointer in device or host device lambda function}} + // cuda-error@-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@-1 {{capture host side class data member by this pointer in device or host device lambda function}} + kernel<<<1,1>>>(lambda1); } };