SYCL :: Basic/accessor/accessor.cpp is failing on PVC/Battlemage with opencl:gpu

[againull]

Describe the bug

See: https://github.com/intel/llvm/actions/runs/13585754417/job/37981521438?pr=17237

2025-02-28T10:32:11.1598939Z -- Testing: 2269 tests, 24 workers --
2025-02-28T10:32:11.1599158Z FAIL: SYCL :: Basic/accessor/accessor.cpp (117 of 2269)
2025-02-28T10:32:11.1599406Z ******************** TEST 'SYCL :: Basic/accessor/accessor.cpp' FAILED ********************
2025-02-28T10:32:11.1599618Z Exit Code: 134
2025-02-28T10:32:11.1599696Z 
2025-02-28T10:32:11.1599760Z Command Output (stdout):
2025-02-28T10:32:11.1599895Z --
2025-02-28T10:32:11.1600135Z exception received: placeholder accessor must be bound by calling handler::require() before it can be used.
2025-02-28T10:32:11.1600930Z exception received: placeholder accessor must be bound by calling handler::require() before it can be used.
2025-02-28T10:32:11.1601345Z exception received: placeholder accessor must be bound by calling handler::require() before it can be used.
2025-02-28T10:32:11.1601725Z exception received: placeholder accessor must be bound by calling handler::require() before it can be used.
2025-02-28T10:32:11.1602219Z SYCL exception caught: A local accessor must not be used in a SYCL kernel function that is invoked via single_task or via the simple form of parallel_for that takes a range parameter.
2025-02-28T10:32:11.1602852Z SYCL exception caught: A local accessor must not be used in a SYCL kernel function that is invoked via single_task or via the simple form of parallel_for that takes a range parameter.
2025-02-28T10:32:11.1603258Z Test passed
2025-02-28T10:32:11.1603535Z exception received: placeholder accessor must be bound by calling handler::require() before it can be used.
2025-02-28T10:32:11.1603964Z exception received: placeholder accessor must be bound by calling handler::require() before it can be used.
2025-02-28T10:32:11.1604382Z exception received: placeholder accessor must be bound by calling handler::require() before it can be used.
2025-02-28T10:32:11.1604743Z exception received: placeholder accessor must be bound by calling handler::require() before it can be used.
2025-02-28T10:32:11.1605322Z SYCL exception caught: A local accessor must not be used in a SYCL kernel function that is invoked via single_task or via the simple form of parallel_for that takes a range parameter.
2025-02-28T10:32:11.1605877Z SYCL exception caught: A local accessor must not be used in a SYCL kernel function that is invoked via single_task or via the simple form of parallel_for that takes a range parameter.
2025-02-28T10:32:11.1606173Z 
2025-02-28T10:32:11.1606227Z --
2025-02-28T10:32:11.1606344Z Command Output (stderr):
2025-02-28T10:32:11.1606484Z --
2025-02-28T10:32:11.1606748Z RUN: at line 2: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu  /__w/llvm/llvm/build-e2e/Basic/accessor/Output/accessor.cpp.tmp.out
2025-02-28T10:32:11.1607158Z + env ONEAPI_DEVICE_SELECTOR=level_zero:gpu /__w/llvm/llvm/build-e2e/Basic/accessor/Output/accessor.cpp.tmp.out
2025-02-28T10:32:11.1607552Z RUN: at line 2: env ONEAPI_DEVICE_SELECTOR=opencl:gpu  /__w/llvm/llvm/build-e2e/Basic/accessor/Output/accessor.cpp.tmp.out
2025-02-28T10:32:11.1607931Z + env ONEAPI_DEVICE_SELECTOR=opencl:gpu /__w/llvm/llvm/build-e2e/Basic/accessor/Output/accessor.cpp.tmp.out
2025-02-28T10:32:11.1608361Z accessor.cpp.tmp.out: /__w/llvm/llvm/llvm/sycl/test-e2e/Basic/accessor/accessor.cpp:1298: int main(): Assertion `Data == 96' failed.
2025-02-28T10:32:11.1608978Z /__w/llvm/llvm/build-e2e/Basic/accessor/Output/accessor.cpp.script: line 3:  5685 Aborted                 (core dumped) env ONEAPI_DEVICE_SELECTOR=opencl:gpu /__w/llvm/llvm/build-e2e/Basic/accessor/Output/accessor.cpp.tmp.out
2025-02-28T10:32:11.1609419Z 
2025-02-28T10:32:11.1609464Z --
2025-02-28T10:32:11.1609528Z 
2025-02-28T10:32:11.1609580Z ********************

[aelovikov-intel]

Reproduced locally on PVC as well. Needed to run entire e2e in background while running this test in a loop in another terminal. Only fails with opencl:gpu and that happens in one of two places, either the one in the logs above or Assertion host_acc[i] == 42' failed.` around line 401. Trying to reduce the test by removing non-impactful checks.

