У меня возникла идея о параллельном сокращении на основе варпа, поскольку все потоки варпа по определению синхронизированы.
Итак, идея заключалась в том, что входные данные можно уменьшить в 64 раза (каждый поток уменьшает два элемента) без необходимости синхронизации.
Как и в исходной реализации Марка Харриса, сокращение применяется на уровне блоков, а данные находятся в общей памяти. http://gpgpu.org/static/sc2007/SC07_CUDA_5_Optimization_Harris.pdf
Я создал ядро, чтобы протестировать его версию и мою версию, основанную на варпе.
Само ядро полностью идентично хранит элементы BLOCK_SIZE в разделяемой памяти и выводит результат по своему уникальному блочному индексу в выходном массиве.
Сам алгоритм работает нормально. Протестировано с полным набором для проверки «подсчета».
Тело функций реализаций:
/**
* 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. Реализация его версии:
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];
}
Ресурсы:
Используются 4 синхреда
Используются 12 операторов if
11 операций чтения + добавления + записи
1 окончательная операция записи
5 использование регистров
Производительность:
среднее время выполнения пяти тестов: ~ 19,54 мс
<сильный>2. Подход на основе деформации: (То же тело функции, что и выше)
/*
* 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];
}
Ресурсы:
1 используется синхропоток
7 операторов if
10 операций чтения и записи
1 окончательная операция записи
5 использование регистра
5 битовых сдвигов
1 добавить
1 подпрограмму
Производительность:
среднее время пяти тестовых прогонов: ~ 20,82 мс
Многократное тестирование обоих ядер на Geforce 8800 GT 512 МБ с 256 МБ значений с плавающей запятой. И работающее ядро с 256 потоками на блок (100 % загрузка).
Версия на основе деформации примерно на 1,28 миллисекунды медленнее.
Если будущие карты позволят увеличить размер блока, подход, основанный на деформации, по-прежнему не будет нуждаться в дальнейшем операторе синхронизации, поскольку максимальное значение составляет 4096, которое уменьшается до 64, а при окончательной деформации уменьшается до 1
Почему не быстрее?, или где ошибка в идее, ядро?
От использования ресурсов подход warp должен быть впереди?
Редактирование 1: ядро исправлено так, что активна только половина потоков, что не приводит к чтению за пределами границ, добавлены новые данные о производительности
