intel / llvm

Intel staging area for llvm.org contribution. Home for Intel LLVM-based projects.
Other
1.25k stars 738 forks source link

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

Open fbarbari opened 1 month ago

fbarbari commented 1 month ago

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

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 commented 1 month ago

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).