<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
<style type="text/css" style="display:none;"> P {margin-top:0;margin-bottom:0;} </style>
</head>
<body dir="ltr">
<div style="font-family: Calibri, Arial, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
This certainly bypasses <code>is_device_ptr</code>, but it's the same operation fundamentally. If you use
<code>LIBOMPTARGET_INFO</code>=1 you can see that in both cases a 64-bit variable is being copied by-value to the device.</div>
<div id="appendonsend"></div>
<hr style="display:inline-block;width:98%" tabindex="-1">
<div id="divRplyFwdMsg" dir="ltr"><font face="Calibri, sans-serif" style="font-size:11pt" color="#000000"><b>From:</b> Itaru Kitayama <itaru.kitayama@gmail.com><br>
<b>Sent:</b> Sunday, October 17, 2021 3:53 AM<br>
<b>To:</b> Huber, Joseph <huberjn@ornl.gov><br>
<b>Cc:</b> openmp-dev <openmp-dev@lists.llvm.org><br>
<b>Subject:</b> Re: [EXTERNAL] [Openmp-dev] Handling vectors with omp_target_alloc_shared()</font>
<div> </div>
</div>
<div class="BodyFragment"><font size="2"><span style="font-size:11pt;">
<div class="PlainText">This worked for me:<br>
<br>
#include <cstdlib><br>
#include <new><br>
#include <limits><br>
#include <iostream><br>
#include <vector><br>
#include <omp.h><br>
#include <array><br>
#include <memory_resource><br>
<br>
extern "C" {<br>
void* llvm_omp_target_alloc_shared(size_t, int);<br>
}<br>
<br>
template <class T><br>
struct Mallocator<br>
{<br>
typedef T value_type;<br>
<br>
Mallocator () = default;<br>
template <class U> constexpr Mallocator (const Mallocator <U>&) noexcept {}<br>
<br>
[[nodiscard]] T* allocate(std::size_t n) {<br>
if (n > std::numeric_limits<std::size_t>::max() / sizeof(T))<br>
throw std::bad_array_new_length();<br>
<br>
//if (auto p = static_cast<T*>(std::malloc(n*sizeof(T)))) {<br>
//if (T* p =<br>
static_cast<T*>(llvm_omp_target_alloc_shared(n*sizeof(T), 0))) {<br>
if (std::uintptr_t p =<br>
(std::uintptr_t)llvm_omp_target_alloc_shared(n*sizeof(T), 0)) {<br>
//report(p, n);<br>
return reinterpret_cast<T*>(p);<br>
}<br>
<br>
throw std::bad_alloc();<br>
}<br>
<br>
void deallocate(T* p, std::size_t n) noexcept {<br>
//report(p, n, 0);<br>
//std::free(p);<br>
//omp_target_free(p, 0);<br>
}<br>
<br>
private:<br>
void report(T* p, std::size_t n, bool alloc = true) const {<br>
std::cout << (alloc ? "Alloc: " : "Dealloc: ") << sizeof(T)*n<br>
<< " bytes at " << std::hex << std::showbase<br>
<< reinterpret_cast<void*>(p) << std::dec << '\n';<br>
}<br>
};<br>
<br>
template <class T, class U><br>
bool operator==(const Mallocator <T>&, const Mallocator <U>&) { return true; }<br>
template <class T, class U><br>
bool operator!=(const Mallocator <T>&, const Mallocator <U>&) { return false; }<br>
<br>
//typedef std::vector<int, Mallocator<int> > vector<int>;<br>
class A {<br>
public:<br>
A();<br>
double get() {<br>
return v_[0][0];<br>
};<br>
<br>
private:<br>
//std::vector<std::array<double,3>, Mallocator< std::array<double,<br>
3> > > v_ = std::vector<std::array<double, 3>,<br>
Mallocator<std::array<double, 3> > >(1024);<br>
std::vector<std::array<double,3>, Mallocator< std::array<double, 3> > > v_;<br>
};<br>
<br>
A::A() {<br>
<br>
v_.reserve(1024);<br>
for (int i=0;i<1024;i++)<br>
v_[i].fill(123.4);<br>
<br>
}<br>
<br>
int main()<br>
{<br>
std:uintptr_t p = (std::uintptr_t)llvm_omp_target_alloc_shared(sizeof(A), 0);<br>
new (reinterpret_cast<A*>(p)) A();<br>
#pragma omp target parallel for<br>
for (int i=0;i < 1024;i++) {<br>
printf("%f\n", reinterpret_cast<A*>(p)->get());<br>
}<br>
<br>
}<br>
<br>
On Fri, Oct 15, 2021 at 8:57 AM Huber, Joseph <huberjn@ornl.gov> wrote:<br>
><br>
> The pointer value should be copied to the kernel by value, I wouldn't think it would have a large amount of overhead. I don't think there's any other way to tell the runtime to use a device pointer instead of trying to create one using the mapping table.
You could try using a single large pointer to a memory region and extracting it from there.<br>
> ________________________________<br>
> From: Itaru Kitayama <itaru.kitayama@gmail.com><br>
> Sent: Thursday, October 14, 2021 5:30 PM<br>
> To: Huber, Joseph <huberjn@ornl.gov><br>
> Cc: openmp-dev <openmp-dev@lists.llvm.org><br>
> Subject: Re: [EXTERNAL] [Openmp-dev] Handling vectors with omp_target_alloc_shared()<br>
><br>
> Joseph<br>
><br>
> I’d like to avoid using the is_device_pre clause as my app deals with a neural network consisting of tens of thousands of neurons.<br>
><br>
> Itaru.<br>
><br>
> On Fri, Oct 15, 2021 at 2:18 Huber, Joseph <huberjn@ornl.gov> wrote:<br>
><br>
> You need to specify is_device_ptr in your region so the runtime knows to just copy the pointer into the device. Otherwise, it will assume it is on the host and try to look it up in the device mapping table and get nothing. You can use env LIBOMPTARGET_INFO=-1
to see the difference in the assumed mapping type.<br>
> ________________________________<br>
> From: Openmp-dev <openmp-dev-bounces@lists.llvm.org> on behalf of Itaru Kitayama via Openmp-dev <openmp-dev@lists.llvm.org><br>
> Sent: Thursday, October 14, 2021 2:37 AM<br>
> To: openmp-dev <openmp-dev@lists.llvm.org><br>
> Subject: [EXTERNAL] [Openmp-dev] Handling vectors with omp_target_alloc_shared()<br>
><br>
> Hi,<br>
> My test code below works on host, but not on device, could someone<br>
> give me advice if there's a workaround for this?<br>
><br>
> #include <cstdlib><br>
> #include <new><br>
> #include <limits><br>
> #include <iostream><br>
> #include <vector><br>
> #include <omp.h><br>
> #include <array><br>
> #include <memory_resource><br>
><br>
> extern "C" {<br>
> void* llvm_omp_target_alloc_shared(size_t, int);<br>
> }<br>
><br>
> template <class T><br>
> struct Mallocator<br>
> {<br>
> typedef T value_type;<br>
><br>
> Mallocator () = default;<br>
> template <class U> constexpr Mallocator (const Mallocator <U>&) noexcept {}<br>
><br>
> [[nodiscard]] T* allocate(std::size_t n) {<br>
> if (n > std::numeric_limits<std::size_t>::max() / sizeof(T))<br>
> throw std::bad_array_new_length();<br>
><br>
> //if (auto p = static_cast<T*>(std::malloc(n*sizeof(T)))) {<br>
> //if (T* p =<br>
> static_cast<T*>(llvm_omp_target_alloc_shared(n*sizeof(T), 0))) {<br>
> if (std::uintptr_t p =<br>
> (std::uintptr_t)llvm_omp_target_alloc_shared(n*sizeof(T), 0)) {<br>
> //report(p, n);<br>
> return reinterpret_cast<T*>(p);<br>
> }<br>
><br>
> throw std::bad_alloc();<br>
> }<br>
><br>
> void deallocate(T* p, std::size_t n) noexcept {<br>
> //report(p, n, 0);<br>
> //std::free(p);<br>
> //omp_target_free(p, 0);<br>
> }<br>
><br>
> private:<br>
> void report(T* p, std::size_t n, bool alloc = true) const {<br>
> std::cout << (alloc ? "Alloc: " : "Dealloc: ") << sizeof(T)*n<br>
> << " bytes at " << std::hex << std::showbase<br>
> << reinterpret_cast<void*>(p) << std::dec << '\n';<br>
> }<br>
> };<br>
><br>
> template <class T, class U><br>
> bool operator==(const Mallocator <T>&, const Mallocator <U>&) { return true; }<br>
> template <class T, class U><br>
> bool operator!=(const Mallocator <T>&, const Mallocator <U>&) { return false; }<br>
><br>
> //typedef std::vector<int, Mallocator<int> > vector<int>;<br>
> class A {<br>
> public:<br>
> A();<br>
> double get() {<br>
> return this->v_[0][0];<br>
> };<br>
> int f() { return reinterpret_cast<int*>(array_)[0]; };<br>
><br>
> private:<br>
> std::vector<std::array<double,3>, Mallocator< std::array<double, 3> > > v_;<br>
> //std::vector<double, Mallocator< double> > v_;<br>
> std::uintptr_t array_;<br>
> };<br>
><br>
> A::A() {<br>
><br>
> v_.reserve(1024);<br>
> for (int i=0;i<1024;i++)<br>
> v_[i].fill(123.4);<br>
><br>
><br>
> }<br>
><br>
> int main()<br>
> {<br>
> A *p = (A*)llvm_omp_target_alloc_shared(sizeof(A), 0);<br>
> new (p) A();<br>
><br>
> std::vector<std::array<double, 3>, Mallocator<std::array<double,<br>
> 3>>> myv(1024);<br>
> myv[0].fill(0.256);<br>
> #pragma omp target parallel for<br>
> for (int i=0;i < 1024;i++) {<br>
> printf("%f\n", p->get());<br>
> //printf("%f\n", myv[0][0]);<br>
> }<br>
><br>
> }<br>
> _______________________________________________<br>
> Openmp-dev mailing list<br>
> Openmp-dev@lists.llvm.org<br>
> <a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev">https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev</a><br>
</div>
</span></font></div>
</body>
</html>