Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

32-45 #4

Merged
merged 14 commits into from
Jan 8, 2025
Merged
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
Q40
DanikKul committed Jan 7, 2025
commit 86a8e8a6903b43f86bf1fb3731886abe776015ec
Binary file added .images/reduction-tree.png
Loading
Sorry, something went wrong. Reload?
Sorry, we cannot display this file.
Sorry, this file is invalid so it cannot be displayed.
38 changes: 38 additions & 0 deletions ANSWERS.md
Original file line number Diff line number Diff line change
@@ -1933,6 +1933,44 @@ void kernel(T* a, T* b, int N)
```

# 40. Алгоритм редукции в CUDA. Пример

Редукция это процесс преобразования множества значений в одно результирующее значение.
Примером может являться сумма/произведение/разность и т.д. всех элементов массива.

Сложность реализации алгоритма на CUDA является необходимость построения дерева редукции
на разделяемой памяти.

![alt text](.images/reduction-tree.png)

Пример редукции на CUDA:

```cuda
__global__
void reduceBase(int* inputData, int* outputData)
{
extern __shared__ int smemData[];
// загружаем данные в разделяемую память
const unsigned tid = threadIdx.x;
smemData[tid] = inputData[blockIdx.x * blockDim.x+ threadIdx.x];
// синхронизируемся
__syncthreads();
// выполняем редукцию над элементами массива
for (unsigned s = 1; s < blockDim.x; s <<= 1)
{
if(tid % (s << 1) == 0)
{
smemData[tid] += smemData[tid + s];
}
// синхронизируемся, чтобы гарантировать
// корректность входных данных
// на следующей итерации цикла
__syncthreads();
}
// сохраняем результат вычислений
if (tid == 0) outputData[blockIdx.x] = smemData[0];
}
```

# 41. Алгоритм свертки в CUDA. Пример
# 42. Алгоритм операции инклюзивного scan в CUDA. Пример
# 43. Алгоритм операции эксклюзивного scan в CUDA. Пример