Сумма массива CUDA - PullRequest
Новая
       15

Сумма массива CUDA

1 голос
/ 20 февраля 2011

У меня есть короткий фрагмент кода, подобный этому: 

typedef struct {
  double sX;
  double sY;
  double vX;
  double vY;
  int rX;
  int rY;
  int mass;
  int species;
  int boxnum;
} particle;

typedef struct {
  double mX;
  double mY;
  double count;
  int rotDir;
  double cX; 
  double cY; 
  int superDir;
} box;
//....
int i;
for(i=0;i<PART_COUNT;i++) {
    particles[i].boxnum = ((((int)(particles[i].sX+boxShiftX))/BOX_SIZE)%BWIDTH+BWIDTH*((((int)(particles[i].sY+boxShiftY))/BOX_SIZE)%BHEIGHT));
}
for(i=0;i<PART_COUNT;i++) {
    //sum the momenta
    boxnum = particles[i].boxnum;
    boxes[boxnum].mX += particles[i].vX*particles[i].mass;
    boxes[boxnum].mY += particles[i].vY*particles[i].mass;
    boxes[boxnum].count++;
}

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

Количество частиц порядка от 10 000 до 10 000 000, а числоящиков стоит на порядок от 1024 до 1048576.

Идеи?

Ответы [ 2 ]

1 голос
/ 21 февраля 2011

В качестве альтернативы вы можете запустить двумерную сетку блоков.

blocks.x = numParticles / threadsPerBlock / repeatPerBlock.

blocks.y = numOfBoxes / 1024;

Каждый блок выполняет атомарные добавления в разделяемой памяти тогда и только тогда, когда boxnum лежит в диапазоне от 1024 * blockIdx.y до 1024 * (blockIdx.y + 1);

Затем следует сокращение вдоль блоков.x

Это может быть, а может и не быть быстрее, чем atomicAdd в глобальной памяти, так как данные читаются как блоки. Сколько раз. Однако это можно исправить, если «частицы» отсортированы по boxnum на этапе сортировки, за которым следует этап разделения.

Может быть несколько других способов сделать это, но, поскольку размер проблемы сильно варьируется, вам может понадобиться написать 2-3 различных метода, оптимизированных для данного диапазона размеров.

1 голос
/ 20 февраля 2011

Вы можете попытаться использовать atomicAdd операции для изменения массива ящиков.Атомарные операции над глобальной памятью очень медленные, но в то же время совершенно невозможно провести какие-либо оптимизации с использованием разделяемой памяти по двум причинам:

  1. При условии, что свойства boxnum частиц particles[0]..particles[n] не упорядочены и не лежат в каких-либо небольших границах (в диапазоне размера блока), вы не можете предсказать, какие блоки загружать из глобальной памяти в общую память.Сначала вам нужно собрать все номера ящиков.
  2. Если вы попытаетесь собрать все номера ящиков, вы не сможете использовать массив со всеми возможными номерами ящиков в качестве индекса, поскольку существует слишком много ящиков, чтобы поместиться в общийобъем памяти.Таким образом, вам придется собирать индексы с очередью (реализованной с помощью массива, указателя на следующий свободный слот и атомарных операций), но тогда у вас все равно будут конфликты, потому что один и тот же номер ящика может появляться несколько раз в вашей очереди.

Вывод: atomicAdd даст вам хотя бы правильное поведение.Попробуйте и проверьте производительность.Если вы не удовлетворены производительностью, подумайте, есть ли другой способ выполнить те же вычисления, которые выиграли бы от совместной памяти.

...