SYCL buffers within buffers

Question

Say I have the following structure containing buffers:

struct SomeAllocatorCode {
  int* rawData;
  size_t rawDataSize;
  cl::sycl::buffer<int> nestedBuffer;
  SomeAllocatorCode(int* rawData, size_t size): rawData(rawData),
  rawDataSize(rawDataSize), nestedBuffer(rawData, cl::sycl::range<1>(rawDataSize)) {}
};

And then later I want to create a buffer like this:

    int* data = new int[64];
    SomeAllocatorCode* allocator = new SomeAllocatorCode(data, 64);
    cl::sycl::buffer<SomeAllocatorCode> topLevelBuffer(allocator, 
         cl::sycl::range<1>(1));

How would I go about reading nestedBuffer from device code? Is it possible to structure data like this? Is it enough to use data accessors as usual and simply get a data accessor to nestedBuffer after accessing topLevelBuffer using a read accessor?

score 1 · Answer 1 · answered Jul 11 '19 at 08:06

1

I'd suggest understanding how memory buffers and accessors work in SYCL and this will help you to adopt the best approach for what you are trying to do. Here are some links to useful resources:

Basics on buffers and accessors

Memory guide

Code sample showing good practice

answered Jul 11 '19 at 08:06

Rod Burns

2,104
13
24

The approach that I've been trying is, adapt the structure hierarchy so that parent structures keep a sycl::buffer of views to child structures. Each child structure does the same thing. The view of a structure is the same as the original structure but with buffers replaced by accessors. This issue with this approach is that accessors are not nullable, and copying to an invalidly initialized one fails at runtime. Is there a way to use placeholder to work around this? What is the best approach for nullable accessors? – user1832287 Jul 21 '19 at 00:09
Accessors being nullable is required because to obtain a device accessor, we need the command group handler. At this point, the buffers must already be defined. To define a buffer, we need to initialize the memory. To initialize the memory of a device accessor, we need the command group handler (unless they can be nullable). Is there any point along this loop that can be adapted to support multiple levels of indirection? It seems like nullable accessors would be the easiest, since their memory could be initialized within command group handlers before device kernels. – user1832287 Jul 21 '19 at 00:15

score 0 · Accepted Answer · answered Jul 30 '19 at 13:11

0

Here is a worked out example of how to approach this using nested placeholder accessors: Possible ComputeCPP SYCL bug reading nested buffers As noted in that answer, following accessors that are located in a device buffer does not work, so multiple levels of indirection are not possible at the moment.

answered Jul 30 '19 at 13:11

user1832287

329
3
11

1

Here is a more complex use of SYCL for reference: https://github.com/BenjaminTrapani/FunGPU/blob/master/Core/PortableMemPool.hpp It handles device-side allocations and avoids the multiple indirection problem by storing everything in a fixed size block of memory. – user1832287 Aug 25 '19 at 23:45

SYCL buffers within buffers

2 Answers2