[SYCL] Add a test for group local memory w/o device code opt

[sergey-semenov]

No description provided.

[Pennycook]

I find this test a little hard to follow, and it's not immediately clear what it's testing. Would something like the below be sufficient as a regression test for the issue here, or is it more complicated than I realize?

int main() {
  queue Q;
  {
    std::vector<int*> VecA(Size, 0);
    std::vector<int*> VecB(Size, 0);
    buffer<int*, 1> BufA{VecA.data(), range<1>(Size)};
    buffer<int*, 1> BufB{VecB.data(), range<1>(Size)};

    Q.submit([&](handler &Cgh) {
      auto AccA = BufA.get_access<access::mode::read_write>(Cgh);
      auto AccB = BufB.get_access<access::mode::read_write>(Cgh);
      Cgh.parallel_for<KernelA>(
          nd_range<1>(range<1>(Size), range<1>(WgSize)), [=](nd_item<1> Item) {
            multi_ptr<int[WgSize], access::address_space::local_space> PtrA =
                group_local_memory_for_overwrite<int[WgSize]>(Item.get_group());
            multi_ptr<int[WgSize], access::address_space::local_space> PtrB =
                group_local_memory_for_overwrite<int[WgSize]>(Item.get_group());
            AccA[GlobalIdx] = PtrA;
            AccB[GlobalIdx] = PtrB;
          });
    });

    auto AccA = BufA.get_access<access::mode::read>();
    auto AccB = BufB.get_access<access::mode::read>();
    for (size_t I = 0; I < Size; ++I) {
      assert(AccA[I] != AccB[I]);
    }
  }
}

[sergey-semenov]

@Pennycook Agreed, the initial version of the test did more than what was necessary here. Updated with the suggested approach.

[mlychkov]

I suppose that failed tests are not related to this patch but could you please check that there is an issue created for those failures?