Skip to content

[SYCL] Compiler cast constant address space to generic address space #40

Closed
@Naghasan

Description

@Naghasan

The following code compiles without any complain

template <typename T, size_t N>
void test(const std::array<T, N> &VA, const std::array<T, N> &VB,
          std::array<T, N> &VC, bool useCst) {
  cl::sycl::queue deviceQueue;
  cl::sycl::range<1> numOfItems{N};
  cl::sycl::buffer<T, 1> bufferA(VA.data(), numOfItems);
  cl::sycl::buffer<T, 1> bufferB(VB.data(), numOfItems);
  cl::sycl::buffer<T, 1> bufferC(VC.data(), numOfItems);

  deviceQueue.submit([&](cl::sycl::handler &cgh) {
    auto accessorA = bufferA.template get_access<sycl_read>(cgh);
    cl::sycl::accessor<T, 1, cl::sycl::access::mode::read,
                       cl::sycl::access::target::constant_buffer,
                       cl::sycl::access::placeholder::false_t>
        accessorB(bufferB, cgh);
    auto accessorC = bufferC.template get_access<sycl_write>(cgh);

    cgh.parallel_for<class Test<T>>(numOfItems,
    [=](cl::sycl::id<1> wiID) {
      T *my_ptr = useCst ? &accessorB.get_pointer()[wiID[0]]
                         : &accessorC.get_pointer()[wiID[0]];
      *my_ptr = accessorA[wiID];
    });
  });
}

Although this force an illegal address space cast as the SPIR-V validator says:

error: line 491: Expected input to have storage class Generic: GenericCastToPtr
  %372 = OpGenericCastToPtr %_ptr_Function_uint %371

Metadata

Metadata

Assignees

No one assigned

    Labels

    bugSomething isn't working

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions