Intel® oneAPI DPC++/C++ Compiler
Talk to fellow users of Intel® oneAPI DPC++/C++ Compiler and companion tools like Intel® oneAPI DPC++ Library, Intel® DPC++ Compatibility Tool, and Intel® Distribution for GDB*
732 Discussions

local_accessor to a shared local memory (with class pointer)

vikinglo
Beginner
521 Views

Hi,

 

Im starting to code with oneapi for GPU, facing some problems on how to load data from GPU global memory to local shared memory. 

 

I have two class: A and B, A have a member X with type of B*, X is init by malloc_device and 

memcpy. class A have a member function (Func0), it need to use X for every work item. It works if X is directly read from global memory, but since X is small, I want to load it to local shared memory for better efficiency. Here is the example code within Func0.
 
inline void A::Fun0(queue myQueue) {
auto self = this;
myQueue.submit([&](handler& h) {
local_accessor<B, 1>
x_local(range(1),h);

h.parallel_for(nd_range<1>(N, range<1>(0)), [=](nd_item<1> item) {
int gid = item.get_global_id(0);
int lid = item.get_local_id(0);

// read X to local shared memory?
if (lid==0) x_local[0] = *self->X;
item.barrier(access::fence_space::local_space);

auto x_ptr = &x_local[0];
// some calculations

});
});
}
 
The code can pass compilor but run into "n illegal memory access was encountered", my guess is this line is not correct: x_local[0] = *self->X; Can someone help me to solve this problem please? Thank you a lot!
0 Kudos
2 Replies
Alex_Y_Intel
Moderator
414 Views

You can reference the material here: https://www.intel.com/content/www/us/en/docs/dpcpp-cpp-compiler/get-started-guide/2024-2/overview.html

If you believe there's a compiler issue, please attach a complete, runnable reproducer to demonstrate your issue, and share the reproducing steps and the compiler version used. 

0 Kudos
vikinglo
Beginner
349 Views

Hi Alex, thanks for the reply, after consulting with local intel engineers, the problem is solved. the problem was how to init Class A on the host side.

0 Kudos
Reply