// RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out
// RUN: %{run} %t.out %if !gpu || linux %{ | FileCheck %s %}

//==------------------ stream.cpp - SYCL stream basic test -----------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//

#include <sycl/detail/core.hpp>

#include <sycl/properties/all_properties.hpp>
#include <sycl/stream.hpp>

#include <cassert>

using namespace sycl;

int main() {
  {
    queue Queue(default_selector_v);
    context Context = Queue.get_context();

    // Check constructor and getters
    size_t sizeInKernel = 0;
    size_t workItemBufferSizeInKernel = 0;
    {
      sycl::buffer<size_t> bufSize(&sizeInKernel, 1);
      sycl::buffer<size_t> bufWorkItemBufferSize(&workItemBufferSizeInKernel,
                                                 1);

      // Check constructor and getters
      Queue.submit([&](handler &CGH) {
        stream Out(1024, 80, CGH);
        assert(Out.size() == 1024);
        assert(Out.get_work_item_buffer_size() == 80);

        sycl::accessor accSize(bufSize, CGH, sycl::write_only);
        sycl::accessor accWorkItemBufferSize(bufWorkItemBufferSize, CGH,
                                             sycl::write_only);

        CGH.single_task<class DummyTask1>([=]() {
          accSize[0] = Out.size();
          accWorkItemBufferSize[0] = Out.get_work_item_buffer_size();
        });
      });
    }
    assert(sizeInKernel == 1024);
    assert(workItemBufferSizeInKernel == 80);

    // Check common reference semantics
    std::hash<stream> Hasher;

    Queue.submit([&](handler &CGH) {
      stream Out1(1024, 80, CGH);
      stream Out2(Out1);

      assert(Out1 == Out2);
      assert(Hasher(Out1) == Hasher(Out2));

      stream Out3(std::move(Out1));

      assert(Out2 == Out3);
      assert(Hasher(Out2) == Hasher(Out3));

      CGH.single_task<class DummyTask2>([=]() {});
    });

    Queue.submit([&](handler &CGH) {
      stream Out(1024, 80, CGH);
      CGH.single_task<class integral>([=]() {
        // String
        Out << "Hello, World!\n";
        // CHECK: Hello, World!

        // Char
        Out << 'a' << '\n';
        // CHECK-NEXT: a

        // Boolean type
        Out << true << endl;
        Out << false << endl;
        // CHECK-NEXT: true
        // CHECK-NEXT: false

        // Integral types
        Out << static_cast<unsigned char>(123) << endl;
        Out << static_cast<signed char>(-123) << endl;
        Out << static_cast<short>(2344) << endl;
        Out << static_cast<signed short>(-2344) << endl;
        Out << 3454 << endl;
        Out << -3454 << endl;
        Out << 3454U << endl;
        Out << 3454L << endl;
        Out << 12345678901245UL << endl;
        Out << -12345678901245LL << endl;
        Out << 12345678901245ULL << endl;
        // CHECK-NEXT: 123
        // CHECK-NEXT: -123
        // CHECK-NEXT: 2344
        // CHECK-NEXT: -2344
        // CHECK-NEXT: 3454
        // CHECK-NEXT: -3454
        // CHECK-NEXT: 3454
        // CHECK-NEXT: 3454
        // CHECK-NEXT: 12345678901245
        // CHECK-NEXT: -12345678901245
        // CHECK-NEXT: 12345678901245

        // Floats
        Out << 33.4f << endl;
        Out << -33.4f << endl;
        Out << -1.0f / 0.0f << endl;
        Out << 1.0f / 0.0f << endl;
        Out << sycl::sqrt(-1.0f) << endl;
        // CHECK-NEXT: 33.4
        // CHECK-NEXT: -33.4
        // CHECK-NEXT: -inf
        // CHECK-NEXT: inf
        // CHECK-NEXT: nan

        // Manipulators for integral types
        Out << dec << 0213 << endl;
        Out << dec << 0x213A << endl;
        Out << oct << 139 << endl;
        Out << hex << 8506 << endl;
        // CHECK-NEXT: 139
        // CHECK-NEXT: 8506
        // CHECK-NEXT: 213
        // CHECK-NEXT: 213a

        Out << oct << showbase << 8506 << ' ' << noshowbase << 8506 << endl;
        Out << hex << showbase << 8506 << ' ' << noshowbase << 8506 << endl;
        Out << dec << showbase << 8506 << ' ' << noshowbase << 8506 << endl;
        // CHECK-NEXT: 020472 20472
        // CHECK-NEXT: 0x213a 213a
        // CHECK-NEXT: 8506 8506

        Out << dec << showpos << 234 << ' ' << noshowpos << 234 << endl;
        Out << hex << showpos << 234 << ' ' << noshowpos << 234 << endl;
        Out << oct << showpos << 234 << ' ' << noshowpos << 234 << endl;
        // CHECK-NEXT: +234 234
        // CHECK-NEXT: ea ea
        // CHECK-NEXT: 352 352

        Out << hex << showpos << -1 << ' ' << noshowpos << -1 << endl;
        Out << oct << showpos << -1 << ' ' << noshowpos << -1 << endl;
        Out << dec << showpos << -1 << ' ' << noshowpos << -1 << endl;
        // CHECK-NEXT: ffffffff ffffffff
        // CHECK-NEXT: 37777777777 37777777777
        // CHECK-NEXT: -1 -1

        // Pointers
        int a = 5;
        int *Ptr = &a;
        Out << Ptr << endl;
        const int *const ConstPtr = &a;
        Out << ConstPtr << endl;
        auto multiPtr = private_ptr<int>(Ptr);
        Out << multiPtr << endl;
        // CHECK-NEXT: 0x{{[0-9a-fA-F]*$}}
        // CHECK-NEXT: 0x{{[0-9a-fA-F]*$}}
        // CHECK-NEXT: 0x{{[0-9a-fA-F]*$}}

        // Vectors
        vec<int, 1> f1(545);
        Out << f1 << endl;
        vec<int, 2> f2(545, 645);
        Out << f2 << endl;
        vec<int, 3> f3(545, 645, 771);
        Out << f3 << endl;
        vec<int, 4> f4(542325, 645, 771, 1024);
        Out << hex << showbase << f4 << endl;
        Out << dec << f4 << endl;
        vec<float, 4> f5(542.3f, 645.3f, 771.6f, 1024.2f);
        Out << f5 << endl;
        // CHECK-NEXT: 545
        // CHECK-NEXT: 545, 645
        // CHECK-NEXT: 545, 645, 771
        // CHECK-NEXT: 0x84675, 0x285, 0x303, 0x400
        // CHECK-NEXT: 542325, 645, 771, 1024
        // CHECK-NEXT: 542.3, 645.3, 771.6, 1024.2

        // Swizzles
        Out << f4.odd() << endl;
        // CHECK-NEXT: 645, 1024

        // SYCL types
        Out << id<1>(23) << endl;
        Out << range<1>(32) << endl;
        Out << id<3>(11, 12, 13) << endl;
        Out << range<3>(11, 12, 13) << endl;
        Out << sycl::nd_range<3>(sycl::range<3>(2, 4, 1),
                                 sycl::range<3>(1, 2, 1))
            << endl;
        // CHECK-NEXT: {23}
        // CHECK-NEXT: {32}
        // CHECK-NEXT: {11, 12, 13}
        // CHECK-NEXT: {11, 12, 13}
        // CHECK-NEXT: nd_range(global_range: {2, 4, 1}, local_range: {1, 2, 1},
        // offset: {0, 0, 0})
      });
    });
    Queue.wait();

    if (Queue.get_device().has(sycl::aspect::fp64)) {
      Queue.submit([&](handler &CGH) {
        stream Out(1024, 80, CGH);
        CGH.single_task<class doubles>([=]() {
          // Double
          Out << 5.2 << endl;
          Out << -5.2 << endl;
          Out << 0.0003 << endl;
          Out << -1.0 / 0.0 << endl;
          Out << 1.0 / 0.0 << endl;
          Out << sycl::sqrt(-1.0) << endl;
        });
      });
      Queue.wait();
    } else {
      // Repeat skipped message same number of times as the number of skipped
      // output lines.
      for (size_t I = 0; I < 6; ++I)
        std::cout << "Skipped double test." << std::endl;
    }
    // CHECK-NEXT: {{(5.2|Skipped double test.)}}
    // CHECK-NEXT: {{(-5.2|Skipped double test.)}}
    // CHECK-NEXT: {{(0.0003|Skipped double test.)}}
    // CHECK-NEXT: {{(-inf|Skipped double test.)}}
    // CHECK-NEXT: {{(inf|Skipped double test.)}}
    // CHECK-NEXT: {{(nan|Skipped double test.)}}

    // Stream in parallel_for
    Queue.submit([&](handler &CGH) {
      stream Out(1024, 80, CGH);
      CGH.parallel_for<class stream_string>(
          range<1>(10), [=](id<1> i) { Out << "Hello, World!" << endl; });
    });
    Queue.wait();
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!

    // nd_item
    Queue.submit([&](handler &CGH) {
      stream Out(1024, 80, CGH);
      CGH.parallel_for<class stream_nd_range>(
          nd_range<3>(range<3>(4, 4, 4), range<3>(2, 2, 2)),
          [=](nd_item<3> item) {
            if (item.get_global_id(0) == 1 && item.get_global_id(1) == 2 &&
                item.get_global_id(2) == 3)
              Out << item << endl;
          });
    });
    Queue.wait();
    // CHECK-NEXT: nd_item(global_id: {1, 2, 3}, local_id: {1, 0, 1})

    Queue.submit([&](handler &CGH) {
      stream Out(1024, 200, CGH);
      CGH.parallel_for_work_group<class stream_h_item>(
          range<3>(1, 1, 1), range<3>(1, 1, 1), [=](group<3> Group) {
            Group.parallel_for_work_item(
                [&](h_item<3> Item) { Out << Item << endl; });
          });
    });
    Queue.wait();
    // CHECK-NEXT: h_item(
    // CHECK-NEXT:   global item(range: {1, 1, 1}, id: {0, 0, 0})
    // CHECK-NEXT:   logical local item(range: {1, 1, 1}, id: {0, 0, 0})
    // CHECK-NEXT:   physical local item(range: {1, 1, 1}, id: {0, 0, 0})
    // CHECK-NEXT: )

    // Multiple streams in command group
    Queue.submit([&](handler &CGH) {
      stream Out1(1024, 80, CGH);
      stream Out2(500, 20, CGH);
      CGH.parallel_for<class multiple_streams>(range<1>(2), [=](id<1> i) {
        Out1 << "Hello, World!" << endl;
        Out2 << "Hello, World!" << endl;
      });
    });
    Queue.wait();
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!
    // CHECK-NEXT: Hello, World!

    // The case when stream buffer is full. To check that there is no problem
    // with end of line symbol when printing out the stream buffer.
    Queue.submit([&](handler &CGH) {
      stream Out(10, 10, CGH);
      CGH.parallel_for<class full_stream_buffer>(
          range<1>(2), [=](id<1> i) { Out << "aaaaaaaaa" << endl; });
    });
    Queue.wait();
    // CHECK-NEXT: aaaaaaaaa

    // Use a big statement size to verify the stream internal implementation can
    // create a big enough flush buffer in global memory to handle this case.
    range<1> global = 16;
    range<1> local = 16;
    Queue.submit([&](handler &cgh) {
      stream ostream(198, 8192, cgh);
      cgh.parallel_for<class test_stream>(
          nd_range<1>(global, local), [=](nd_item<1> it) {
            ostream << "global id " << it.get_global_id(0)
                    << stream_manipulator::endl;
          });
    });
    Queue.wait();
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
    // CHECK: global id {{[0-9]+}}
  }

  return 0;
}