Intel® oneAPI Data Parallel C++
Support for Intel® oneAPI DPC++ Compiler, Intel® oneAPI DPC++ Library, Intel ICX Compiler , Intel® DPC++ Compatibility Tool, and GDB*

STL algorithm inside kernel

KonradKusiak
Beginner
1,369 Views

Hello. I am trying to use STL sorting algorithm inside the sycl kernel. I initially tried using thrust::sort, then I was experimenting with oneDPL. The problem is that I am trying to use the accessor to local memory. Let me present my problem:

I reserved the chunk of local memory and passed the accessor to it to the kernel. So now I have the local accessor to this memory to play with inside the kernel. This was done like this:

 

 

template <typename T>
using local_accessor = ::sycl::accessor<
    T,
    1,
    ::sycl::access::mode::read_write,
    ::sycl::access::target::local>;

q->submit([&](::sycl::handler& h){
        auto localMem = 
            local_accessor<int>(16, h); 
        SortingKernel kernel(..., localMem); // ... indicate other arguments

        h.parallel_for<class sorting_kernel>(calculatedNdRange, kernel);
}

 

 

 

Now, inside the kernel, after filling in the local memory with some values, I want to sort chunks of it. For example:

 

 

std::sort(m_local_acc.get_pointer() + 2, m_local_acc.get_pointer() + length, myfunction);

 

 

 Now, the problem is that the type that I am getting after calling get_pointer() on the local accessor is the multi_ptr type. And the error that I am getting is that the multi_ptr type does not have the iterator_traits<>::value_type. 

Now, I know that the accessor has the value_type, but I don't know how to start then the sorting from the location of accessor + 2. Because I can't add an int type to the accessor type. 

 

So in general, I think my question would be, how do I sort a vector or array in the local memory, inside the sycl kernel using STL algorithm or similar. 

The error that I am getting when I try using the get_pointer():

 

 

error: no type named 'value_type' in 'std::iterator_traits<sycl::multi_ptr<traccc::triplet, sycl::access::address_space::local_space>>'
      typedef typename iterator_traits<_RandomAccessIterator>::value_type

 

 

 

Labels (1)
0 Kudos
6 Replies
DitiD_Intel
Moderator
1,298 Views

Hi,

Please find the source code below for using std::sort. The std::sort call will implicitly launch a kernel based on the execution policy to sort the given array/vector.

#include <oneapi/dpl/execution>
#include <oneapi/dpl/algorithm>
#include <oneapi/dpl/iterator>
#include <CL/sycl.hpp>
#include <vector>
using std::vector;
int main() {
sycl::range<1> R(4);
vector<int> vec1;
int element;
for (int i=0;i<4;i++)
{
        std::cin>>element;
        vec1.push_back(element);
}
sycl::buffer<int,1> buf(vec1.data(),R);
  auto buf_begin = oneapi::dpl::begin(buf);
  auto buf_end   = oneapi::dpl::end(buf);
  std::sort(oneapi::dpl::execution::dpcpp_default, buf_begin, buf_end);
  auto acc_buf = buf.get_access<sycl::access::mode::read>();
  for (int i=0;i<4; i++)
        std::cout<< acc_buf[i];
  return 0;
}
​


Please note that for using local memory, you need to develop your own kernel since std::sort cannot be used inside the kernel. oneDPL works with the Intel® oneAPI DPC++/C++ Compiler to provide high-productivity APIs to developers, which can minimize Data Parallel C++ (DPC++) programming efforts across devices for high performance parallel applications.

 

For more information on oneDPL, please refer to the below link.

 

https://oneapi-src.github.io/oneDPL/onedpl_gsg.html

 

Thanks and Regards,

Ditipriya.

 

0 Kudos
KonradKusiak
Beginner
1,285 Views

Hi Ditipriya, I can see the source code now. Thank you

 

If I can follow up on my question...

My problem addresses a slightly different issue. I would like to use the sorting algorithm exactly how you did but inside the kernel. So that would mean that it would be nested parallelism. So exactly as you've said, I need to write my own kernel for using local memory.

Then inside this kernel, can I use the std::sort algorithm somehow with iterators pointing at the values in the local memory? 

 

Regards,

 

Konrad

 

 

0 Kudos
KonradKusiak
Beginner
1,277 Views

If there is a way in which I can sort a vector in local memory by doing shuffle operations, I would be interested in such an example also.

 

Kind Regards,

 

Konrad

0 Kudos
KonradKusiak
Beginner
1,294 Views

Hello Ditipriya, thanks for the answer.

 

I don't see any source code though that you've mentioned. Could you try posting it again?

 

Thank you.

0 Kudos
DitiD_Intel
Moderator
1,193 Views

Hi,


We are working on your issue internally, we will get back to you soon.


Thanks and Regards,

Ditipriya.


0 Kudos
KonradKusiak
Beginner
1,181 Views

Hello Ditipriya,

 

In my research on this issue, I found this: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc

which I will try to implement. 

As you mentioned before, the std::sort doesn't work inside the kernel. I tried thrust sort in sequential mode and one dpl sorting algorithms, also in sequential mode. All of those options fail to compile because SYCL kernels don't support recursion which was surprising to me.

I will probably then go with the option that I found but nevertheless if you know of any sorting algorithms similar to thrust::sort or std::sort that could be used inside the SYCL kernel, feel free to let me know. Otherwise, I guess this issue can be closed.

 

KInd Regards,

 

Konrad

0 Kudos
Reply