topic local_accessor to a shared local memory (with class pointer) in Intel® oneAPI DPC++/C++ Compiler

local_accessor to a shared local memory (with class pointer)

vikinglo — Fri, 05 Jul 2024 10:07:09 GMT

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!

Re: local_accessor to a shared local memory (with class pointer)

Alex_Y_Intel — Mon, 08 Jul 2024 22:22:15 GMT

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.

Re: local_accessor to a shared local memory (with class pointer)

vikinglo — Tue, 16 Jul 2024 08:45:08 GMT

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.