Работа с массивом в CUDA

Господа, вопрос таков:
В теле CUDA-ядра необходимо заполнить "рваный" трёхмерный массив, то есть массив неровный ... Существует ли вообще такая возможность? Знаю, что нити сами по себе лежат в блоках (сетке), которые по своей сути являются правильными, прямоугольными массивами! Интересует - может можно обыграть такое положение вещей и всё-таки реализовать нужную мне "фишку"?

Forums: 

Работа с массивом в CUDA

да и вообще хотелось бы знать как быстро работать с многомерными массивами... Все примеры из SDK показывают обращение с одномерными массивами, хотя работа подразумевается зачастую с матрицами! Как выделить память под двухмерный массив - знаю, чо через cudaMallocPitch, но интересует как сделать это через CUDA DRIVER API

Секрет работы очень простой -

Секрет работы очень простой - coalesced access. Т.е. все нити варпа одномоментно идут к подряд расположенным в памяти элементам, подряд, в правильном порядке и с правильным выравниванием (в доке все написано)

Если размер массива неподходящий (т.е не в первой строке не получается выравнивание) - ну значит надо округлить.

давайте на конкретном примере, если не затруднит

Давайте, с основ начнём! т.к. только начинаю разбираться со всей этой "кухней", определенные вещи - тьма!
предположим, для 2-хмерного массива... на CPU выделяем под него место:

  1. int** a = new int*[n];
  2. for(int i=0; i < n; i++)
  3.   a[i] = new int[k];

как такой массив скопировать в глобальную память GPU?
(просьба прикреплять код на Runtime либо Driver)

По процедуре: Если код

По процедуре:
Если код обрамлять <code>.... </code> то будет еще и красиво. Я, пользуясь админскими правами, вам поправил :)

По смыслу:
ну так точно так же. cudaMalloc() на каждую строчку.

Другой вопрос, что и на CPU и на GPU разумнее аллоцировать одним куском, а если они у вас строки неравной длины - то завести отдельно массив индексов строк.

проверьте правильность

Т.е. если я правильно понял, делаем следующее (CUDA.NET) :

  1. CUdeviceptr[] d_a = new CUdeviceptr[n];
  2. for (int i = 0; i < n; i++)
  3.       d_a[i] = cuda.CopyHostToDevice<int>(a[i]);

,
создаём массив указателей и каждому из них ставим в соответствие строку (указатель на строку) нашего массива. А когда параметр передаём на ядро, передаём массив указателей в качестве параметра:
 cuda.SetParameter(func, 0, (uint)d_a.Pointer);
Верно???

В CUDA.NET я не силен. В CUDA

В CUDA.NET я не силен.

В CUDA просто я бы еще перед копированием на девайс - аллоцировал бы на дейвайсе. Возможно, в .NET аллоцируется само, не знаю.

если не сложно, напишите

Аллоцирование массива что подразумевает под собой? я так полагаю происходит выделение памяти на девайсе и создание указателя на эту область... Вопрос - зачем тогда что-то копировать, достаточно лишь параметры ядру передать(в частности указатель на массив)? И если не сложно, напишите код... слова словами, а лучше всё же код посмотреть как всё это реализовывается!

Да все точно так

Да все точно так же.
cudaMalloc, копирование, а потом пачку указателей тоже на Device.

Но я не понимаю - зачем. Почему не иметь один массив и к нему еще таблицу смещений строк?

В моём случае, по моему

В моём случае, по моему мнению, разумнее использовать 2-хмерный массив, исходя из алгоритмических особенностей ... На постоянке нужно будет обращаться к конкретным строкам массива и пробегать по их содержимому! Если вы предлагаете писать 2 массива одномерных по принципу: 1-ый - длины строк, 2-ой содержимое всех строк последовательно записанное, то честно признаться не могу прикинуть будет ли лучше или хуже от этого! Пока что задачу перед собой не ставлю реализовать "быстрый" код, для начала хотелось бы просто переложить на CUDA, там уже и к оптимизации подбираться буду!

Нет, я предлагаю писать два

Нет, я предлагаю писать два массива
- один смещения начала строк. Назовем его offsets
- второй - в общей длинной строке - все строки подряд (возможно, с нужным выравниванием)

И обращаться к A[offsets[i]+j] вместо A[i][j]

В обоих случаях один индирект и особой разницы по производительности не будет. Но
1) offsets - общий на девайсе и на хосте а не два комплекта указателей.
2) оно, на самом деле, более явно.
3) Нет множества аллокаций-деаллокаций.

Естественно, если массив фиксированного размера с фиксированными длинами строк, то это все не нужно, мы обсуждаем именно переменную длину.

ну что ж попробуем ...

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

доброго времени суток... В

доброго времени суток... В общем с массивами разобрался - переделал всё нужное многомерное в одномерное, код конечно стал массивнее, но всё как параметры ядру передаётся и процедуры описанные в ядре вроде бы как выполняются... Проблема очередная : теперь при выполнении функции копирования массивов из DEVICE на HOST начинает ругаться! в чём могут быть проблемы???

  1.  cuda.CopyDeviceToHost<float>(d_Q, Q);

выслушаю очень внимательно любые предположения!!!

И, да, ядро оперирует с

И, да, ядро оперирует с адресами-данными на GPU. Которые туда попадают из RAM путем явного копировния/

А вы представляете как

А вы представляете как работать с многомерными массивами не на gpu, а на cpu?

И что значит "рваный", много нулей? Или вообще недоступные ячейки?
Если речь идёт об разряженных массивах, (где много нулей), можно придумать формат и для них (естественно где нули не хранятся), также есть уже много придуманных (для 2D) http://en.wikipedia.org/wiki/Sparse_matrix#Storing_a_sparse_matrix

А как он "порван", по

А как он "порван", по перфорации или разрыв неровно пошел?

А если серьезно, то вопрос сформулирован так, что я его не понимаю.

Интересует - может можно

Интересует - может можно обыграть такое положение вещей и всё-таки реализовать нужную мне "фишку"

"Фишку" в студию!