Параллелизация кода с зависимостью

Question

У меня есть код выглядит следующим образом:

Массива a содержит информацию о группе, это элементы принадлежат. то есть элемент i относится к группе a[i]. Каждая группа может содержать два элемента. Я хочу сохранить эту информацию в b, длина которой составляет 2*number групп. Таким образом, значение в b[j] и b[j+1] предоставит мне элементы, которые относятся к группе j/2 (целочисленное деление) и j.

void assign(int *a, int *b){ 
    for(int i = 0; i<N; i++){ 
     int group = a[i]; 
     int posit=2*group; 
     if(b[2*i]!=0){ 
      posit++; 
     } 
     b[posit] = i; 
    } 
} 

N as is clear length of a. 

По умолчанию значение в b [] равно нулю и указывает на отсутствие элементов.

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

Answer 1

В общем, вы можете попытаться просто использовать алгоритм параллельного сортирования на массиве, содержащем пары {i, a[i]}. Может быть, есть более быстрый общий метод, который я не вижу на данный момент ...

В CUDA специально, однако, вы можете воспользоваться тем фактом, что когда у вас есть 2 конфликтующих нити, пишущих 32-битное слово в то же место памяти - один гарантированно будет успешным (вы не знаете, какой из них). Формально он действует как машина Arbitrary CRCW.

Таким образом, вы можете решить вашу проблему в 2 ядра вызовов:

__global__ void assign1(int* a, int* b, int elementCount) { 
    int idx = threadIdx.x + blockIdx.x*blockDim.x; 
    if (idx<elementCount) { 
     int group = a[idx]; 
     b[2*group] = idx; 
    } 
} 

__global__ void assign2(int* a, int* b, int elementCount) { 
    int idx = threadIdx.x + blockIdx.x*blockDim.x; 
    if (idx<elementCount) { 
     int group = a[idx]; 
     if (b[2*group] != idx) 
      b[2*group+1] = idx; 
    } 
} 

__host__ void assign(int* dev_a, int* dev_b, int elementCount) { 
    int gridSize = elementCount/512+1; 
    int blockSize = 512; 
    assign1<<<gridSize,blockSize>>>(dev_a, dev_b, elementCount); 
    assign2<<<gridSize,blockSize>>>(dev_a, dev_b, elementCount); 
} 

В assign1 до 2 нити записи в одной и той же ячейке памяти b[2*group]. Один из этих потоков гарантированно будет успешным. В assign2 поток, который не выполнил запись, повторяет процесс с b[2*group+1].

Этот подход может быть повторен, если в группе было до 3 или 4 элементов, но по мере того, как число становится выше, оно может быстро перестать быть выполнимым.