#define DPCT_USM_LEVEL_NONE
#include <CL/sycl.hpp>
#include <dpct/dpct.hpp>

#define BUFFER_SIZE 32

typedef struct params{
	int buffer[BUFFER_SIZE];
} params;

dpct::constant_memory<params, 0> input_const1;
dpct::constant_memory<params, 0> input_const2;

void kernel(params input1, params input2, int* output) {
  output[0] = input1.buffer[0]; 
  output[1] = input2.buffer[0];
}

int main() {
  params params1;
  params params2;
  for (int i = 0; i < BUFFER_SIZE; ++i) {
    params1.buffer[i] = 123;
    params2.buffer[i] = 321;
  }

  dpct::dpct_memcpy(input_const1.get_ptr(), &params1, sizeof(params));
  dpct::dpct_memcpy(input_const2.get_ptr(), &params2, sizeof(params));

  printf("Device: %s\n",
    dpct::get_default_queue().get_device().get_info<sycl::info::device::name>().c_str());
  printf("Expected Output: %d, %d\n", params1.buffer[0], params2.buffer[0]);

  int output[2] = { 0, 0 };

  {
    sycl::buffer<int, 1> output_buffer(output, 2);

    dpct::get_default_queue().submit([&](sycl::handler &cgh) {
      auto input_const_acc1 = input_const1.get_access(cgh);
      auto input_const_acc2 = input_const2.get_access(cgh);
      auto output_acc = output_buffer.get_access<sycl::access::mode::write>(cgh);
      
      cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
        [=](sycl::nd_item<3> item_ct1) {
          kernel(input_const_acc1, input_const_acc2, output_acc.get_pointer());
        });
    });
  }

  printf("Real Output: %d, %d\n", output[0], output[1]);

  return 0;
}