Miałem pomysł na redukcję równoległą opartą na warp, ponieważ wszystkie wątki warp są zsynchronizowane z definicji.

Pomysł polegał na tym, że dane wejściowe można zmniejszyć o czynnik 64 (każdy wątek redukuje dwa elementy) bez potrzeby synchronizacji.

Podobnie jak w przypadku oryginalnej implementacji Marka Harrisa redukcja jest stosowana na poziomie bloku, a dane znajdują się w pamięci współdzielonej.http://gpgpu.org/static/sc2007/SC07_CUDA_5_Optimization_Harris.pdf

Stworzyłem jądro do testowania jego wersji i mojej wersji opartej na warp.
Samo jądro całkowicie identycznie przechowuje elementy BLOCK_SIZE w pamięci współdzielonej i wyprowadza swój wynik w unikalnym indeksie bloku w tablicy wyjściowej.

Sam algorytm działa dobrze. Testowany z pełną tablicą do testowania „liczenia”.

Ciało funkcji implementacji:

/**
 * Performs a parallel reduction with operator add 
 * on the given array and writes the result with the thread 0
 * to the given target value
 *
 * @param inValues T* Input float array, length must be a multiple of 2 and equal to blockDim.x
 * @param targetValue float 
 */
__device__ void reductionAddBlockThread_f(float* inValues,
    float &outTargetVar)
{
    // code of the below functions
}

1. Wdrożenie jego wersji:

if (blockDim.x >= 1024 && threadIdx.x < 512)
    inValues[threadIdx.x] += inValues[threadIdx.x + 512];
__syncthreads();
if (blockDim.x >= 512 && threadIdx.x < 256)
    inValues[threadIdx.x] += inValues[threadIdx.x + 256];
__syncthreads();
if (blockDim.x >= 256 && threadIdx.x < 128)
    inValues[threadIdx.x] += inValues[threadIdx.x + 128];
__syncthreads();
if (blockDim.x >= 128 && threadIdx.x < 64)
    inValues[threadIdx.x] += inValues[threadIdx.x + 64];
__syncthreads();

//unroll last warp no sync needed
if (threadIdx.x < 32)
{
    if (blockDim.x >= 64) inValues[threadIdx.x] += inValues[threadIdx.x + 32];
    if (blockDim.x >= 32) inValues[threadIdx.x] += inValues[threadIdx.x + 16];
    if (blockDim.x >= 16) inValues[threadIdx.x] += inValues[threadIdx.x + 8];
    if (blockDim.x >= 8) inValues[threadIdx.x] += inValues[threadIdx.x + 4];
    if (blockDim.x >= 4) inValues[threadIdx.x] += inValues[threadIdx.x + 2];
    if (blockDim.x >= 2) inValues[threadIdx.x] += inValues[threadIdx.x + 1];

    //set final value
    if (threadIdx.x == 0)
        outTargetVar = inValues[0];
}

Zasoby:

Użyto 4 synchronizatorów
12 jeśli używane są instrukcje
11 operacje odczytu + dodawania + zapisu
1 końcowa operacja zapisu
5 użycie rejestru

Wydajność:

pięć testów średnio: ~ 19,54 ms

2. Podejście oparte na warp: (Ta sama funkcja jak powyżej)

/*
 * Perform first warp based reduction by factor of 64
 *
 * 32 Threads per Warp -> LOG2(32) = 5
 *
 * 1024 Threads / 32 Threads per Warp = 32 warps
 * 2 elements compared per thread -> 32 * 2 = 64 elements per warp
 *
 * 1024 Threads/elements divided by 64 = 16
 * 
 * Only half the warps/threads are active
 */
if (threadIdx.x < blockDim.x >> 1)
{
    const unsigned int warpId = threadIdx.x >> 5;
    // alternative threadIdx.x & 31
    const unsigned int threadWarpId = threadIdx.x - (warpId << 5);
    const unsigned int threadWarpOffset = (warpId << 6) + threadWarpId;

    inValues[threadWarpOffset] += inValues[threadWarpOffset + 32];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 16];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 8];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 4];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 2];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 1];
}

// synchronize all warps - the local warp result is stored
// at the index of the warp equals the first thread of the warp
__syncthreads();

// use first warp to reduce the 16 warp results to the final one
if (threadIdx.x < 8)
{
    // get first element of a warp
    const unsigned int warpIdx = threadIdx.x << 6;

    if (blockDim.x >= 1024) inValues[warpIdx] += inValues[warpIdx + 512];
    if (blockDim.x >= 512) inValues[warpIdx] += inValues[warpIdx + 256];
    if (blockDim.x >= 256) inValues[warpIdx] += inValues[warpIdx + 128];
    if (blockDim.x >= 128) inValues[warpIdx] += inValues[warpIdx + 64];

    //set final value
    if (threadIdx.x == 0)
        outTargetVar = inValues[0];
}

Zasoby:

1 użyty sync
7 jeśli oświadczenia
10 operacji zapisu do odczytu
1 końcowa operacja zapisu
5 użycie rejestru

5-bitowe zmiany
1 dodać
1 sub

Wydajność:

pięć testów średnio: 20,82 ms

Wielokrotne testowanie obu jąder na aGeforce 8800 GT 512 mb z256 mb wartości pływaka. I uruchom jądro z256 wątków na blok (100% obłożenia).

Wersja oparta na warp to ~1.28 milisekundy wolniej.

Jeśli przyszłe karty pozwolą na większe rozmiary bloków, podejście oparte na warp nadal nie będzie wymagało dalszych instrukcji synchronizacji, ponieważ maksimum to 4096, które zostaje zredukowane do 64, które zostaje zredukowane przez ostatnią warp do 1

Dlaczego nie jest szybszy? Lub gdzie tkwi wada pomysłu, jądro?

Od wykorzystania zasobów podejście „warp” powinno być przed nami?

Edit1: Poprawiono jądro, że tylko połowa wątków jest aktywna, nie powodując niezwiązanych odczytów, dodano nowe dane dotyczące wydajności