SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute

[AmirIpma]

Cannot call member functions size and get_max_statement_size of sycl::stream

Sample code:

#include <sycl.hpp>
#include <vector>

struct test_move {
    static constexpr std::size_t result_count = 2;

    template <typename Storage, typename T, typename VectorType>
    static void run(VectorType& vec, T& t) {
        std::size_t i = 0;

        Storage storage(t);

        // Move constructible
        T s0(std::move(t));
        vec[i++] = storage.check(s0);

        // Move assignable
        T s1 = std::move(s0);  // NB: move s0, not t
        vec[i++] = storage.check(s1);
    }
};

struct stream_kernel {
  void operator()() const {}
};

struct storage {
    std::size_t size;
    std::size_t max_statement_size;

    explicit storage(const sycl::stream& stream)
        : size(stream.size()),
        max_statement_size(stream.get_max_statement_size()
        ) {
    }

    bool check(const sycl::stream& stream) const {
        return (
            stream.size() == size &&
            stream.get_max_statement_size() == max_statement_size
        );
    }
};

template <typename storage, typename T, typename InitFunc>
void check_kernel(InitFunc init_func, const std::string& type_name) {
    std::vector<int> results(test_move::result_count, false);
    {
        sycl::buffer<int> buffer(results.data(), sycl::range<1>{test_move::result_count});
        sycl::queue queue;
        queue.submit([&](sycl::handler& cgh) {
            auto accessor = buffer.template get_access<sycl::access_mode::write>(cgh);
            T t = init_func(cgh);
            // use non-simple parallel_for to be able to use local_accessor
            cgh.parallel_for<stream_kernel>(
            sycl::nd_range<1>{sycl::range<1>{1}, sycl::range<1>{1}},
            [=](sycl::nd_item<1> nd_item) {
                auto ptr = accessor.begin();
                test_move::run<storage>(ptr, t);  // last since invalidates instance
                ptr += test_move::result_count;
            });
        });
    }
}

int main() {
    check_kernel<storage, sycl::stream>(
        [&](sycl::handler& cgh) {
        return sycl::stream{0, 0, cgh};
        },
        "stream");

    return 0;
}

Build:
clang++ -fsycl -Wdeprecated-declarations stream_semantics.cpp -o sample

Output:

stream_semantics.cpp:39:20: error: SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute
            stream.size() == size &&
                   ^
/users/aipmagax/dpcpp_compiler/bin/../include/sycl/stream.hpp:760:10: note: 'size' declared here
  size_t size() const noexcept;
         ^
stream_semantics.cpp:37:10: note: called by 'check'
    bool check(const sycl::stream& stream) const {
         ^
stream_semantics.cpp:40:20: error: SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute
            stream.get_max_statement_size() == max_statement_size
                   ^
/users/aipmagax/dpcpp_compiler/bin/../include/sycl/stream.hpp:770:10: note: 'get_max_statement_size' declared here
  size_t get_max_statement_size() const;
         ^
stream_semantics.cpp:37:10: note: called by 'check'
    bool check(const sycl::stream& stream) const {
         ^
2 warnings and 2 errors generated.

• OS: Linux
• DPC++ version: clang version 16.0.0 (/netbatch/donb2930466_00/runDir/jenkins-dir/workspace/Tools_SH/SYCLOS_Nightly/RHEL/llvm.src/clang 7589ac9)
  Target: x86_64-unknown-linux-gnu
  Thread model: posix
• Dependencies version: oclcpuexp-2022.14.8.0.04