This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG26e492e134c0: [HIP] Warn capture this pointer in device
lambda (authored by yaxunl).
Herald added a project: clang.
Repository:
rG LLVM Github Monorepo
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@-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);
@@ -22,19 +27,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}}
+ // 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);
}
};
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>;
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@-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);
@@ -22,19 +27,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}}
+ // 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);
}
};
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>;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
