Обща цел

Трябва да направя няколко съкращения на двустранен граф, представен от два плътни масива за върхове и плътен масив, указващ дали ребро присъства между двете. Да кажем, че два масива са a0[] и a1[] и всички ръбове вървят като e[i0][i1] (т.е. от елементи в a0 до елементи в a1).

Има ~100+100 върха и ~100*100 ръба, така че всяка нишка отговаря за един ръб.

Задача 1: максимално намаление

За всеки връх в a0 искам да намеря максимума от всички върхове (в a1), свързани с него, и след това същото в обратен ред: след като присвоих резултата на масив b0, за всеки връх в a1, искам да намеря максимума b0[i0] на свързаните върхове.

За да направя това, аз:

1) зареждане в споделена памет

    #define DC_NUM_FROM_SHARED 16
    #define DC_NUM_TO_SHARED 16
    __global__ void max_reduce_down(
            Value* value1
        , Value* max_value_in_connected
        , int r0_size, int r1_size
        , bool** connected
        )
    {
        int id_from;
        id_from = blockIdx.x * blockDim.x + threadIdx.x;
        id_to   = blockIdx.y * blockDim.y + threadIdx.y;
        bool within_bounds = (id_from < r0_size) && (id_to < r1_size);

        //load into shared memory
        __shared__ Value value[DC_NUM_TO_SHARED][DC_NUM_FROM_SHARED]; //FROM is the inner (consecutive) dimension
        if(within_bounds)
            value[threadIdx.y][threadIdx.x] = connected[id_to][id_from]? value1[id_to] : 0;
        else
            value[threadIdx.y][threadIdx.x] = 0;
        __syncthreads();

        if(!within_bounds)
            return;

2) намалявам

for(int stride = DC_NUM_TO_SHARED/2; threadIdx.y < stride; stride >>= 1)
{
    value[threadIdx.y][threadIdx.x] = max(value[threadIdx.y][threadIdx.x], dc[threadIdx.y + stride][threadIdx.x]);
    __syncthreads();
}

3) пишете обратно

max_value_connected[id_from] = value[0][threadIdx.x];

Задача 2: най-добрият k

Подобен проблем, но намаляването е само за върхове в a0, трябва да намеря k най-добрите кандидати са избрани от свързани в a1 (k е ~5).

1) Инициализирам споделения масив с нула елементи с изключение на 1-во място

int id_from, id_to;
id_from = blockIdx.x * blockDim.x + threadIdx.x;
id_to   = blockIdx.y * blockDim.y + threadIdx.y;

__shared Value* values[MAX_CHAMPS * CHAMPS_NUM_FROM_SHARED * CHAMPS_NUM_TO_SHARED]; //champion overlaps
__shared int* champs[MAX_CHAMPS * CHAMPS_NUM_FROM_SHARED * CHAMPS_NUM_TO_SHARED]; // overlap champions


bool within_bounds = (id_from < r0_size) && (id_to < r1_size);
int i = threadIdx.y * CHAMPS_NUM_FROM_SHARED + threadIdx.x;
if(within_bounds)
{
    values[i] = connected[id_to][id_from] * values1[id_to];
    champs[i] = connected[id_to][id_from] ? id_to : -1;
}
else
{
    values[i] = 0;
    champs[i] = -1;
}

for(int place = 1; place < CHAMP_COUNT; place++)
{
    i = (place * CHAMPS_NUM_TO_SHARED + threadIdx.y) * CHAMPS_NUM_FROM_SHARED + threadIdx.x;
    values[i] = 0;
    champs[i] = -1;
}
if(! within_bounds)
    return;
__syncthreads();

2) намалете го

for(int stride = CHAMPS_NUM_TO_SHARED/2; threadIdx.y < stride; stride >>= 1)
{
    merge_2_champs(values, champs, CHAMP_COUNT, id_from, id_to, id_to + stride);
    __syncthreads();
}

3) запишете резултатите обратно

for(int place = 0; place < LOCAL_DESIRED_ACTIVITY; place++)
    champs0[place][id_from] = champs[place * CHAMPS_NUM_TO_SHARED * CHAMPS_NUM_FROM_SHARED + threadIdx.x];

Проблем

Как да подредя (транспонирам) елементите в споделения масив, така че достъпът до паметта да използва по-добре кеша? Има ли значение в този момент или има много повече, което мога да спечеля от други оптимизации? Ще бъде ли по-добре да транспонирам матрицата на ръбовете, ако трябва да оптимизирам за Задача 2? (доколкото разбрах в 1 задача има симетрия, така че няма значение).

P.S.

Забавих разгръщането на цикли и извършването на първата редуцираща итерация по време на зареждането, тъй като смятах, че е твърде сложно да се направи, преди да проуча по-прости начини.

За Задача 2 би било хубаво да не се зареждат нулеви елементи, тъй като масивът никога няма да има нужда да расте и да започне да се свива едва след като са направени log k стъпки. Това би го направило k пъти по-компактен в споделената памет! Но се страхувам от получената математика на индекса.

Синтаксис и коректност

Необичайните типове са просто дефинирани ints/chars/etc - AFAIK, в графичните процесори има смисъл да ги компактифицирате колкото е възможно повече. Все още не съм стартирал кода, няма нужда да проверявам за грешки при индексирането.

Освен това използвам CUDA, но се интересувам и от гледна точка на OpenCL, тъй като смятам, че най-доброто решение трябва да е същото и така или иначе ще използвам OpenCL в бъдеще.