Skip to content

Commit

Permalink
[SYCL][E2E] Fix local memory optimized out in AddressSanitizer test g…
Browse files Browse the repository at this point in the history
…roup_local_memory.cpp

Since intel#16356 local memory is lowered
to global variable before AddressSanitizerPass. The local memory access
is optimized out before AddressSanitizerPass.
This PR updates the test so that it won't be optimized out.
  • Loading branch information
wenju-he committed Feb 11, 2025
1 parent 36ce10e commit 0b1a6b9
Showing 1 changed file with 6 additions and 1 deletion.
Original file line number Diff line number Diff line change
Expand Up @@ -16,20 +16,25 @@ constexpr std::size_t group_size = 8;
int main() {
sycl::queue Q;

auto *array = sycl::malloc_device<int>(N, Q);

Q.submit([&](sycl::handler &h) {
h.parallel_for<class MyKernel>(
sycl::nd_range<1>(N, group_size), [=](sycl::nd_item<1> item) {
sycl::multi_ptr<int[N], sycl::access::address_space::local_space>
ptr = sycl::ext::oneapi::group_local_memory<int[N]>(
item.get_group());
auto &ref = *ptr;
ref[item.get_local_linear_id() * 2 + 4] = 42;
size_t lid = item.get_local_linear_id();
ref[lid * 2 + 4] = 42;
array[item.get_global_linear_id()] = ref[lid];
// CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Local Memory
// CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}} LID(6, 0, 0) GID({{.*}}, 0, 0)
// CHECK: #0 {{.*}} {{.*group_local_memory.cpp}}:[[@LINE-3]]
});
});

Q.wait();
sycl::free(array, Q);
return 0;
}

0 comments on commit 0b1a6b9

Please sign in to comment.