Static Analysis - access violations

[danchitnis]

Hi, is there a way to detect buffer access violations in SYCL code? Running the code below results in different behaviour with a segmentation fault on HOST and CPU, but not on GPU(CUDA).

(@FMarno for finding the issue and the code)

#include <CL/sycl.hpp>
#include <iostream>

int main()
{
	constexpr size_t buffer_size = 10;
	sycl::default_selector selector;
	sycl::queue q{selector};
	auto device = q.get_device();
	std::cout << "running on device: " << device.get_info<sycl::info::device::name>() << '\n';

	sycl::buffer<float, 1> too_small{sycl::range<1>{buffer_size}};

	q.submit([&](sycl::handler &h)
			 {
				 auto acc_buffer = too_small.template get_access<sycl::access::mode::write>(h);
				 h.parallel_for(sycl::range<1>{buffer_size}, [=](sycl::id<1> id)
								{ acc_buffer[id] = id[0]; });
			 });

	q.submit([&](sycl::handler &h)
			 {
				 auto acc_buffer = too_small.template get_access<sycl::access::mode::read_write>(h);
				 h.parallel_for(sycl::range<1>{buffer_size + 100}, [=](sycl::id<1> id)
								{ acc_buffer[id] = acc_buffer[id] + 1; });
			 });

	auto host_acc = too_small.template get_access<sycl::access::mode::read>();
	for (size_t i = 0; i < host_acc.size(); ++i)
	{
		std::cout << host_acc[i] << ' ';
	}
	std::cout << '\n';
}

For CPU:

clang++ -fsycl small_buffer.cpp -o small
SYCL_DEVICE_FILTER=CPU ./small 

Results:

running on device: Intel(R) Core(TM) i7-9700 CPU @ 3.00GHz
1 2 3 4 5 6 7 8 9 10 
Segmentation fault

For GPU:

clang++ -fsycl small_buffer.cpp -o small -fsycl-targets=nvptx64-nvidia-cuda-sycldevice
SYCL_DEVICE_FILTER=GPU ./small 

Results:

running on device: NVIDIA RTX A4000
1 2 3 4 5 6 7 8 9 10 

If I run scan-build it will also fail due to an error.:

scan-build: The analyzer encountered problems on some source files.
scan-build: Preprocessed versions of these sources were deposited in '/tmp/scan-build-2021-08-20-163621-8285-1/failures'.
scan-build: Please consider submitting a bug report using these files:
scan-build:   http://clang-analyzer.llvm.org/filing_bugs.html

[zjin-lcf]

Is cuda-memcheck useful for you ?