Урок 25: параллельный scan — удваивающий kernel
На прошлом уроке мы увидели, что вычисляет prefix sum, и идею удвоения: на каждом раунде каждый элемент прибавляет соседа на расстоянии d, и d удваивается (1, 2, 4...). Здесь мы построим это как настоящий CUDA-kernel. Мы работаем внутри одного block с массивом в разделяемой памяти (__shared__), где
Каждый thread владеет одной позицией в разделяемом массиве. На каждом шаге он сначала читает значение на d позиций левее и откладывает его в сторону, ждёт, пока все прочитают (__syncthreads), только затем прибавляет его к себе, и снова ждёт, пока все запишут, перед следующим шагом. Именно ожидания сохраняют порядок и предотвращают перезапись.
- Hillis-Steele
- Удваивающий scan-kernel за log2(n) шагов: на шаге d каждый thread с tid>=d прибавляет значение из tid-d. d удваивается на каждом шаге.
- tile в разделяемой памяти
- Массив __shared__ для всего block. scan загружает в него входные данные, работает на месте на протяжении всех шагов и в конце записывает в out.
- барьеры чтения/записи
- __syncthreads() после чтения (чтобы все прочитали до любой записи) и после записи (чтобы записи были видны на следующем шаге).
- число шагов log2(n)
- Цикл for(d=1; d<n; d*=2) выполняется log2(n) раз, потому что d удваивается на каждом шаге, пока не достигнет n.
