[Openmp-dev] Use of llvm_omp_target_alloc_shared API
Itaru Kitayama via Openmp-dev
openmp-dev at lists.llvm.org
Sat Aug 7 00:03:15 PDT 2021
1 dimensional case seems to be working.
On Sat, Aug 7, 2021 at 3:31 PM Itaru Kitayama <itaru.kitayama at gmail.com> wrote:
>
> 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]);
> }
>
> }
More information about the Openmp-dev
mailing list