Использование барьеров в SYCL - PullRequest
Новая
       63

Использование барьеров в SYCL

1 голос
/ 17 октября 2019

Я делаю матричное умножение в SYCL и у меня есть рабочий код, в котором я использовал только диапазон параллельно для вместо использования nd_range параллельно для . Теперь я хочу использовать в нем барьеры и, насколько я понимаю, барьеры можно использовать только с nd_range, верно? Я прилагаю часть моего кода, пожалуйста, скажите мне, если это можно сделать без nd_range или какие изменения я должен сделать с nd_range. Спасибо 

queue.submit([&](cl::sycl::handler &cgh) {
    auto A = A_sycl.get_access<cl::sycl::access::mode::read>(cgh);
    auto B = B_sycl.get_access<cl::sycl::access::mode::read>(cgh);
    auto C = C_sycl.get_access<cl::sycl::access::mode::write>(cgh);

    cgh.parallel_for<class test>(
        cl::sycl::range<2>(4, 4), [=](cl::sycl::id<2> id) {
        c_access[id] = A[id] * Y[id.get(1)];
    });

});

1 Ответ

2 голосов
/ 18 октября 2019

Использование nd_range позволяет явно указать локальный диапазон. Чтобы иметь возможность разместить барьер рабочей группы в вашем ядре, вам также потребуется использовать nd_item вместо id , чтобы получить доступ к большему количеству местоположений и размеров идентификаторов, таких как globalи локальный идентификатор, диапазон группы и локальный диапазон, а также примитив синхронизации барьера.

Затем вы можете установить барьер после окончания чтения / записи в локальную память устройства (используя локальный метод доступа только для устройства).

Принимая во внимание, что использование range и id не может дать вам ни одной из этих функций. Он предназначен только для упрощения настройки группы команд и написания ядер глобальной памяти, где вы хотите, чтобы среда выполнения определяла размеры рабочих групп для вас, и у вас есть простой способ индексировать ваши рабочие элементы в отличие от традиционного подхода OpenCL, гдевы всегда должны явно определять NDRange (nd_range в SYCL) независимо от того, насколько просты или сложны ваши ядра.

Вот простой пример, предполагающий, что вы хотите запустить двумерное ядро. 

myQueue.submit([&](cl::sycl::handler& cgh) {
    auto A_ptr = A_buf.get_access<cl::sycl::access::mode::read>(cgh);
    auto B_ptr = B_buf.get_access<cl::sycl::access::mode::read_write>(cgh);
    auto C_ptr = C_buf.get_access<cl::sycl::access::mode::write>(cgh);
    // scratch/local memory for faster memory access to compute the results
    cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
                       cl::sycl::access::target::local>
        C_scratch(range<1>{size}, cgh);

    cgh.parallel_for<example_kernel>(
        cl::sycl::nd_range<2>(range<2>{size >> 3, size >> 3},   // 8, 8
                              range<2>{size >> 4, size >> 4}),  // 4, 4
        [=](cl::sycl::nd_item<2> item) {
          // get the 2D x and y indices
          const auto id_x = item.get_global_id(0);
          const auto id_y = item.get_global_id(1);
          // map the 2D x and y indices to a single linear,
          // 1D (kernel space) index
          const auto width =
              item.get_group_range(0) * item.get_local_range(0);
          // map the 2D x and y indices to a single linear,
          // 1D (work-group) index
          const auto index = id_x * width + id_y;
          // compute A_ptr * B_ptr into C_scratch
          C_scratch[index] = A_ptr[index] * B_ptr[index];
          // wait for result to be written (sync local memory read_write)
          item.barrier(cl::sycl::access::fence_space::local_space);
          // output result computed in local memory
          C_ptr[index] = C_scratch[index];
        });
  });

Я использую одномерное представление данных хоста и буферов SYCL, что объясняет сопоставление двумерных индексов с одним линейным одномерным индексом.

Надеюсь, это объяснение поможет применить эти концепции в вашем случае.

...