// Can the iGPU and CPU run relu at the same time? They share DDR5. This runs both
// concurrently on separate buffers and compares combined BW to each alone.
#include <sycl/sycl.hpp>
#include <omp.h>
#include <immintrin.h>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <chrono>
#include <algorithm>
using namespace sycl;

static void cpu_relu(float* __restrict o, const float* __restrict in, size_t n){
  const __m256 z=_mm256_setzero_ps();
  #pragma omp parallel for schedule(static)
  for(size_t j=0;j<n;j+=8) _mm256_stream_ps(o+j,_mm256_max_ps(_mm256_loadu_ps(in+j),z));
  _mm_sfence();
}

template<class F> static double best_of(F f,int reps=8){
  for(int w=0;w<2;w++) f();
  double best=1e30;
  for(int r=0;r<reps;r++){ auto a=std::chrono::high_resolution_clock::now(); f();
    auto b=std::chrono::high_resolution_clock::now();
    best=std::min(best,std::chrono::duration<double>(b-a).count()); }
  return best;
}

int main(){
  queue q{gpu_selector_v};
  printf("iGPU: %s + CPU: %d OpenMP threads\n",
         q.get_device().get_info<info::device::name>().c_str(), omp_get_max_threads());
  const size_t ng=64ull*1024*1024, nc=64ull*1024*1024;
  float* gin =malloc_device<float>(ng,q); float* gout=malloc_device<float>(ng,q);
  std::vector<float> hg(ng); for(size_t i=0;i<ng;++i) hg[i]=(i&1)?-1.f:1.f;
  q.memcpy(gin,hg.data(),ng*4).wait();
  float* cin =(float*)aligned_alloc(64,nc*4); float* cout=(float*)aligned_alloc(64,nc*4);
  for(size_t i=0;i<nc;++i) cin[i]=(i&1)?-1.f:1.f;

  auto gpu_submit=[&](){ q.parallel_for(range<1>(ng),[=](id<1> id){ size_t k=id[0]; float x=gin[k]; gout[k]=x>0.f?x:0.f; }); };
  const double gbg=2.0*ng*4/1e9, gbc=2.0*nc*4/1e9;

  double tg = best_of([&](){ gpu_submit(); q.wait(); });          // iGPU alone
  double tc = best_of([&](){ cpu_relu(cout,cin,nc); });            // CPU alone
  double tb = best_of([&](){ gpu_submit(); cpu_relu(cout,cin,nc); q.wait(); });  // BOTH at once

  double bw_g=gbg/tg, bw_c=gbc/tc, bw_both=(gbg+gbc)/tb;
  printf("\niGPU alone        : %5.1f GB/s\n", bw_g);
  printf("CPU alone         : %5.1f GB/s\n", bw_c);
  printf("BOTH concurrently : %5.1f GB/s combined\n", bw_both);
  printf("  if independent  : %5.1f GB/s (= sum)\n", bw_g+bw_c);
  printf("  scaling vs sum  : %.0f%%  (100%% = no contention, ~50%% = fully memory-bound)\n",
         100.0*bw_both/(bw_g+bw_c));
  free(gin,q); free(gout,q); free(cin); free(cout);
  return 0;
}