// The test checks that ESIMD headers don't break the ODR:
// two SYCL sources including ESIMD headers can be compiled and linked into a
// single executable w/o linker complaining about multiple symbol definitions.
// Template functions must have the same instantiation in both sources to cause
// ODR problems potentially - esimd::min is used for that purpose.
//
// CUDA/HIP do not support intrinsics generated by the ESIMD compilation path:
//
// RUN: %clangxx -fsycl -fsycl-targets=spir64-unknown-unknown -DSOURCE1 -c %s -o %t1.o
// RUN: %clangxx -fsycl -fsycl-targets=spir64-unknown-unknown -DSOURCE2 -c %s -o %t2.o
// RUN: %clangxx -fsycl -fsycl-targets=spir64-unknown-unknown %t1.o %t2.o -o %t.exe

#include <iostream>
#include <sycl/ext/intel/esimd.hpp>
#include <sycl/sycl.hpp>

using namespace sycl;
using namespace sycl::ext::intel;
using namespace sycl::ext::intel::esimd;

#define VLEN 8

#ifdef SOURCE1
void run_kernel2(queue &, int *);

int main() {
  queue q;
  int *data = static_cast<int *>(
      malloc_shared(sizeof(int) * VLEN, q.get_device(), q.get_context()));
  for (int i = 0; i < VLEN; i++) {
    data[i] = 4;
  }

  try {
    q.submit([&](handler &cgh) {
      cgh.single_task<class my_kernel0>([=]() SYCL_ESIMD_KERNEL {
        simd<int, VLEN> v0({0, 1, 2, 3, 4, 5, 6, 7});
        simd<int, VLEN> v1 = block_load<int, VLEN>(data);
        v0 = esimd::min(v0, v1); // v0 becomes 0,1,2,3,4,4,4,4
        block_store(data, v0);
      });
    });
    q.wait();
    run_kernel2(q, data);
  } catch (sycl::exception const &e) {
    std::cout << "SYCL exception caught: " << e.what() << '\n';
    return 2;
  }
  int gold[] = {0, 1, 2, 3, 3, 2, 1, 0};

  if (std::memcmp(gold, data, sizeof(int) * VLEN)) {
    std::cout << "FAILED\n";
    for (int i = 0; i < VLEN; i++)
      std::cout << " " << data[i];
    std::cout << "\n";
    return 1;
  }

  std::cout << "Passed\n";
  return 0;
}
#elif defined SOURCE2
void run_kernel2(queue &q, int *data) {
  try {
    q.submit([&](handler &cgh) {
      cgh.single_task<class my_kernel1>([=]() SYCL_ESIMD_KERNEL {
        simd<int, VLEN> v0({7, 6, 5, 4, 3, 2, 1, 0});
        simd<int, VLEN> v1 =
            block_load<int, VLEN>(data); // v1 = 0,1,2,3,4,4,4,4
        v0 = esimd::min(v0, v1);         // v0 becomes 0,1,2,3,3,2,1,0
        block_store(data, v0);
      });
    });
    q.wait();
  } catch (sycl::exception const &e) {
    std::cout << "SYCL exception caught: " << e.what() << '\n';
  }
}
#endif