Unused `sycl::accessor` detected only when created with `buffer.get_access(...)`

[fbarbari]

Describe the bug

It appears that icpx is able to detect unused accessors only when they are created with buffer.get_access(...). For example, creating them with sycl::accessor{buffer, ...} does not generate any warning.

To reproduce

I have this simple program:

#include <iostream>
#include <vector>
#include <numeric>

#include <sycl/sycl.hpp>

int main() {
	const size_t n = 10;
	std::vector<int> x(n);

	std::iota(x.begin(), x.end(), 1);

	sycl::queue q{sycl::default_selector_v};
	
	{
		auto x_buf = sycl::buffer<int, 1>{x.data(), sycl::range<1>{n}};

		q.submit([&](sycl::handler& cgh) {
			sycl::accessor x_acc = x_buf.get_access(cgh, sycl::read_write);
			auto unused_acc = sycl::accessor{x_buf, cgh, sycl::read_write};

			cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx) {
				x_acc[idx] *= 2;
			});
		});

		q.wait_and_throw();
	}

	for(size_t i{0}; i < n; i++) {
		std::cout << " x[" << i << "] = " << x.at(i) << std::endl;
	}

	return 0;
}

with an unused accessor which doesn't get detected when I compile with icpx --std=c++17 -fsycl -Wall -Wextra -Wpedantic -Werror -Wshadow main.cpp -o main.x.

However, If I change it to:

#include <iostream>
#include <vector>
#include <numeric>

#include <sycl/sycl.hpp>

int main() {
	const size_t n = 10;
	std::vector<int> x(n);

	std::iota(x.begin(), x.end(), 1);

	sycl::queue q{sycl::default_selector_v};
	
	{
		auto x_buf = sycl::buffer<int, 1>{x.data(), sycl::range<1>{n}};

		q.submit([&](sycl::handler& cgh) {
			sycl::accessor x_acc = x_buf.get_access(cgh, sycl::read_write);
			sycl::accessor unused_acc = x_buf.get_access(cgh, sycl::read_write);

			cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx) {
				x_acc[idx] *= 2;
			});
		});

		q.wait_and_throw();
	}

	for(size_t i{0}; i < n; i++) {
		std::cout << " x[" << i << "] = " << x.at(i) << std::endl;
	}

	return 0;
}

It gets detected and the compilation fails with:

main.cpp:20:19: error: unused variable 'unused_acc' [-Werror,-Wunused-variable]
   20 |                         sycl::accessor unused_acc = x_buf.get_access(cgh, sycl::read_write);
      |                                        ^~~~~~~~~~
1 error generated.

Environment

• OS: Ubuntu 22.04.3
• Target device and vendor: Intel(R) Core(TM) i5-8265U CPU @ 1.60GHz
• DPC++ version:

Intel(R) oneAPI DPC++/C++ Compiler 2024.2.1 (2024.2.1.20240711)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/intel/oneapi/compiler/2024.2/bin/compiler
Configuration file: /opt/intel/oneapi/compiler/2024.2/bin/compiler/../icpx.cfg

• Dependencies version:

[opencl:cpu][opencl:0] Intel(R) OpenCL, Intel(R) Core(TM) i5-8265U CPU @ 1.60GHz OpenCL 3.0 (Build 0) [2024.18.7.0.11_160000]
[opencl:gpu][opencl:1] Intel(R) OpenCL Graphics, Intel(R) UHD Graphics 620 OpenCL 3.0 NEO  [23.30.26918.9]

Platforms: 2
Platform [#1]:
    Version  : OpenCL 3.0 LINUX
    Name     : Intel(R) OpenCL
    Vendor   : Intel(R) Corporation
    Devices  : 1
        Device [#0]:
        Type       : cpu
        Version    : OpenCL 3.0 (Build 0)
        Name       : Intel(R) Core(TM) i5-8265U CPU @ 1.60GHz
        Vendor     : Intel(R) Corporation
        Driver     : 2024.18.7.0.11_160000
        Aspects    : cpu fp16 fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations usm_system_allocations usm_atomic_host_allocations usm_atomic_shared_allocations atomic64 ext_oneapi_srgb ext_oneapi_native_assert ext_intel_legacy_image ext_oneapi_ballot_group ext_oneapi_fixed_size_group ext_oneapi_opportunistic_group ext_oneapi_tangle_group
        info::device::sub_group_sizes: 4 8 16 32 64
Platform [#2]:
    Version  : OpenCL 3.0 
    Name     : Intel(R) OpenCL Graphics
    Vendor   : Intel(R) Corporation
    Devices  : 1
        Device [#1]:
        Type       : gpu
        Version    : OpenCL 3.0 NEO 
        Name       : Intel(R) UHD Graphics 620
        Vendor     : Intel(R) Corporation
        Driver     : 23.30.26918.9
        Aspects    : gpu fp16 fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations atomic64 ext_oneapi_srgb ext_intel_device_id ext_intel_legacy_image ext_intel_esimd ext_oneapi_ballot_group ext_oneapi_fixed_size_group ext_oneapi_opportunistic_group ext_oneapi_tangle_group
        info::device::sub_group_sizes: 8 16 32
default_selector()      : gpu, Intel(R) OpenCL Graphics, Intel(R) UHD Graphics 620 OpenCL 3.0 NEO  [23.30.26918.9]
accelerator_selector()  : No device of requested type available. Please chec...
cpu_selector()          : cpu, Intel(R) OpenCL, Intel(R) Core(TM) i5-8265U CPU @ 1.60GHz OpenCL 3.0 (Build 0) [2024.18.7.0.11_160000]
gpu_selector()          : gpu, Intel(R) OpenCL Graphics, Intel(R) UHD Graphics 620 OpenCL 3.0 NEO  [23.30.26918.9]
custom_selector(gpu)    : gpu, Intel(R) OpenCL Graphics, Intel(R) UHD Graphics 620 OpenCL 3.0 NEO  [23.30.26918.9]
custom_selector(cpu)    : cpu, Intel(R) OpenCL, Intel(R) Core(TM) i5-8265U CPU @ 1.60GHz OpenCL 3.0 (Build 0) [2024.18.7.0.11_160000]
custom_selector(acc)    : No device of requested type available. Please chec...

Additional context

No response

[AlexeySachkov]

accessor constructor has a side effect, because it calls non-const methods on a buffer object which was passed to it by reference. So, the variable cannot be safely removed, because it would change the state of the program.

buffer::get_access also invokes the very same accessor constructor, but it then returns it and the resulting unused_acc is probably initialized differently, i.e. via copy constructor of accessor which doesn't have the same side effect, so that resulting variable can be safely removed.

That is my understanding as to why it could be happening. The results showcased in the description do look confusing, but I'm not 100% if that's a bug due to the said side effect (which is legal, because buffer is passed to accessor as a non-const reference in accordance with the SYCL 2020 spec).