приведение методов доступа к указателям C++ в коде ядра (особенно (int (*)[Nelem])
Среда: Ubuntu 18.04, OneAPI beta 6
Полный код приведен ниже, но вот ошибка:
#dpcpp -O2 -g -o so2 so2.cpp -lOpenCL -lsycl
so2.cpp:64:38: error: cannot cast from type 'global_ptr<int>' (aka 'multi_ptr<int, access::address_space::global_space>') to pointer type 'int (*)[nelem]'
int (*xptr)[nelem] = (int (*)[nelem])xaccessor.get_pointer();
^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
1 error generated.
Небольшое объяснение, если вам интересно, почему....
При разработке параллельного кода данных я часто разрабатываю то, что Intel называла "элементарными функциями". Они написаны для работы с одним элементом приложения (то, что SYCL называет рабочим элементом). Я всегда считал, что это проще сделать с базовой средой разработки ПО, легко тестировать и, в более общем смысле, использовать повторно (скаляр, SIMD, CUDA и т. Д.).
После тестирования одного элемента перейти к параллельному использованию данных довольно просто, расширив вызывающий код без необходимости переписывать / повторно тестировать функции:
int x[NELEM]
fn1(x, NELEM)
становится
int x[NPROC][NELEM]
for (int p=0; p<NPROC; p++) fn1(x[p], NELEM);
В ядре SYCL fn1(x[item.get_linear_id()], NELEM); было бы всем, что мне нужно, без необходимости переписывать функции для понимания идентификаторов и / или средств доступа.
Проблема SYCL с приведенным выше кодом заключается в том, что в ядре C++ я не могу преобразовать указатель доступа в 2D-указатель. Это допустимо в приложении C++ (см. Код выше).
Возможно, это плохой способ исправить код, но он упрощает разработку / тестирование кода, который работает для скалярного кода и параллельного кода данных, и сохраняет библиотеки в некоторой степени переносимыми. Он также позволяет обойти ограничение SYCL 3 размерностью буферов / средств доступа.
Во всяком случае, мне любопытно, что подумает настоящий программист SYCL.
Полный код для примера игрушки:
#include <CL/sycl.hpp>
#include <cstdio>
namespace sycl = cl::sycl;
const int Nproc=3;
const int Nelem=4;
/** elemental function **/
void
fn1(int *h, int n)
{
for (int i=0; i<n; i++) h[i] = 10*h[i]+2*i;
}
int
main(int argc, char *argv[])
{
/** Make some memory **/
int x1d[Nproc * Nelem];
for (int j=0; j<Nproc; j++) {
for (int i=0; i<Nelem; i++) x1d[j*Nelem+i] = 10*j+i;
}
printf("1D\n");
for (int i=0; i<Nelem; i++) {
printf("%d : ", i);
for (int j=0; j<Nproc; j++) printf("%d ", x1d[j*Nelem+i]);
printf("\n");
}
/** Reshape it into 2D **/
int (*x2d)[Nelem] = (int (*)[Nelem])x1d;
for (int j=0; j<Nproc; j++) fn1(x2d[j], Nelem);
printf("2D\n");
for (int i=0; i<Nelem; i++) {
printf("%d : ", i);
for (int j=0; j<Nproc; j++) printf("%d ", x2d[j][i]);
printf("\n");
}
/** SYCL setup **/
sycl::device dev = sycl::default_selector().select_device();
std::cout << "Device: "
<< "name: " << dev.get_info<sycl::info::device::name>() << std::endl
<< "vendor: " << dev.get_info<sycl::info::device::vendor>() << std::endl;
sycl::queue q(dev);
{
sycl::buffer<int, 1> xbuffer(x1d, sycl::range<1> {Nproc*Nelem});
q.submit([&](sycl::handler& cgh) {
int nelem = Nelem;
auto xaccessor = xbuffer.get_access<sycl::access::mode::read_write, sycl::access::target::global_buffer>(cgh);
cgh.parallel_for<class k0>(
sycl::range<1> {Nproc},
[=] (sycl::item<1> item) {
int idx = item.get_linear_id();
#if 0
int *xptr = (int *)xaccessor.get_pointer(); // doing this does work so we _can_ get a real pointer
fn1(xptr + nelem*idx, nelem);
#else
int (*xptr)[nelem] = (int (*)[nelem])xaccessor.get_pointer();
//int *ptr = (int *)xaccessor.get_pointer(); // splitting it into two doesn't work either
//int (*xptr)[nelem] = (int (*)[nelem])ptr;
fn1(xptr[idx], nelem);
#endif
}
);
}
);
}
printf("2D SYCL\n");
for (int i=0; i<Nelem; i++) {
printf("%d : ", i);
for (int j=0; j<Nproc; j++) printf("%d ", x1d[j*Nelem+i]);
printf("\n");
}
}
Изменить 1:
За комментарием Иллухада я попытался конкретизировать некоторые альтернативы.
Сначала кажется, что две прокомментированные строки должны делать то, что он предлагает:
int *ptr = (int *)xaccessor.get_pointer();
int (*xptr)[nelem] = (int (*)[nelem])ptr;
но на самом деле это дает эту ошибку:
error: cannot initialize a variable of type 'int (*)[nelem]' with an rvalue of type 'int (*)[nelem]'
int (*xptr)[nelem] = (int (*)[nelem])ptr;
^ ~~~~~~~~~~~~~~~~~~~
добавление "get()" в конец get_pointer дает то же самое.
Любопытно, что устранение части ошибки "инициализация":
int *ptr = (int *)xaccessor.get_pointer().get();
int (*xptr)[nelem];
xptr = (int (*)[nelem])ptr;
Выдает забавную ошибку:
error: incompatible pointer types assigning to 'int (*)[nelem]' from 'int (*)[nelem]'
xptr = (int (*)[nelem])ptr;
^~~~~~~~~~~~~~~~~~~
Так что, если / когда у кого-то будет время, мне все равно любопытно...
1 ответ
Краткий ответ: не проблема SYCL;)
Исходя из вашего редактирования 1, ясно, что если строки
int *ptr = (int *)xaccessor.get_pointer();
int (*xptr)[nelem] = (int (*)[nelem])ptr;
вызвать ошибку преобразования во второй строке, это не может быть проблемой DPC++/SYCL, поскольку задействованы только варианты указателей int, и здесь не происходит ничего, связанного с SYCL.
На самом деле проблема в том, что nelemне является константой времени компиляции. Итак, следующая не-SYCL тестовая программа
int main(){
int nelem = 10;
int* ptr = nullptr;
int (*xptr)[nelem] = (int (*)[nelem])ptr;
}
воспроизводит вашу проблему при компиляции с помощью обычного clang или gcc с -pedantic. Однако по умолчанию gcc поддерживает массивы переменной длины как расширение в C++, поэтому код компилируется, даже если он недействителен C++.
Ваша проблема решена поворотом nelemв константу времени компиляции, как того требует C++. Массивы переменной длины являются частью новых версий C, но не являются частью C++.