[Openmp-dev] Use of llvm_omp_target_alloc_shared API

[Itaru Kitayama via Openmp-dev]

Hi,
Is this supposed to work? It does not on A100.

#include <cstdlib>
#include <new>
#include <limits>
#include <iostream>
#include <vector>
#include <omp.h>

extern "C" {
  void* llvm_omp_target_alloc_shared(size_t, int);
}

template <class T>
struct Mallocator
{
  typedef T value_type;

  Mallocator () = default;
  template <class U> constexpr Mallocator (const Mallocator <U>&) noexcept {}

  [[nodiscard]] T* allocate(std::size_t n) {
    if (n > std::numeric_limits<std::size_t>::max() / sizeof(T))
      throw std::bad_array_new_length();

    //if (auto p = static_cast<T*>(std::malloc(n*sizeof(T)))) {
    if (T* p = static_cast<T*>(llvm_omp_target_alloc_shared(n*sizeof(T), 0))) {
      report(p, n);
      return p;
    }

    throw std::bad_alloc();
  }

  void deallocate(T* p, std::size_t n) noexcept {
    report(p, n, 0);
    std::free(p);
  }

private:
  void report(T* p, std::size_t n, bool alloc = true) const {
    std::cout << (alloc ? "Alloc: " : "Dealloc: ") << sizeof(T)*n
      << " bytes at " << std::hex << std::showbase
      << reinterpret_cast<void*>(p) << std::dec << '\n';
  }
};

template <class T, class U>
bool operator==(const Mallocator <T>&, const Mallocator <U>&) { return true; }
template <class T, class U>
bool operator!=(const Mallocator <T>&, const Mallocator <U>&) { return false; }

int main()
{
  std::vector<std::vector<int, Mallocator<int> >,
Mallocator<std::vector<int, Mallocator<int> > > > v1;
  v1.reserve(1024);

  auto p = v1.data();
  v1[0][0] = 999;
  v1[1][0] = 1;
#pragma omp target parallel for device(0) is_device_ptr(p)
  for (int i=0;i < 2;i++) {
    printf("v1 %d %d\n", i, p[0][0]);
  }

}