sycl / dpC++ аксессор против global_ptr в объекте функции ядра

Question

Со следующим кодом игрушки с использованием Intel OneAPI beta6.

#include <CL/sycl.hpp>
#include <iostream>

namespace sycl = cl::sycl;

const int SIZE=1;

class Increment_accessor {
  public:
    Increment_accessor(sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer> ptr_) : ptr {ptr_} {}
    void operator()(sycl::item<1> item) {
      ptr[item.get_linear_id()]++;
    }
  private:
    sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer> ptr;
};

class Increment_pointer {
  public:
    Increment_pointer(sycl::global_ptr<int> ptr_) : ptr {ptr_} {} 
    void operator()(sycl::item<1> item) {
      ptr[item.get_linear_id()]++;
    }
  private:
    sycl::global_ptr<int> ptr;
};

int 
main(int argc, char *argv[])
{
  sycl::device dev = sycl::default_selector().select_device();
  sycl::queue q(dev);
  int hbuffer[SIZE] = {};

  {
    sycl::buffer<int, 1> hbuf(hbuffer, sycl::range<1> {SIZE});
    q.submit([&](sycl::handler& cgh) {
        auto harray = hbuf.get_access<sycl::access::mode::read_write, sycl::access::target::global_buffer>(cgh);
        // !!! Uncomment _one_ of the following lines to compile !!!
        //Increment_accessor increment {harray};
        //Increment_pointer increment {harray};
        //Increment_pointer increment {harray.get_pointer()};
        cgh.parallel_for<class kernel1>(
            sycl::range<1> {SIZE}, 
            increment
        );
      }
      ); 
  }

  for (int i=0; i<SIZE; i++) std::cout << "hbuffer[" << i << "]= " << hbuffer[i] << std::endl;
}

Вопрос: почему версии Increment_pointer "неправильные"? Нет ошибок компиляции / выполнения. Вы просто не получаете увеличенный hbuffer в конце. (Я играл с некоторыми похожими версиями, где ptr в operator() заканчивается равным 0x0).

Я все еще учусь думать в "SYCL", поэтому подробные объяснения приветствуются.

Answer 1

Если я правильно понял, вы спрашиваете, почему ваш код работает при использовании Increment_accessor но ломается при использовании Increment_pointer. Или, говоря в более общем смысле, можно ли построить функцию ядра, которая принимает указатели в качестве аргументов вместо средств доступа?

Спецификация SYCL не очень ясна по этому поводу, но раздел 4.7.6.3 дает подсказку:

Средство доступа SYCL может быть средством доступа к устройству, в этом случае оно обеспечивает доступ к данным в рамках функции ядра SYCL, или средством доступа узла, в этом случае оно обеспечивает немедленный доступ к узлу. Если аксессор имеет цель доступа access::target::global_buffer, access::target::constant_buffer, access::target::local,access::target::image или access::target::image_array, тогда он считается аксессор устройства, поэтому может использоваться только в функции ядра SYCL

Таким образом, аксессоры устройства недействительны на хосте. Теперь в ваших версиях, основанных на указателях, вы вызываетеget_pointer()(или полагаться на неявное преобразование метода доступа в указатель, что будет эквивалентно). Но в этот момент вы все еще находитесь в области действия группы команд, а не внутри ядра, то есть кода внутриparallel_for incrementядро. Область действия группы команд всегда оценивается на хосте в SYCL, потому что созданные там средства доступа сообщают среде выполнения SYCL, как строить граф задач и какие узлы зависимостей в графе задач SYCL имеют.

Поэтому мы можем свести вопрос к тому, звонит ли get_pointer()уже квалифицируется как использование аксессуара устройства. Если это так, это означает, что средство доступа к устройству используется за пределами ядра SYCL, что нарушает указанный раздел спецификации и делает этот код незаконным.

Я бы сказал, что это призвание get_pointer() уже соответствует "использованию аксессуара". get_pointer()требует, например, чтобы у аксессора уже было допустимое выделение памяти устройства, на которое он может указывать. Но в области действия группы команд возможно, что это распределение еще даже не существует, поскольку эффективная среда выполнения SYCL может лениво выполнять необходимые распределения в памяти устройства в фоновом режиме прямо перед тем, как они понадобятся. Но в ходе оценки команды группы, команда группа даже не в полной мере представлена во время выполнения SYCL, так как это, по сути, до сих пор находится в стадии их представление. Следовательно, реализация SYCL, как правило, не может гарантировать, чтоget_pointer() уже работает на данном этапе.

Поэтому лучше всего предположить, что средства доступа к устройствам на хосте представляют собой просто описание данных, к которым осуществляется доступ, и приобретают большее значение только как механизм, который фактически разрешает доступ к данным при использовании внутри ядер.

Теперь к более общему вопросу, возможно ли, чтобы ядра принимали указатели в качестве аргументов вместо средств доступа: в SYCL 1.2.1, насколько я понимаю, невозможно иметь средства доступа и каким-то образом преобразовывать их в указатели на узле из-за проблем описано выше. Но если вы вообще не используете аксессоры, вы можете использовать расширение Intel Unified shared memory (USM). Этого расширения пока нет в обычном SYCL 1.2.1, но оно доступно в Intel oneAPI DPC++. USM позволяет явно создавать доступные устройства распределения, которые управляются с помощью указателей. Затем вы можете использовать эти указатели непосредственно в своих ядрах.