Multiple accessor to the same buffer with different ddr_bank property doesn't work

[lforg37]

When running a program with multiple kernels having arguments mapped on different DDR bank the program fail due to buffer allocated on wrong DDR bank.

To Reproduce
Run the sycl/test/on-device/xocc/simple_tests/ddr_bank_test.cpp test.

// RUN: %clangxx -fsycl -std=c++20 -fsycl-targets=%sycl_triple %s -o %t.out 2>&1 | FileCheck %s
// RUN: %ACC_RUN_PLACEHOLDER %t.out

#include <sycl/sycl.hpp>
#include <sycl/ext/xilinx/fpga.hpp>

int main() {
  cl::sycl::buffer<cl::sycl::cl_int, 1> Buffer(4);
  cl::sycl::queue Queue;
  const cl::sycl::cl_int buf_size = Buffer.size();

  Queue.submit([&](cl::sycl::handler &cgh) {
    sycl::ext::oneapi::accessor_property_list PL{sycl::ext::xilinx::ddr_bank<1>};
    // CHECK:{{.*}}:DDR[1]
    sycl::accessor Accessor(Buffer, cgh, sycl::write_only, PL);
    cgh.single_task<class SmallerTesta>([=]{
          for (cl::sycl::cl_int i = 0 ; i < buf_size ; ++i)
            Accessor[i] = i;
        });
  });
  Queue.submit([&](cl::sycl::handler &cgh) {
    sycl::ext::oneapi::accessor_property_list PL{sycl::ext::xilinx::ddr_bank<3>};
    // CHECK:{{.*}}:DDR[3]
    sycl::accessor Accessor(Buffer, cgh, sycl::write_only, PL);
    cgh.single_task<class SmallerTestb>([=]{
          for (cl::sycl::cl_int i = 0 ; i < buf_size ; ++i)
            Accessor[i] = i;
        });
  });
  Queue.submit([&](cl::sycl::handler &cgh) {
    sycl::ext::oneapi::accessor_property_list PL{sycl::ext::xilinx::ddr_bank<0>};
    // CHECK:{{.*}}:DDR[0]
    sycl::accessor Accessor(Buffer, cgh, sycl::write_only, PL);
    cgh.single_task<class SmallerTestc>([=]{
          for (cl::sycl::cl_int i = 0 ; i < buf_size ; ++i)
            Accessor[i] = i;
        });
  });
}

There is only one buffer created on the device, where three copies are required : one on each target ddr bank.

[lforg37]

Opening a discussion on intel repo to compare with what is done for their intel::property::buffer_location property.