[aelovikov-intel]

Test reduction update:

a.cpp:

#include <iostream>                                                   
#include <sycl/detail/core.hpp>                                       
                                                                      
int main() {                                                          
  sycl::queue queue;                                                  
  sycl::queue{}.single_task([]() {});                                 
                                                                      
  sycl::buffer<int, 1> buf{3};                                        
                                                                      
  queue.submit([&](sycl::handler &cgh) {                              
    sycl::accessor dev_acc{buf, cgh};                                 
                                                                      
    cgh.parallel_for(sycl::range<1>{3},                               
                     [=](sycl::id<1> index) { dev_acc[index] = 42; });
  });                                                                 
                                                                      
  sycl::host_accessor host_acc{buf};                                  
  for (int i = 0; i != 3; ++i)                                        
    std::cout << host_acc[i] << " ";                                  
  std::cout << std::endl;                                             
  for (int i = 0; i != 3; ++i)                                        
    assert(host_acc[i] == 42);                                        
                                                                      
  return 0;                                                           
}                                                                     

heavy.cpp:

#include <sycl/sycl.hpp>                             
                                                     
int main() {                                         
  constexpr int N = 64*1024*1024;                    
  sycl::queue q;                                     
  auto *p = sycl::malloc_device<float>(N, q);        
  for (int iter = 0; iter < 100; ++iter)             
    q.parallel_for(sycl::range<1>{N}, [=](auto ndi) {
       for (int i = 0; i < 100; ++i)                 
         p[i] = sycl::sin(i + p[i]);                 
     }).wait();                                      
}                                                    

$ clang++ -fsycl heavy.cpp -o heavy
$ clang++ -fsycl a.cpp 
$ ./heavy &
# Immediately after, while `heavy` is still running:
$ while ONEAPI_DEVICE_SELECTOR=opencl:gpu ./a.out > log.txt 2>&1 ; do true ; done ; cat log.txt
Aborted (core dumped)                                                                          
0 0 0                                                                                          
a.out: a.cpp:22: int main(): Assertion `host_acc[i] == 42' failed.                             

[aelovikov-intel]

Reproducible with USM as well:

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

int main() {
  sycl::queue queue;

  sycl::queue{}.single_task([]() {});
  // queue.single_task([]() {});

  auto *p = sycl::malloc_shared<int>(3, queue);

  auto e = queue.parallel_for(sycl::range<1>{3},
                              [=](sycl::id<1> index) { p[index] = 42; });

  e.wait();
  // queue.wait();

  for (int i = 0; i != 3; ++i)
    std::cout << p[i] << " ";
  std::cout << std::endl;
  for (int i = 0; i != 3; ++i)
    assert(p[i] == 42);

  return 0;
}

Note that using any of the commented-out lines in place of the one preceding it, results in stable pass.

[aelovikov-intel]

I've moved internal bug report associated with this GH issue to OpenCL GPU RT component.