// { dg-do run }
// { dg-additional-options "-std=c++23" }
// { dg-additional-options -DMEM_SHARED { target offload_device_shared_as } }

#include <stdlib.h>
#include <time.h>
#include <flat_set>
#include <algorithm>

// MAX should be less than N to ensure that some duplicates occur.
#define N 4000
#define MAX 1000

void init (int data[])
{
  for (int i = 0; i < N; ++i)
    data[i] = rand () % MAX;
}

bool validate (int sum, int data[])
{
  int total = 0;
  for (int i = 0; i < N; ++i)
    total += data[i];
  return sum == total;
}

int main (void)
{
  int data[N];
  std::flat_multiset<int> set;
  int sum = 0;

  srand (time (NULL));
  init (data);

#ifndef MEM_SHARED
  #pragma omp target data map (to: data[ :N]) map (alloc: set)
#endif
    {
      #pragma omp target
	{
#ifndef MEM_SHARED
	  new (&set) std::flat_multiset<int> ();
#endif
	  for (int i = 0; i < N; ++i)
	    set.insert (data[i]);
	}

      #pragma omp target teams distribute parallel for reduction (+:sum)
	for (int i = 0; i < MAX; ++i)
	  sum += i * set.count (i);

#ifdef OMP_USM
      #pragma omp target
	/* Restore the object into pristine state.  In particular, deallocate
	   any memory allocated during device execution, which otherwise, back
	   on the host, we'd SIGSEGV on, when attempting to deallocate during
	   destruction of the object.  */
	__typeof__ (set){}.swap (set);
#endif
#ifndef MEM_SHARED
      #pragma omp target
	set.~flat_multiset ();
#endif
    }

  bool ok = validate (sum, data);
  return ok ? 0 : 1;
}