Skip to content

[DeviceSanitizer] Fix wrong test for nullptr detection #15467

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Sep 24, 2024
Merged

[DeviceSanitizer] Fix wrong test for nullptr detection #15467

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
417fcb1
fix wrong test
AllanZyne Sep 23, 2024
6ba0490
fix
AllanZyne Sep 23, 2024
5fd4ac1
fix O0 case
AllanZyne Sep 23, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
20 changes: 7 additions & 13 deletions libdevice/sanitizer_utils.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -120,22 +120,20 @@ __SYCL_PRIVATE__ void *ToPrivate(void *ptr) {
return __spirv_GenericCastToPtrExplicit_ToPrivate(ptr, 7);
}

inline bool ConvertGenericPointer(uptr &addr, uint32_t &as) {
inline void ConvertGenericPointer(uptr &addr, uint32_t &as) {
auto old = addr;
if ((addr = (uptr)ToPrivate((void *)old))) {
as = ADDRESS_SPACE_PRIVATE;
} else if ((addr = (uptr)ToLocal((void *)old))) {
as = ADDRESS_SPACE_LOCAL;
} else if ((addr = (uptr)ToGlobal((void *)old))) {
as = ADDRESS_SPACE_GLOBAL;
} else {
if (__AsanDebug)
__spirv_ocl_printf(__generic_to_fail, old);
return false;
// FIXME: I'm not sure if we need to check ADDRESS_SPACE_CONSTANT,
// but this can really simplify the generic pointer conversion logic
as = ADDRESS_SPACE_GLOBAL;
addr = old;
}
if (__AsanDebug)
__spirv_ocl_printf(__generic_to, old, addr, as);
return true;
}

inline uptr MemToShadow_CPU(uptr addr) {
Expand All @@ -144,9 +142,7 @@ inline uptr MemToShadow_CPU(uptr addr) {

inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
if (as == ADDRESS_SPACE_GENERIC) {
if (!ConvertGenericPointer(addr, as)) {
return 0;
}
ConvertGenericPointer(addr, as);
}

if (as == ADDRESS_SPACE_GLOBAL) { // global
Expand Down Expand Up @@ -233,9 +229,7 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {

inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
if (as == ADDRESS_SPACE_GENERIC) {
if (!ConvertGenericPointer(addr, as)) {
return 0;
}
ConvertGenericPointer(addr, as);
}

if (as == ADDRESS_SPACE_GLOBAL) { // global
Expand Down
9 changes: 1 addition & 8 deletions sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,21 +11,14 @@

#include <sycl/detail/core.hpp>

#include <sycl/ext/oneapi/experimental/address_cast.hpp>

int main() {
sycl::queue Q;
constexpr std::size_t N = 4;
int *array = nullptr;

Q.submit([&](sycl::handler &h) {
h.parallel_for<class MyKernel>(
sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) {
auto private_array =
sycl::ext::oneapi::experimental::static_address_cast<
sycl::access::address_space::private_space>(array);
private_array[0] = 0;
});
sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) { array[0] = 0; });
Q.wait();
});
// CHECK: ERROR: DeviceSanitizer: null-pointer-access on Unknown Memory
Expand Down
Loading